如何在 CUDA 上对 struct 应用原子操作?
Posted
技术标签:
【中文标题】如何在 CUDA 上对 struct 应用原子操作?【英文标题】:How do I apply atomic operation for struct on CUDA? 【发布时间】:2020-10-16 22:13:07 【问题描述】:让struct定义如下:
typedef struct S
float x;
float y;
T;
而操作struct_add
定义如下:
__device__ T struct_add(T a1, T a2)
T result;
result.x = a1.x + a2.x;
result.y = a1.y + a2.y;
如果我想以原子方式应用struct_add
,我该如何在 CUDA 中实现呢?例如a
、b
、c
需要用struct_add
求和,结果需要存放在d
中。 (其中a
、b
、c
和d
的类型为T)
听说不推荐通过 while 循环进行“锁定和访问控制”。有什么合适的方法来实现吗?
【问题讨论】:
【参考方案1】:CUDA 没有提供涵盖任意结构原子更新的通用原子方法。一些可能性:
因为您特别想更新两个相邻的 32 位项目,您可以使用通用的 64 位原子操作,它是here 描述的变体。
另一种选择是您已经提到的,基本上是实现critical section。
最后,另一种可能的方法可能是parallel reduction,尽管这并不完全类似于原子使用
根据上面的建议 1,这里是对代码 from this answer 的修改,它可能表明您可以如何使用 64 位原子:
$ cat t56.cu
#include <stdio.h>
#define DSIZE 512
#define nTPB 256
#define cudaCheckErrors(msg) \
do \
cudaError_t __err = cudaGetLastError(); \
if (__err != cudaSuccess) \
fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \
msg, cudaGetErrorString(__err), \
__FILE__, __LINE__); \
fprintf(stderr, "*** FAILED - ABORTING\n"); \
exit(1); \
\
while (0)
typedef union
float floats[2];
unsigned long long int ulong; // for atomic update
my_atomics;
__device__ my_atomics test;
__device__ unsigned long long int my_atomicAdd_2floats(unsigned long long int* address, float val0, float val1)
my_atomics loctest;
unsigned long long old = *address;
do
loctest.ulong = old;
my_atomics loc;
loc.floats[0] = val0 + loctest.floats[0];
loc.floats[1] = val1 + loctest.floats[1];
old = atomicCAS(address, loctest.ulong, loc.ulong);
while (old != loctest.ulong);
return old;
__global__ void min_test(const float* data)
int idx = (blockDim.x * blockIdx.x) + threadIdx.x;
if (idx < DSIZE)
my_atomicAdd_2floats(&(test.ulong), data[idx], (float)idx);
int main()
float *d_data, *h_data;
my_atomics my_init;
my_init.floats[0] = 0.0f;
my_init.floats[1] = 0.0f;
h_data = (float *)malloc(DSIZE * sizeof(float));
if (h_data == 0) printf("malloc fail\n"); return 1;
cudaMalloc((void **)&d_data, DSIZE * sizeof(float));
cudaCheckErrors("cm1 fail");
for (int i = 0; i < DSIZE; i++) h_data[i] = 1.0f;
cudaMemcpy(d_data, h_data, DSIZE*sizeof(float), cudaMemcpyHostToDevice);
cudaCheckErrors("cmcp1 fail");
cudaMemcpyToSymbol(test, &(my_init.ulong), sizeof(unsigned long long int));
cudaCheckErrors("cmcp2 fail");
min_test<<<(DSIZE+nTPB-1)/nTPB, nTPB>>>(d_data);
cudaDeviceSynchronize();
cudaCheckErrors("kernel fail");
cudaMemcpyFromSymbol(&(my_init.ulong), test, sizeof(unsigned long long int));
cudaCheckErrors("cmcp3 fail");
printf("device float0 result = %f\n", my_init.floats[0]);
printf("device float1 result = %f\n", my_init.floats[1]);
float host_val0 = 0.0f;
float host_val1 = 0.0f;
for (int i=0; i<DSIZE; i++)
host_val0 += h_data[i];
host_val1 += (float)(i);
printf("host float0 result = %f\n", host_val0);
printf("host float1 result = %f\n", host_val1);
return 0;
$ nvcc -arch=sm_35 -o t56 t56.cu -Wno-deprecated-gpu-targets
$ cuda-memcheck ./t56
========= CUDA-MEMCHECK
device float0 result = 512.000000
device float1 result = 130816.000000
host float0 result = 512.000000
host float1 result = 130816.000000
========= ERROR SUMMARY: 0 errors
$
我不保证上述代码没有缺陷。我建议在使用前仔细测试。
【讨论】:
如果结构中有浮点数[4],如何扩展? 不能,因为 4 个浮点数总共是 128 位,而 CUDA 硬件原子机制目前停止在 64 位。对于较大的结构更新,其他 2 个建议之一可能是您想要考虑的:并行缩减方案(可能是最好的)或关键部分。或者重新设计您的算法,使其不需要 4 次同时原子更新。 如果你在做atomicAdd
,我不清楚为什么你需要一个“耦合”的原子操作,比如这里讨论的。在我看来,对于您的 floats[4]
案例,您可以执行 4 次单独的 float
atomicAdd
操作。以上是关于如何在 CUDA 上对 struct 应用原子操作?的主要内容,如果未能解决你的问题,请参考以下文章