instruction meanings

[Jianbin Fang]

Hi all,

When reading the assembly code in svn/test/, I find that you have used
some special-purpose registers: SR_ClockLo, SR_Tid_x. How can know the
meaning of these special registers? Do you have any references for
that?

Also, I find that some instructions, not listed in the cuobjdump, are
used, e.g., LOP.AND, SHR, etc. Fortunately, these instructions are
listed in the asfermi wiki page, which means these instructions have
to be detected by ourselves?

Thanks,

Jianbin

[Hou Yunqing]

Hi Jianbin,

Here's for the special registers: http://code.google.com/p/asfermi/wiki/OpcodeMiscellaneous#S2R

You get their meaning by 

1. reading NV docs. The information is scattered all over so there's no fixed place to find it all

2. compare your CUDA source code to the output of cuobjdump, and

3. do some guess work, the names are somewhat self-explanatory

SR_ClockLo is the (lower 32 bits) scheduler clock number, which is incremented at every scheduler clock.

SR_Tid_X really is threadIdx.x

As for the instructions, I think I had gotten them all here for Fermi: http://code.google.com/p/asfermi/wiki/nanb

And yes, the information in that page was obtained through manual/semi-automated work

I'd be surprised if there's anything I had missed. I did get all the strings in the cuobjdump executable dumped, and checked through all the things that looked like instruction names and I think I had them all listed in the nanb page.

Regards,

Yunqing