cuda - 如何以原子方式在多个内存位置设置值?
问题描述
CUDA Programming Guide 说任何原子操作都可以使用 来实现atomicCAS()
,并给出了一个原子双加的例子:
__device__ float single(double *address,double val)
{
unsigned long long int *address_as_ull =(unsigned long long int*)address;
unsigned long long int assumed;
unsigned long long int old = *address_as_ull;
do
{
assumed = old;
old = atomicCAS(address_as_ull,assumed,__double_as_longlong(val + __longlong_as_double(assumed)));
}while(assumed !=old);
return __longlong_as_double(old);
}
现在,我面临的问题是:
我想写一个可以原子地操作两个变量地址的函数。
例如: 原子添加大约两个变量
输入
double *address_1, int *address_2
double val_1,int val_2
结果
*address_1 = *address_1+val_1;
*address_2 = *address_2+val_2;
我该如何处理这个问题?谢谢。
解决方案
一般来说,你不能这样做。硬件不支持对内存中的多个位置进行原子更改。如果两个变量都足够小以适合单个原子操作的大小,则您可以避免这种情况 - 如果您总共有超过 8 个字节,则此方法将失败。你会遇到“奶太多”的问题。
您可以做的一件事是使用某种同步协议来访问这两个值。例如,您可以使用只有一个线程可以获得的互斥锁,以安全地知道在该线程处理这些值时没有其他人正在更改这些值。请参阅:避免花很长时间来完成“牛奶过多”的场景。
当然,这在 GPU 设置中是相当昂贵的。您可能最好执行以下操作之一(通过增加好感度的顺序):
- 使用指向更大数组的指针或索引,而不是原子地更改结构,而是原子地切换指针。这解决了并发问题,但使访问速度变慢。
- 更改您的算法,以便可以分离访问并且不必原子地发生。
- 进一步更改您的算法,以避免多个线程写入单个复杂数据结构。
推荐阅读
- java - Android记录的来电和去电是静音的
- python - 将 numpy 数组的每一行保存到文件的单独行中,而不是将其分开
- react-native - 是否可以在 React Native 中实现 NFC 支付?
- pandas - 如何将多个变量数组添加到 pandas 数据框中的列?
- command - 为什么我不断收到无法运行程序错误?
- reactjs - TypeScript:仅传播多余的属性
- reactjs - 无法修复错误:找不到模块'webpack-cli/bin/config-yargs'
- html - Bootstrap 5 表单上的水平和垂直中心
- c# - 为什么我的应用程序在 300 条记录后停止工作?
- reactjs - 为什么我不能在我的项目中实现 React Helmet?