Cooperative_groups::this_grid() is not valid on my Volta architecture GPU. How to globally synchronize

[chaosangel]

I want to synchronize among all blocks and threads in my kernel and want to use cooperate groups for this:

#include <iostream>
#include <cooperative_groups.h>

using namespace cooperative_groups;

__global__ void testKernel() {
    grid_group block = this_grid();
    printf("I can synchronize: %d\n", block.is_valid());
    block.sync();
}

int main() {
    testKernel<<<2, 1>>>();
    cudaError_t syncErr = cudaDeviceSynchronize();
    std::cout << "Sync error: " << cudaGetErrorString(syncErr) << std::endl;
    return 0;
}

Compiling it with:
nvcc -rdc=true -arch=sm_86 -rdc=true -o test test.cu
Does give the following output:

I can synchronize: 0
I can synchronize: 0
Sync error: unspecified launch failure
Reset error: no error

Why is this group not vallid and how can I synchronize all threads of all blocks.

Here are my stats from deviceQuery

CUDA Device Query (Runtime API) version (CUDART static linking)

Detected 1 CUDA Capable device(s)

Device 0: "NVIDIA RTX A500 Laptop GPU"
  CUDA Driver Version / Runtime Version          12.2 / 12.5
  CUDA Capability Major/Minor version number:    8.6
  Total amount of global memory:                 3905 MBytes (4094427136 bytes)
  (016) Multiprocessors, (128) CUDA Cores/MP:    2048 CUDA Cores
  GPU Max Clock rate:                            1537 MHz (1.54 GHz)
  Memory Clock rate:                             6001 Mhz
  Memory Bus Width:                              64-bit
  L2 Cache Size:                                 1048576 bytes
  Maximum Texture Dimension Size (x,y,z)         1D=(131072), 2D=(131072, 65536), 3D=(16384, 16384, 16384)
  Maximum Layered 1D Texture Size, (num) layers  1D=(32768), 2048 layers
  Maximum Layered 2D Texture Size, (num) layers  2D=(32768, 32768), 2048 layers
  Total amount of constant memory:               65536 bytes
  Total amount of shared memory per block:       49152 bytes
  Total shared memory per multiprocessor:        102400 bytes
  Total number of registers available per block: 65536
  Warp size:                                     32
  Maximum number of threads per multiprocessor:  1536
  Maximum number of threads per block:           1024
  Max dimension size of a thread block (x,y,z): (1024, 1024, 64)
  Max dimension size of a grid size    (x,y,z): (2147483647, 65535, 65535)
  Maximum memory pitch:                          2147483647 bytes
  Texture alignment:                             512 bytes
  Concurrent copy and kernel execution:          Yes with 2 copy engine(s)
  Run time limit on kernels:                     Yes
  Integrated GPU sharing Host Memory:            No
  Support host page-locked memory mapping:       Yes
  Alignment requirement for Surfaces:            Yes
  Device has ECC support:                        Disabled
  Device supports Unified Addressing (UVA):      Yes
  Device supports Managed Memory:                Yes
  Device supports Compute Preemption:            Yes
  Supports Cooperative Kernel Launch:            Yes
  Supports MultiDevice Co-op Kernel Launch:      Yes
  Device PCI Domain ID / Bus ID / location ID:   0 / 1 / 0
  Compute Mode:
     < Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) >

deviceQuery, CUDA Driver = CUDART, CUDA Driver Version = 12.2, CUDA Runtime Version = 12.5, NumDevs = 1

[striker159]

As described in the programming guide here 1. Introduction — CUDA C Programming Guide ,
kernels with grid synchronization cannot be launched with <<< >>> syntax but must be launched using cudaLaunchCooperativeKernel instead.

[chaosangel]

Amazing, thank you very much.