CUDA实现inter-block间同步——基于锁的块间同步

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>

#define BLOCK_NUM 5
#define THREAD_NUM 32

__device__ volatile int g_mutex;

__global__ void gpu_sync(int goalVal)
{
	int tid = blockIdx.x *blockDim.x + threadIdx.x;

	for (int loop = 0; loop < 10; loop++)
	{
		if (tid % 3 == 0)
		{
			int x = 52;
			int y = 10;
			int z = x * y;
			int k = x % y;
			int l = x ^ y;
		}
		printf("%d(loop#):%d(BLOCK#)%d(TREADZ#)\n", loop, blockIdx.x, threadIdx.x);
		if (threadIdx.x == 0)
		{
			atomicAdd((int*)&g_mutex, 1);

			// only when all blocks add 1 go g_mutex
			// will g_mutex equal to goalVal
			while (g_mutex != goalVal)
			{
				// Do nothing here
			}
		}
		__syncthreads();
		printf("Yeah!\n");
		if (threadIdx.x == 0)
		{
			atomicExch((int*)&g_mutex, 0);
		}
	}
}

__global__ void gpu_sync_nonlock(int goalVal)
{
	int tid = blockIdx.x *blockDim.x + threadIdx.x;

	for (int loop = 0; loop < 10; loop++)
	{
		if (tid%3==0)
		{
			int x = 52;
			int y = 10;
			int z = x * y;
			int k = x % y;
			int l = x ^ y;
		}
		printf("%d(loop#):%d(BLOCK#)%d(TREADZ#)\n", loop, blockIdx.x, threadIdx.x);
		__syncthreads();
		printf("Yeah!\n");
	}
}

int main()
{

	gpu_sync << <BLOCK_NUM, THREAD_NUM >> > (BLOCK_NUM);

    return 0;
}

__syncthreads只能实现intra-block间的同步：

CUDA 基于锁的同步的基本思想是使用一个全局互斥量变量来计算到达同步点的线程块的数量。如下代码所示，在 barrier 函数 __gpu_sync() 中，在一个块完成它的计算之后，它的一个线程 (这里人为设置为 0 号线程，我们称之为主导线程) 将自动地向 g_mutex 添加 1 （原子操作）。然后，主导线程将重复将 g_mutex 和一个目标值 goalVal 进行比较。如果 g_mutex 等于 goalVal，那么就意味着同步完成，即每个线程块都可以进行下一阶段的计算。在该设计中，当第一次调用 barrier 函数时，将 goalVal 设置为内核中的块数 N 。然后，当连续调用 barrier 函数时，goalVal的值每次递增 N 。这种设计比保持 goalVal 常量并在每个 barrier 之后重新设置 g_mutex 更有效，因为前者节省了指令的数量并避免了条件分支 。
参考链接