title: HIP-Clang Inline Assembly
body: <p>What is the Hip-Clang equivalent of this CUDA function?</p>
  <pre><code>__device__ __forceinline__ uint32_t add_cc(uint32_t a, uint32_t b)
  {
     uint32_t r;
     asm volatile (&quot;add.cc.u32 %0, %1, %2;&quot; : &quot;=r&quot;(r) : &quot;r&quot;(a), &quot;r&quot;(b));
     return r;
  }
  </code></pre>
  <p>I'm porting a CUDA project to HIP-Clang that contains inline PTX assembly. The function is used to implement multi-precision addition in the NVIDIA GPU. I tried:</p>
  <pre><code>asm volatile (&quot;add.cc.u32 %0, %1, %2;&quot; : &quot;=r&quot;(r) : &quot;r&quot;(a), &quot;r&quot;(b)); //invalid instruction
  asm volatile (&quot;V_ADD_CO_U32 %0, %1, %2;&quot; : &quot;=r&quot;(r) : &quot;r&quot;(a), &quot;r&quot;(b)); //invalid operand for instruction
  asm volatile (&quot;V_ADD_CO_U32 %0, %1, %2;&quot; : &quot;=v&quot;(r) : &quot;v&quot;(a), &quot;v&quot;(b)); //operands are not valid for this GPU or mode
  </code></pre>
  <p>The target hardware is RX 6800. AMD clang version 14.0.0.<br />
  Is RDNA2 the correct <a href="https://developer.amd.com/wp-content/resources/RDNA2_Shader_ISA_November2020.pdf" rel="nofollow noreferrer">instruction set reference</a>?<br />
  Is this <a href="https://llvm.org/docs/AMDGPUUsage.html#assembler" rel="nofollow noreferrer">LLVM user guide to AMDGPU backend</a> an applicable reference?</p>