title: cublas matrix inversion from device
body: <p>I am trying to run a matrix inversion from the device.  This logic works fine if called from the host.</p>
  
  <p>Compilation line is as follows (Linux):</p>
  
  <pre><code>nvcc -ccbin g++ -arch=sm_35 -rdc=true simple-inv.cu -o simple-inv -lcublas_device -lcudadevrt
  </code></pre>
  
  <p>I get the following warning that I cannot seem to resolve. (My GPU is Kepler.  I don't know why it is trying to link to Maxwell routines. I have Cuda 6.5-14):</p>
  
  <pre><code>nvlink warning : SM Arch ('sm_35') not found in '/usr/local/cuda/bin/../targets/x86_64-linux/lib/libcublas_device.a:maxwell_sm50_sgemm.o'
  </code></pre>
  
  <p>The program runs with:</p>
  
  <pre><code>handle 0 n = 3
  simple-inv.cu:63 Error [an illegal memory access was encountered]
  </code></pre>
  
  <p>The test program is as follows:</p>
  
  <pre><code>#include &lt;stdio.h&gt;
  #include &lt;stdlib.h&gt;
  #include &lt;math.h&gt;
  #include &lt;cuda_runtime.h&gt;
  #include &lt;cublas_v2.h&gt;
  
  #define PERR(call) \
    if (call) {\
     fprintf(stderr, "%s:%d Error [%s] on "#call"\n", __FILE__, __LINE__,\
        cudaGetErrorString(cudaGetLastError()));\
     exit(1);\
    }
  #define ERRCHECK \
    if (cudaPeekAtLastError()) { \
      fprintf(stderr, "%s:%d Error [%s]\n", __FILE__, __LINE__,\
         cudaGetErrorString(cudaGetLastError()));\
      exit(1);\
    }
  
  __global__ void
  inv_kernel(float *a_i, float *c_o, int n)
  { 
    int p[3], info[1], batch;
    cublasHandle_t hdl;
    cublasStatus_t status = cublasCreate_v2(&amp;hdl);
    printf("handle %d n = %d\n", status, n);
  
    info[0] = 0;
    batch = 1;
    float *a[] = {a_i};
    const float *aconst[] = {a_i};
    float *c[] = {c_o};
    // See
    // http://docs.nvidia.com/cuda/pdf/CUDA_Dynamic_Parallelism_Programming_Guide.pdf
    //http://stackoverflow.com/questions/27094612/cublas-matrix-inversion-from-device
  
    status = cublasSgetrfBatched(hdl, n, a, n, p, info, batch);
    __syncthreads();
    printf("rf %d info %d\n", status, info[0]);
    status = cublasSgetriBatched(hdl, n, aconst, n, p,
        c, n, info, batch);
    __syncthreads();
    printf("ri %d info %d\n", status, info[0]);
  
    cublasDestroy_v2(hdl);
    printf("done\n");
  }
  static void
  run_inv(float *in, float *out, int n)
  {
    float *a_d, *c_d;
  
    PERR(cudaMalloc(&amp;a_d, n*n*sizeof(float)));
    PERR(cudaMalloc(&amp;c_d, n*n*sizeof(float)));
    PERR(cudaMemcpy(a_d, in, n*n*sizeof(float), cudaMemcpyHostToDevice));
  
    inv_kernel&lt;&lt;&lt;1, 1&gt;&gt;&gt;(a_d, c_d, n);
  
    cudaDeviceSynchronize();
    ERRCHECK;
  
    PERR(cudaMemcpy(out, c_d, n*n*sizeof(float), cudaMemcpyDeviceToHost));
    PERR(cudaFree(a_d));
    PERR(cudaFree(c_d));
  }
  
  int
  main(int argc, char **argv)
  {
    float c[9];
    float a[] = {
      1,   2,   3,
      0,   4,   5,
      1,   0,   6 };
  
    run_inv(a, c, 3);
    return 0;
  }
  </code></pre>
  
  <p>I have followed the guide at <a href="http://docs.nvidia.com/cuda/cublas/index.html#device-api" rel="nofollow">http://docs.nvidia.com/cuda/cublas/index.html#device-api</a> section 2.1.9, but I suspect I have overlooked something.</p>
  
  <p>Note: Edited on 11/24 to use correct pointer inputs.  This still reports illegal memory access inside the kernel.</p>