title: cooperative_groups::this_grid() causes any CUDA API call to return 'unknown error'
body: <p>Following the same steps in <a href="https://github.com/NVIDIA/cuda-samples/blob/master/Samples/conjugateGradientMultiBlockCG/conjugateGradientMultiBlockCG.cu#L436" rel="nofollow noreferrer">CUDA samples</a> to launch a kernel and sync across the grid using <code>cooperative_groups::this_grid().sync()</code> causes any CUDA API call to fails. While using 
   <code>cooperative_groups::this_thread_block().sync()</code> works fine and gives correct results.</p>
  
  <p>I used the following code and <code>CMakeLists.txt</code> (cmake version 3.11.1) to test it using CUDA 10 on TITAN V GPU (Driver Version 410.73) with Ubuntu 16.04.5 LTS. The code is also available on <a href="https://github.com/Ahdhn/test_cg" rel="nofollow noreferrer">github</a> in order to make it easy to reproduce the error.</p>
  
  <p>The code reads an array and then reverses it (from <code>[0 1 2 ... 9]</code> to <code>[9 8 7 ... 0]</code>). In order to do this, each thread reads a single element from the array, sync, and then writes its element to the right destination. The code can be easily modified to ensure that <code>this_thread_block().sync()</code> works fine. Simply change <code>arr_size</code> to be less 1024 and use <code>cg::thread_block barrier = cg::this_thread_block();</code> instead. </p>
  
  <p><strong>test_cg.cu</strong></p>
  
  <pre><code>#include &lt;cuda_runtime_api.h&gt;
  #include &lt;stdio.h&gt;
  #include &lt;stdint.h&gt;
  #include &lt;cstdint&gt;
  #include &lt;numeric&gt;
  #include &lt;cuda.h&gt;
  #include &lt;cooperative_groups.h&gt;
  namespace cg = cooperative_groups;
  
  //********************** CUDA_ERROR
  inline void HandleError(cudaError_t err, const char *file, int line) {
      //Error handling micro, wrap it around function whenever possible
      if (err != cudaSuccess) {
          printf("\n%s in %s at line %d\n", cudaGetErrorString(err), file, line);
  
  #ifdef _WIN32
          system("pause");
  #else
          exit(EXIT_FAILURE);
  #endif
      }
  }
  #define CUDA_ERROR( err ) (HandleError( err, __FILE__, __LINE__ ))
  //******************************************************************************
  
  
  //********************** cg kernel 
  __global__ void testing_cg_grid_sync(const uint32_t num_elements,
      uint32_t *d_arr){
      uint32_t tid = threadIdx.x + blockDim.x*blockIdx.x;
  
      if (tid &lt; num_elements){
  
          uint32_t my_element = d_arr[tid];
  
          //to sync across the whole grid 
          cg::grid_group barrier = cg::this_grid();
  
          //to sync within a single block 
          //cg::thread_block barrier = cg::this_thread_block();
  
          //wait for all reads 
          barrier.sync();
  
          uint32_t tar_id = num_elements - tid - 1;
  
          d_arr[tar_id] = my_element;
      }
  }
  //******************************************************************************
  
  
  //********************** execute  
  void execute_test(const int sm_count){
  
      //host array 
      const uint32_t arr_size = 1 &lt;&lt; 20; //1M 
      uint32_t* h_arr = (uint32_t*)malloc(arr_size * sizeof(uint32_t));
      //fill with sequential numbers
      std::iota(h_arr, h_arr + arr_size, 0);
  
      //device array 
      uint32_t* d_arr;
      CUDA_ERROR(cudaMalloc((void**)&amp;d_arr, arr_size*sizeof(uint32_t)));
      CUDA_ERROR(cudaMemcpy(d_arr, h_arr, arr_size*sizeof(uint32_t),
          cudaMemcpyHostToDevice));
  
      //launch config
      const int threads = 512;
  
      //following the same steps done in conjugateGradientMultiBlockCG.cu 
      //cuda sample to launch kernel that sync across grid 
      //https://github.com/NVIDIA/cuda-samples/blob/master/Samples/conjugateGradientMultiBlockCG/conjugateGradientMultiBlockCG.cu#L436
  
      int num_blocks_per_sm = 0;
      CUDA_ERROR(cudaOccupancyMaxActiveBlocksPerMultiprocessor(&amp;num_blocks_per_sm,
          (void*)testing_cg_grid_sync, threads, 0));
  
      dim3 grid_dim(sm_count * num_blocks_per_sm, 1, 1), block_dim(threads, 1, 1);
  
      if(arr_size &gt; grid_dim.x*block_dim.x){
           printf("\n The grid size (numBlocks*numThreads) is less than array size.\n");
           exit(EXIT_FAILURE);
      }
      printf("\n Launching %d blocks, each containing %d threads", grid_dim.x,
          block_dim.x);
  
      //argument passed to the kernel     
      void *kernel_args[] = {
          (void *)&amp;arr_size,
          (void *)&amp;d_arr, };
  
  
      //finally launch the kernel 
      cudaLaunchCooperativeKernel((void*)testing_cg_grid_sync,
          grid_dim, block_dim, kernel_args);
  
  
      //make sure everything went okay
      CUDA_ERROR(cudaGetLastError());
      CUDA_ERROR(cudaDeviceSynchronize());
  
  
      //get results on the host 
      CUDA_ERROR(cudaMemcpy(h_arr, d_arr, arr_size*sizeof(uint32_t),
          cudaMemcpyDeviceToHost));
  
      //validate 
      for (uint32_t i = 0; i &lt; arr_size; i++){
          if (h_arr[i] != arr_size - i - 1){
              printf("\n Result mismatch in h_arr[%u] = %u\n", i, h_arr[i]);
              exit(EXIT_FAILURE);
          }
      }
  }
  //******************************************************************************
  
  int main(int argc, char**argv) {
  
      //set to Titan V
      uint32_t device_id = 0;
      cudaSetDevice(device_id);
  
      //get sm count 
      cudaDeviceProp devProp;
      CUDA_ERROR(cudaGetDeviceProperties(&amp;devProp, device_id));
      int sm_count = devProp.multiProcessorCount;
  
      //execute 
      execute_test(sm_count);
  
      printf("\n Mission accomplished \n");
      return 0;
  }
  </code></pre>
  
  <p><strong>CMakeLists.txt</strong></p>
  
  <pre><code>cmake_minimum_required(VERSION 3.8 FATAL_ERROR)
  
  set(PROJECT_NAME "test_cg")
  project(${PROJECT_NAME} LANGUAGES CXX CUDA)  
  
  #default build type is Release
  if (CMAKE_BUILD_TYPE STREQUAL "")
      set(CMAKE_BUILD_TYPE Release)
  endif ()
  
  SET(CUDA_SEPARABLE_COMPILATION ON)
  
  ########## Libraries/flags Starts Here ######################
  find_package(CUDA REQUIRED)
  include_directories("${CUDA_INCLUDE_DIRS}")
  set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS}; -lineinfo; -std=c++11; -expt-extended-lambda; -O3; -use_fast_math; -rdc=true;)
  set(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS};-gencode=arch=compute_70,code=sm_70) #for TITAN V
  set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS}")
  set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -m64 -Wall -std=c++11")
  ########## Libraries/flags Ends Here ######################
  
  
  ########## inc/libs/exe/features Starts Here ######################
  set(CMAKE_INCLUDE_CURRENT_DIR ON)
  CUDA_ADD_EXECUTABLE(${PROJECT_NAME} test_cg.cu)
  target_compile_features(${PROJECT_NAME} PUBLIC cxx_std_11)
  set_target_properties(${PROJECT_NAME} PROPERTIES POSITION_INDEPENDENT_CODE  ON)
  set_target_properties(${PROJECT_NAME} PROPERTIES CUDA_SEPARABLE_COMPILATION ON)    
  target_link_libraries(${PROJECT_NAME} ${CUDA_LIBRARIES} ${CUDA_cudadevrt_LIBRARY})
  ########## inc/libs/exe/features Ends Here ######################
  </code></pre>
  
  <p>Running this code gives:</p>
  
  <p><code>unknown error in /home/ahdhn/test_cg/test_cg.cu at line 67</code> </p>
  
  <p>This is the first line that uses <code>cudaMalloc</code>. I made sure that the code is compiled for the correct architecture by querying <code>__CUDA_ARCH__</code> from the device and the results is 700. Kindly let me know if you spot me doing something wrong in the code or the <code>CMakeLists.txt</code> file. </p>