# PTX assembly language reference does one exist, or plans to release one?

**URL:** <https://forums.developer.nvidia.com/t/ptx-assembly-language-reference-does-one-exist-or-plans-to-release-one/465>\
**Category:** CUDA Programming and Performance\
**Created:** [March 25, 2007, 1:56am UTC](https://forums.developer.nvidia.com/t/ptx-assembly-language-reference-does-one-exist-or-plans-to-release-one/465 "2007-03-25T01:56:51Z")\
**Posts on this page:** 7\
**Page:** 1

<div class="post-metadata">

**Author:** ![jhanweck](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@jhanweck](https://forums.developer.nvidia.com/u/jhanweck)\
**Post date:** [March 25, 2007, 1:56am UTC](https://forums.developer.nvidia.com/t/ptx-assembly-language-reference-does-one-exist-or-plans-to-release-one/465/1 "2007-03-25T01:56:51Z")

</div>

Are there plans to release a PTX assembly language manual? Or, does one exist somewhere?

It would be helpful for general debugging… Also, I can see writing assembly to hand optimize register use, and so forth.

That said, would it be possible to allow “inlined” assembly in Cuda code? As device functions, perhaps? (I seem to recall some old PC C compilers – maybe Turbo C? – allowed something like this.)

---

<div class="post-metadata">

**Author:** ![Stewie](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Stewie](https://forums.developer.nvidia.com/u/Stewie)\
**Post date:** [March 26, 2007, 3:08pm UTC](https://forums.developer.nvidia.com/t/ptx-assembly-language-reference-does-one-exist-or-plans-to-release-one/465/2 "2007-03-26T15:08:59Z")

</div>

I believe the .ptx file is pre-optimized code, not sure that hand optimizing the register usage there would have an impact.  
(e.g., I wrote some code that used 100 registers, but the final .cubin file says it only used 21)

---

<div class="post-metadata">

**Author:** ![Mark\_Harris](https://sea2.discourse-cdn.com/nvidia/user_avatar/forums.developer.nvidia.com/mark_harris/32/23175_2.png) [@Mark\_Harris](https://forums.developer.nvidia.com/u/Mark_Harris)\
**Post date:** [March 27, 2007, 1:11pm UTC](https://forums.developer.nvidia.com/t/ptx-assembly-language-reference-does-one-exist-or-plans-to-release-one/465/3 "2007-03-27T13:11:05Z")

</div>

Stewie is correct. PTX is a virtual machine assembly language, so it does not optimize for exact hardware instruction set or register usage. This way we can target different GPUs with the same assembly-level interface. The code is then further optimized at load time by the GPU driver.

Mark

---

<div class="post-metadata">

**Author:** ![jhanweck](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@jhanweck](https://forums.developer.nvidia.com/u/jhanweck)\
**Post date:** [March 28, 2007, 2:45pm UTC](https://forums.developer.nvidia.com/t/ptx-assembly-language-reference-does-one-exist-or-plans-to-release-one/465/4 "2007-03-28T14:45:26Z")

</div>

Interesting… good to know.

Still, a PTX cheat sheet would be nice to have (e.g., for debugging).

For example: does mul.lo.u32 multiply (unsigned) the low 16 bits of the operands together into an 32-bit word?

In this snippet, which is generating an index into global memory array of ints, parm\_output. The index is threadIdx.x + blockIdx.x \* blockDim.x.

$r15 gets the product blockIdx.x \* blockDim.x using mul.lo.u32, which makes sense because the block dimensions are limited to 65535.

$r16 gets the final index, adding threadIdx.x to the product above. This value could very well be larger than 2^16.

$r17 gets the offset into global memory, again using mul.lo.u32 to multiply the index ($r16) by sizeof(int) (4).

But since $r16 could be greater than 2^16, is mul.lo.u32 appropriate? Or do I misunderstand mul.lo.u32?

```auto
# parm_output[threadIdx.x + blockIdx.x * blockDim.x] = intval;

ld.param.u32  $r11, %parm_output;	#  id:30 %parm_output+0x0

cvt.u32.u16  $r12, %tid.x;    	#  

cvt.u32.u16  $r13, %ctaid.x;  	#  

cvt.u32.u16  $r14, %ntid.x;    #  

mul.lo.u32  $r15, $r13, $r14;  #  blockIdx.x * blockDim.x

add.u32  $r16, $r12, $r15;    	#  threadIdx.x + (blockIdx.x * blockDim.x)

mul.lo.u32  $r17, $r16, 4;    	#  ... * sizeof(int)  ???

add.u32  $r18, $r11, $r17;    	#  

st.global.s32  [$r18+0], $r10;	#  id:34

```

---

<div class="post-metadata">

**Author:** ![prkipfer](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@prkipfer](https://forums.developer.nvidia.com/u/prkipfer)\
**Post date:** [March 29, 2007, 10:29am UTC](https://forums.developer.nvidia.com/t/ptx-assembly-language-reference-does-one-exist-or-plans-to-release-one/465/5 "2007-03-29T10:29:04Z")

</div>

I think the .lo. is not an indicator for the number of bits used. For example the \_\_mul24 function is translated into mul24.lo.s32 opcode.

Peter

---

<div class="post-metadata">

**Author:** ![jhanweck](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@jhanweck](https://forums.developer.nvidia.com/u/jhanweck)\
**Post date:** [March 29, 2007, 1:56pm UTC](https://forums.developer.nvidia.com/t/ptx-assembly-language-reference-does-one-exist-or-plans-to-release-one/465/6 "2007-03-29T13:56:50Z")

</div>

> [@](#):
>
> I think the .lo. is not an indicator for the number of bits used. For example the \_\_mul24 function is translated into mul24.lo.s32 opcode.
> 
> Peter
> 
> [snapback]177294[/snapback]

Hence, the need for a cheatsheet… :P

---

<div class="post-metadata">

**Author:** ![Zvi](https://developer.download.nvidia.com/images/forums/profile-default-devtalk-84.png) [@Zvi](https://forums.developer.nvidia.com/u/Zvi)\
**Post date:** [March 29, 2009, 10:25am UTC](https://forums.developer.nvidia.com/t/ptx-assembly-language-reference-does-one-exist-or-plans-to-release-one/465/7 "2009-03-29T10:25:48Z")

</div>

> [@](#):
>
> Are there plans to release a PTX assembly language manual? Or, does one exist somewhere?

See PTX: Parallel Thread Execution, ISA Version 1.1

and PTX: Parallel Thread Execution, ISA Version 1.2.