Is that possible to disable threads in a block with cooperative groups?

Accelerated Computing CUDA CUDA Programming and Performance
1

Hello everyone,

Is that possible to disable some of the threads within a thread group using cooperative groups?

For instance, I have following code.

__global__ kernel() {
	__shared__ res;
	if(threadIdx.x == 0)
		res = a[threadIdx.x]++;
	__syncthreads();
	
	// all threads use res 
}

Can I implement same thing like that?

__global__ kernel() {
	disableThreads();
	res = a[threadIdx.x]++;
	enableThreads();

	// all threads use res 
}
2

I cannot see the difference. If you don’t want to synchronize threads for just one working thread, it may be a good choice to use constant memory (or kernel parameter, probably easier). As long as the results of some threads will be used by other threads, synchronization is mandatory to avoid data racing. To guarantee right result, the function of disableThreads()/enableThreads() must have this kind of functionality as well. The only way to get rid of this is to make the result independent.