In my .cu file, I have these asm statements
asm volatile("bar.sync 0;");
...
asm volatile("mov.u64 %0, %%clock64;" : "=l"(start) :: "memory");
asm volatile("ld.global.ca.u64 data, [%0];"::"l"(ptr+offset):"memory");
...
asm volatile("bar.sync 0;");
When I view the sass code, I see something like this
/0160/ BSYNC B0 ; /* 0x0000000000007941 /
/ 0x000fea0003800000 /
…
/03b0/ BSYNC B0 ; / 0x0000000000007941 /
/ 0x000fea0003800000 /
/03c0/ CS2R R2, SR_CLOCKLO ; / 0x0000000000027805 /
/ 0x000fd00000015000 /
/03d0/ IADD3 R6, P0, R8, R4, RZ ; / 0x0000000408067210 /
/ 0x000fca0007f1e0ff /
/03e0/ IMAD.X R5, R9, 0x1, R5, P0 ; / 0x0000000109057824 /
/ 0x000fe200000e0605 /
/03f0/ LEA R4, P0, R6, c[0x0][0x180], 0x3 ; / 0x0000600006047a11 /
/ 0x000fc800078018ff /
/0400/ LEA.HI.X R5, R6, c[0x0][0x184], R5, 0x3, P0 ; / 0x0000610006057a11 /
/ 0x000fd000000f1c05 /
/0410/ LDG.E.64.STRONG.CTA R4, [R4] ; / 0x0000000004047381 */
I would like to know if BSYNC is the corresponding instruction for bar.sync 0 or it is something else.