title: Auto selection of free device in exclusive compute mode (when cudaSetDevice is not used) isn't working
body: <p>I recently got access to 4 GPUs (Tesla M2070) and I am trying to get previously working CUDA to use them well. My code runs on one GPU only and I would like to run several instances concurrently. I would like the code to automatically pick an open GPU whenever it is run. </p>
  
  <p>I believe that this should be doable with no problem. The <a href="https://docs.nvidia.com/cuda/cuda-c-programming-guide/#device-selection" rel="nofollow">Cuda Toolkit Documentation</a> says, "If no call to cudaSetDevice() is made, the current device is device 0." However these two threads (thread <a href="https://devtalk.nvidia.com/default/topic/415723/choosing-cuda-device-programmatically/" rel="nofollow">1</a> and <a href="https://devtalk.nvidia.com/default/topic/521594/how-to-get-fastest-free-gpu-cutgetmaxgflopsdeviceid-always-same-device/" rel="nofollow">2</a>) and later in the <a href="https://docs.nvidia.com/cuda/cuda-c-programming-guide/#compute-modes" rel="nofollow">Cuda Toolkit Documentation</a> it is indicated that if the GPUs are set in a compute exclusive mode (exclusive thread or process) that the runtime API should choose a free GPU (instead of always 0) if cudaSetDevice is not called.</p>
  
  <p>This is exactly the situation I am in. The GPUs are set to exclusive thread (see nvidia-smi output below) and I do not call cudaSetDevice (see code below) however when I run two instances of my code one process runs on device zero and the other fails. Both report to be set to device zero. </p>
  
  <p>Does anyone have ideas on why a free device isn't automatically being selected or how to fix it.</p>
  
  <p>Thank you,</p>
  
  <p>Greg</p>
  
  <p>Here is the output of the successful process:</p>
  
  <pre><code>GPU:
  
  ************** CUDA active **************
  
  
  == mean and stdev of uniformly distributed (0,1] random numbers ==
  
  Device calculation:
  
  Device 0: "Tesla M2070"
  mean  = 0.500187
  stdev = 0.282953
  
  *****************************************
  </code></pre>
  
  <p>Here is the output of the failed code:</p>
  
  <pre><code>GPU:
  PottsCUDA: d_kernel_ran.cu:60: void d_ran(int, long int): Assertion `cudaRet==cudaSuccess' failed.
  
  ************** CUDA active **************
  
  
  == mean and stdev of uniformly distributed (0,1] random numbers ==
  
  Device calculation:
  
  Device 0: "Tesla M2070"
  /opt/gridengine/default/spool/compute-0-4/job_scripts/27865: line 47: 45678 Aborted                 (core dumped) ./PottsCUDA
  </code></pre>
  
  <p>Here is the output of nvidia-smi:</p>
  
  <pre><code>Mon Oct  5 15:48:54 2015       
  +------------------------------------------------------+                       
  | NVIDIA-SMI 346.89     Driver Version: 346.89         |                       
  |-------------------------------+----------------------+----------------------+
  | GPU  Name        Persistence-M| Bus-Id        Disp.A | Volatile Uncorr. ECC |
  | Fan  Temp  Perf  Pwr:Usage/Cap|         Memory-Usage | GPU-Util  Compute M. |
  |===============================+======================+======================|
  |   0  Tesla M2070         Off  | 0000:49:00.0     Off |                    0 |
  | N/A   N/A    P0    N/A /  N/A |      9MiB /  5375MiB |      0%    E. Thread |
  +-------------------------------+----------------------+----------------------+
  |   1  Tesla M2070         Off  | 0000:4A:00.0     Off |                    0 |
  | N/A   N/A    P0    N/A /  N/A |      9MiB /  5375MiB |      0%    E. Thread |
  +-------------------------------+----------------------+----------------------+
  |   2  Tesla M2070         Off  | 0000:4B:00.0     Off |                    0 |
  | N/A   N/A    P0    N/A /  N/A |      9MiB /  5375MiB |      0%    E. Thread |
  +-------------------------------+----------------------+----------------------+
  |   3  Tesla M2070         Off  | 0000:4C:00.0     Off |                    0 |
  | N/A   N/A    P0    N/A /  N/A |      9MiB /  5375MiB |      0%    E. Thread |
  +-------------------------------+----------------------+----------------------+
  
  +-----------------------------------------------------------------------------+
  | Processes:                                                       GPU Memory |
  |  GPU       PID  Type  Process name                               Usage      |
  |=============================================================================|
  |  No running processes found                                                 |
  +-----------------------------------------------------------------------------+
  </code></pre>
  
  <p>Here is (an approximately) minimum working example of my code. I am using cuda 7.0 and compiling with arch=compute_20 and code=sm_20.</p>
  
  <p>main.ccp:</p>
  
  <pre><code>#include "d_kernel_ran.h"
  #include &lt;stdio.h&gt; 
  #include &lt;stdlib.h&gt; 
  int  main() 
  {   
  
      // == CUDA check ==
      printf("\n");
      printf("************** CUDA active **************\n");
      printf("\n");
  
      {
          printf("\n");
          printf("== mean and stdev of uniformly distributed (0,1] random numbers ==\n");
          printf("\n");
  
          unsigned long seed_gpu = 170786;
          long ncycles = 10000000;
  
          // - device
          printf("Device calculation:\n");
          d_ran(seed_gpu, ncycles);
          printf("\n");
      }
  
      printf("*****************************************\n\n");
      fflush(stdout);
  
  
      return 0;
  }   
  </code></pre>
  
  <p>d_kernel_ran.cu:</p>
  
  <pre><code>#include "d_kernel_ran.h"
  #include &lt;stdio.h&gt; 
  #include &lt;assert.h&gt;
  #include &lt;cuda.h&gt;
  #include &lt;curand_kernel.h&gt;
  
  __global__ void ranInit( curandState* state, int seed, long N )
  {
      int id = threadIdx.x + blockDim.x * blockIdx.x;
      if (id &lt; N) {
          if (seed &lt; 0) {
              seed = -seed;
          }
          curand_init ( seed, id, 0, &amp;state[id] );
      }
  } 
  
  __global__ void ranGen( curandState* globalState, long N ) 
  {
      int id = threadIdx.x + blockDim.x * blockIdx.x;
      if (id &lt; N) {
          curandState localState = globalState[id];
          float RANDOM = curand_uniform( &amp;localState );
          globalState[id] = localState; 
      }
  }
  
  __global__ void MCMC( curandState* globalState, float* d_ranArray, long N ) 
  {
      int id = threadIdx.x + blockDim.x * blockIdx.x;
      if (id &lt; N) {
          curandState localState = globalState[id];
          float RANDOM = curand_uniform( &amp;localState );
          globalState[id] = localState;
  
          d_ranArray[id] = RANDOM;
      }
  }
  
  void d_ran(int seed, long ncycles) {
  
      long N = ncycles;
  
      // initialization of device arrays
      float* d_ranArray; 
  
      int memsize= N*sizeof(float);
      cudaError_t cudaRet;
  
      cudaRet = cudaMalloc((void**)&amp;d_ranArray, memsize);
      int dev = 0;
      cudaGetDevice(&amp;dev);
      cudaDeviceProp deviceProp;
      cudaGetDeviceProperties(&amp;deviceProp, dev);
      printf("\nDevice %d: \"%s\"\n", dev, deviceProp.name);
      assert(cudaRet==cudaSuccess);
  
  
      // initialization of host arrays
      float* h_ranArray;
  
      h_ranArray = (float*) malloc(memsize);
      assert(h_ranArray);
  
      for (int i=0; i&lt;N; i++) { 
          h_ranArray[i]=(float)0;
      }
  
  
      // copy host -&gt; device
      cudaRet = cudaMemcpy(d_ranArray, h_ranArray, memsize, cudaMemcpyHostToDevice); 
      assert(cudaRet==cudaSuccess); 
  
  
      // initialize blocks
      long threadsPerBlock=512;
  
      long blocksPerGrid = N/threadsPerBlock;
      if (N%threadsPerBlock) blocksPerGrid++;
  
      dim3 dim_grid, dim_block;
  
      dim_grid.x=blocksPerGrid;
      dim_grid.y=1;
      dim_grid.z=1;
  
      dim_block.x=threadsPerBlock;
      dim_block.y=1;
      dim_block.z=1;
  
  
      // setup rngs
      curandState* devStates;
      cudaMalloc( &amp;devStates, N*sizeof( curandState ) );
      ranInit &lt;&lt;&lt; dim_grid, dim_block &gt;&gt;&gt; ( devStates, seed, N );
  
  
      // generate random numbers
      ranGen &lt;&lt;&lt; dim_grid, dim_block &gt;&gt;&gt; ( devStates, N );
  
  
      // perform sum
      MCMC &lt;&lt;&lt; dim_grid, dim_block &gt;&gt;&gt; ( devStates, d_ranArray, N );
  
  
      // copy device -&gt; host
      cudaRet = cudaMemcpy(h_ranArray, d_ranArray, memsize, cudaMemcpyDeviceToHost);
      assert(cudaRet==cudaSuccess); 
  
  
      // free device memory
      cudaRet = cudaFree(devStates); 
      assert(cudaRet==cudaSuccess);
      cudaRet = cudaFree(d_ranArray); 
      assert(cudaRet==cudaSuccess); 
  
  
      // printing result
      float h_agg = (float)0;
      float h_agg2 = (float)0;
      for (int i=0; i&lt;N; i++) { 
          h_agg += h_ranArray[i];
          h_agg2 += h_ranArray[i]*h_ranArray[i];
          #ifdef DEBUG
          printf("%f\n",h_ranArray[i]);
          #endif
      }
  
      float mean = h_agg / (float)N;
      float stdev = std::pow(( h_agg2 / (float)N - mean*mean ),(float)0.5);
  
      printf("mean  = %f\n",mean);
      printf("stdev = %f\n",stdev);
  
      // free host memory
      free(h_ranArray);
  
  }
  </code></pre>
  
  <p>d_kernel_ran.h merely has the declaration of d_ran.</p>