title: How to use LLVM IR Builder to insert cudaPrefetch Instruction
body: <p>I am trying to insert an <strong>cuda Prefetch</strong> instruction in LLVM IR</p>
  <pre class="lang-cpp prettyprint-override"><code>// The origin cuda prefetch API
  cudaMemPrefetchAsync ( const void* devPtr, size_t count, int  dstDevice, cudaStream_t stream = 0 )
  
  // LLVM IR relevent to cudaMemPrefetchAsync
  // %9 = alloca float*, align 8
  // %16 = bitcast float** %9 to i8**
  // %17 = call i32 @cudaMallocManaged(i8** nonnull %16, i64 4096, i32 1)
  // %20 = load i8*, i8** %16, align 8, !tbaa !7
  // %21 = call i32 @cudaMemPrefetchAsync(i8* %20, i64 4096, i32 0, %struct.CUstream_st* null)
  </code></pre>
  <p>Here's what I've done</p>
  <pre class="lang-cpp prettyprint-override"><code>// Get function declaration of cudaMemPrefetchAsync
  if (mod_func-&gt;getName().str().find(cudaMemPrefetchAsyncName) != string::npos) {
      cuda_prefetch_func = &amp;(*mod_func);
  }
  
  // Generate prefetch data size and device ID
  IntegerType* int_type32 = IntegerType::get(context, 32);
  IntegerType* int_type64 = IntegerType::get(context, 64);
  ConstantInt* device_id = ConstantInt::get(int_type32, 0, false);
  ConstantInt* alloca_size = ConstantInt::get(int_type64, 4096, false);
  
  // Get variable address from cudaMallocManaged function parameter
  if(fname == cudaMallocManagedName) {
    val = cast&lt;CallInst&gt;(&amp;inst)-&gt;getArgOperand(0);
  }
  
  Value* args[] = { val, alloca_size, device_id};
  CallInst *MI = builder.CreateCall(cuda_prefetch_func, args);
  </code></pre>
  <p>But I don'know how to generate the <code>%struct.CUstream_st* null</code> and the address of the <strong>devPtr</strong> is incompatible, since cudaMallocManaged is <code>i8 **</code> and cudaMemPrefetchAsync is <code>i8 *</code>.</p>
  <p>Undoubtedly, I'll get the following error message.</p>
  <pre><code>Incorrect number of arguments passed to called function!
  %1 = call i32 @cudaMemPrefetchAsync(i8** %20, i64 4096, i32 0)
  </code></pre>
  <p>How can I generate <code>%struct.CUstream_st* null</code> and bitcast <code>i8 **</code> to <code>i8 *</code> (or where should I get the i8 * pointer).</p>