Recogonize lines with predicate (e.g., `@%p1`, `@!%p1`) as instructions:
- `@%p1 ld.global.v4.b32 { %r9, %r10, %r11, %r12 }, [ %rd3 + 0 ];`
- `@%p2 bra $L__BB0_28;`
- `@!%p2 bra $L__BB0_28;`
| Before | After |
| --- | --- |
| <img width="580" height="320" alt="image"
src="https://github.com/user-attachments/assets/c3913b09-ce04-44c5-a4c0-fa69805cf1d7"
/> | <img width="566" height="323" alt="image"
src="https://github.com/user-attachments/assets/08308b3d-df00-4cf1-8321-203c374c27c3"
/> |
From
https://docs.nvidia.com/cuda/parallel-thread-execution/#instruction-statements:
> Instructions have an optional guard predicate which controls
conditional execution. The guard predicate follows the optional label
and precedes the opcode, and is written as @p, where p is a predicate
register. The guard predicate may be optionally negated, written as @!p.
Implementation-wise, added `(@!?%\w+\s+)?` to existing default
`hasOpcodeRe` (`/^\s*(%[$.A-Z_a-z][\w$.]*\s*=\s*)?[A-Za-z]/`)