为什么atomicAdd没有实现双打? [英] Why has atomicAdd not been implemented for doubles?
问题描述
为什么不将 atomicAdd()
作为CUDA 4.0或更高版本的一部分显式实现?
Why hasnt atomicAdd()
for doubles been implemented explicitly as a part of CUDA 4.0 or higher?
从 CUDA计划指南4.1 的附录F第97页
atomicAdd的以下版本已经实现。
From the appendix F Page 97 of the CUDA programming guide 4.1 the following versions of atomicAdd have been implemented.
int atomicAdd(int* address, int val);
unsigned int atomicAdd(unsigned int* address,
unsigned int val);
unsigned long long int atomicAdd(unsigned long long int* address,
unsigned long long int val);
float atomicAdd(float* address, float val)
一个小实现的atomicAdd双重如下
,我刚刚开始在我的项目中使用。
The same page goes on to give a small implementation of atomicAdd for doubles as follows which I have just started using in my project.
__device__ double atomicAdd(double* address, double val)
{
unsigned long long int* address_as_ull =
(unsigned long long int*)address;
unsigned long long int old = *address_as_ull, assumed;
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);
}
为什么不将上述代码定义为CUDA的一部分?
Why not define the above code as a part of CUDA ?
推荐答案
编辑:自CUDA 8起,双精度 atomicAdd()
在CUDA中使用SM_6X(Pascal)GPU中的硬件支持实现。
As of CUDA 8, double-precision atomicAdd()
is implemented in CUDA with hardware support in SM_6X (Pascal) GPUs.
目前,没有CUDA设备支持如你所说,它可以用 atomicAdd
c> atomicCAS
对64位整数,但是有一个非平凡的性能成本。
Currently, no CUDA devices support As you noted, it can be implemented in terms of atomicAdd
for double
in hardware.atomicCAS
on 64-bit integers, but there is a non-trivial performance cost for that.
因此,CUDA软件团队选择将正确的实现文档作为开发人员的选项记录,而不是将其作为CUDA标准库的一部分。这样,开发人员不会在不知不觉中选择他们不明白的性能成本。
Therefore, the CUDA software team chose to document a correct implementation as an option for developers, rather than make it part of the CUDA standard library. This way developers are not unknowingly opting in to a performance cost they don't understand.
Aside:我不认为这个问题应该被封闭为不建设性 。我认为这是一个完全有效的问题,+1。
Aside: I don't think this question should be closed as "not constructive". I think it's a perfectly valid question, +1.
这篇关于为什么atomicAdd没有实现双打?的文章就介绍到这了,希望我们推荐的答案对大家有所帮助,也希望大家多多支持IT屋!