I have b number of blocks and each block has t number of threads. I can use
__syncthreads()
to synchronize the threads that are in a particular block. for example
__global__ void aFunction()
{
for(i=0;i<10;i++)
{
//execute something
__syncthreads();
}
}
But my problem is to synchronize all the threads in all the blocks. How can I do this?
Cooperative groups have some requirements, such as needing to launch your kernel via cudaLaunchCooperativeKernel. Which makes it not a good solution for simple projects.
An easy alternative is using atomics with bitfields, like so:
// A global var with 64 bits can track 64 blocks,
// use an array if you need to track more blocks
__device__ uint64_t CompleteMask;
//This is where we put in all the smarts
//from the CPU reference solver
__global__ void doWork() {
atomicAnd(&CompleteMask, 0);
//do lots of work
const auto SollMask = (1 << gridDim.x) - 1;
if (ThreadId() == 0) {
while ((atomicOr(&CompleteMask, 1ULL << blockIdx.x)) != SollMask) { /*do nothing*/ }
}
if (ThreadId() == 0 && 0 == blockIdx.x) {
printf("Print a single line for the entire process")
}
}
Because every block is assigned its own bit in the mask, they can never interfere. If you have more than 64 blocks, use an array to track the bits and atomicAdd to track the count like so:
// A global var with 64 bits can track 64 blocks,
// use an array if you need to track more blocks
__device__ int CompleteMask[2];
__device__ int CompleteSuperMask;
__global__ void doWork() {
for (auto i = 0; i < 2; i++) { atomicAnd(&CompleteMask[i], 0); }
atomicAnd(&CompleteSuperMask, 0);
//do lots of work
int SollMask[3];
SollMask[0] = -1;
SollMask[1] = (1 << (gridDim.x % 32)) - 1;
SollMask[2] = (1 << (gridDim.x / 32)) - 1;
const auto b = blockIdx.x / 32;
while (atomicOr(&CompleteMask[b], (1U << (blockIdx.x % 32))) != SollMask[b]) { /*do nothing*/ }
while (atomicOr(&CompleteSuperMask, (1U << b)) != SollMask[2]) { /*do nothing*/ }
if (threadIdx.x == 0 && blockIdx.x == 0) {
printf("Print a single line for the entire process");
}
}