title: CUDA 5.0: CUBIN and CUBLAS_device, compute capability 3.5
body: <p>I'm trying to compile a kernel that uses dynamic parallelism to run CUBLAS to a cubin file.
  When I try to compile the code using the command</p>
  
  <pre><code>nvcc -cubin -m64 -lcudadevrt -lcublas_device -gencode arch=compute_35,code=sm_35 -o test.cubin -c test.cu
  </code></pre>
  
  <p>I get <code>ptxas fatal   : Unresolved extern function 'cublasCreate_v2</code></p>
  
  <p>If I add the <code>-rdc=true</code> compile option it compiles fine, but when I try to load the module using cuModuleLoad I get error 500: CUDA_ERROR_NOT_FOUND. From cuda.h:</p>
  
  <pre><code>/**
   * This indicates that a named symbol was not found. Examples of symbols
   * are global/constant variable names, texture names, and surface names.
   */
  CUDA_ERROR_NOT_FOUND                      = 500,
  </code></pre>
  
  <p><strong>The kernel code:</strong></p>
  
  <pre><code>#include &lt;stdio.h&gt;
  #include &lt;cublas_v2.h&gt;
  extern "C" {
  __global__ void a() {
      cublasHandle_t cb_handle = NULL;
      cudaStream_t stream;
      if( threadIdx.x == 0 ) {
          cublasStatus_t status = cublasCreate_v2(&amp;cb_handle);
          cublasSetPointerMode_v2(cb_handle, CUBLAS_POINTER_MODE_HOST);
          if (status != CUBLAS_STATUS_SUCCESS) {
              return;
          }
          cudaStreamCreateWithFlags(&amp;stream, cudaStreamNonBlocking);
          cublasSetStream_v2(cb_handle, stream);
      }
      __syncthreads();
      int jp;
      double A[3];
      A[0] = 4.0f;
      A[1] = 5.0f;
      A[2] = 6.0f;
      cublasIdamax_v2(cb_handle, 3, A, 1, &amp;jp );
  }
  }
  </code></pre>
  
  <p>NOTE: The scope of <code>A</code> is local, so the data at the pointer given to <code>cublasIdamax_v2</code> is undefined, and so <code>jp</code> ends up as a more or less random value in this code. The correct way to do it would be to have <code>A</code> in global memory.</p>
  
  <p><strong>Host code:</strong></p>
  
  <pre><code>#include &lt;stdio.h&gt;
  #include &lt;cuda.h&gt;
  #include &lt;cuda_runtime_api.h&gt;
  
  int main() {
      CUresult error;
      CUdevice cuDevice;
      CUcontext cuContext;
      CUmodule cuModule;
      CUfunction testkernel;
      // Initialize
      error = cuInit(0);
      if (error != CUDA_SUCCESS) printf("ERROR: cuInit, %i\n", error);
      error = cuDeviceGet(&amp;cuDevice, 0);
      if (error != CUDA_SUCCESS) printf("ERROR: cuInit, %i\n", error);
      error = cuCtxCreate(&amp;cuContext, 0, cuDevice);
      if (error != CUDA_SUCCESS) printf("ERROR: cuCtxCreate, %i\n", error);
      error = cuModuleLoad(&amp;cuModule, "test.cubin");
      if (error != CUDA_SUCCESS) printf("ERROR: cuModuleLoad, %i\n", error);
      error = cuModuleGetFunction(&amp;testkernel, cuModule, "a");
      if (error != CUDA_SUCCESS) printf("ERROR: cuModuleGetFunction, %i\n", error);
      return 0;
  }
  </code></pre>
  
  <p>The host code is compiled using <code>nvcc -lcuda test.cpp</code>.
  If I replace the kernel with a simple kernel (below) and compile it without <code>-rdc=true</code>, it works fine.</p>
  
  <p><strong>Simple working kernel</strong></p>
  
  <pre><code>#include &lt;stdio.h&gt;
  extern "C" {
  __global__ void a() {
      printf("hello\n");
  }
  }
  </code></pre>
  
  <p>Thanks in advance</p>
  
  <ul>
  <li>Soren</li>
  </ul>