title: How to use LLVM IR call mma instruction in PTX
body: <p>What's the proper way to use <code>asm</code> in LLVM IR?</p>
  <p>I tried to add mma intruction which talked <a href="https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-instructions-mma" rel="nofollow noreferrer">here</a> using <code>asm</code> in llvm IR, like the code below.</p>
  <p>However the intrinsic part appearently has something since it generated PTX code like:</p>
  <pre><code>//begin inline asm
  mma.sync.aligned.m16n8k8.row.col.f16.f16.f16.f16 {%0,%1},{%2,%3},{%4},{%5,%6};
  // end inline asm`
  </code></pre>
  <p>It seems the registers are not subsituted into <code>%0</code>, <code>%1</code> ...</p>
  <p>I am not sure whether I return the proper data type of this kind of multiple return value instruction. I made it a vector of <code>&lt;2 x i16&gt;</code>. There is no explaination in the document.</p>
  <pre><code>define &lt;2 x i16&gt; @thefunction(&lt;2 x i16&gt; %0, &lt;1 x i16&gt; %1, &lt;2 x i16&gt; %2) {
  BB:
    %a0 = extractelement &lt;2 x i16&gt; %0, i32 0
    %a1 = extractelement &lt;2 x i16&gt; %0, i32 1
    %b0 = extractelement &lt;1 x i16&gt; %1, i32 0
    %c0 = extractelement &lt;2 x i16&gt; %2, i32 0
    %c1 = extractelement &lt;2 x i16&gt; %2, i32 1
    %value = call &lt;2 x i16&gt; asm &quot;mma.sync.aligned.m16n8k8.row.col.f16.f16.f16.f16 {%0,%1},{%2,%3},{%4},{%5,%6};&quot;, &quot;=r,=r,r,r,r,r,r&quot;(i16 %a0, i16 %a1, i16 %b0, i16 %c0, i16 %c1)
    ret &lt;2 x i16&gt; %value
  }
  </code></pre>