Cooperative groups in CUDA

user_8386009 0
Viewed 3504

since the CUDA 9 release apparently it is possible to group different threads and blocks into the same group so you can manage them together. That`s very useful for me because I need to launch a kernel with several blocks and wait until all of them are synchronized (cudaThreadSynchronize() is not worthy for me because after the threads are synchronized I have to continue working in my kernel).

What I have thought is to include these blocks of threads into the same group and wait until all of them are synchronized, as the examples of Nvdia main page suggest.

They do something like this:

__device__ int reduce_sum(thread_group g, int *temp, int val)
{
    int lane = g.thread_rank();

// Each iteration halves the number of active threads
// Each thread adds its partial sum[i] to sum[lane+i]
for (int i = g.size() / 2; i > 0; i /= 2)
{
    temp[lane] = val;
    g.sync(); // wait for all threads to store
    if(lane<i) val += temp[lane + i];
    g.sync(); // wait for all threads to load
}

My problem is how to group these blocks into the g group. This is how I originally launched my kernel:

asap << <5, 1000 >> > (cuda_E2, cuda_A2, cuda_temp, Nb, *binM, Nspb);

Any time that I try to use thread_group the compiler says that it is undefied. I'm using the cooperative_groups.h header.

Does anyone know how to deal with this? Thanks in advance.

cuda
1 Answers

Quote from the documentation:

Cooperative Groups requires CUDA 9.0 or later. To use Cooperative Groups, include the header file:

#include <cooperative_groups.h>

and use the Cooperative Groups namespace:

using namespace cooperative_groups;

Then code containing any intra-block Cooperative Groups functionality can be compiled in the normal way using nvcc.

The namespace is what you are missing.

user_6818650
Related