具有 long long int 的 CUDA atomicAdd()答案

【问题标题】：CUDA atomicAdd() with long long int具有 long long int 的 CUDA atomicAdd()
【发布时间】：2013-06-22 13:40:43
【问题描述】：

每当我尝试将atomicAdd 与(*int, int) 以外的任何东西一起使用时，我都会收到此错误：

error: no instance of overloaded function "atomicAdd" matches the argument list

但我需要使用比int 更大的数据类型。这里有什么解决方法吗？

设备查询：

/usr/local/cuda/samples/1_Utilities/deviceQuery/deviceQuery Starting...

 CUDA Device Query (Runtime API) version (CUDART static linking)

Detected 1 CUDA Capable device(s)

Device 0: "GeForce GTX 680"
  CUDA Driver Version / Runtime Version          5.0 / 5.0
  CUDA Capability Major/Minor version number:    3.0
  Total amount of global memory:                 4095 MBytes (4294246400 bytes)
  ( 8) Multiprocessors x (192) CUDA Cores/MP:    1536 CUDA Cores
  GPU Clock rate:                                1084 MHz (1.08 GHz)
  Memory Clock rate:                             3004 Mhz
  Memory Bus Width:                              256-bit
  L2 Cache Size:                                 524288 bytes
  Max Texture Dimension Size (x,y,z)             1D=(65536), 2D=(65536,65536), 3D=(4096,4096,4096)
  Max Layered Texture Size (dim) x layers        1D=(16384) x 2048, 2D=(16384,16384) x 2048
  Total amount of constant memory:               65536 bytes
  Total amount of shared memory per block:       49152 bytes
  Total number of registers available per block: 65536
  Warp size:                                     32
  Maximum number of threads per multiprocessor:  2048
  Maximum number of threads per block:           1024
  Maximum sizes of each dimension of a block:    1024 x 1024 x 64
  Maximum sizes of each dimension of a grid:     2147483647 x 65535 x 65535
  Maximum memory pitch:                          2147483647 bytes
  Texture alignment:                             512 bytes
  Concurrent copy and kernel execution:          Yes with 1 copy engine(s)
  Run time limit on kernels:                     Yes
  Integrated GPU sharing Host Memory:            No
  Support host page-locked memory mapping:       Yes
  Alignment requirement for Surfaces:            Yes
  Device has ECC support:                        Disabled
  Device supports Unified Addressing (UVA):      Yes
  Device PCI Bus ID / PCI location ID:           1 / 0
  Compute Mode:
     < Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) >

deviceQuery, CUDA Driver = CUDART, CUDA Driver Version = 5.0, CUDA Runtime Version = 5.0, NumDevs = 1, Device0 = GeForce GTX 680

【问题讨论】：

您可以使用unsigned long long int 还是必须是long long int？如果您可以使用未签名的版本，它应该可以工作。如果您必须使用签名的 64 位版本，您可以对the documentation 中给出的示例进行变体，以使用 atomicCAS 进行任意原子访问。如果您需要帮助，请做出相应的回应，我可以举个例子。

标签： cuda int add atomic long-integer

【解决方案1】：

我的猜测是错误的编译标志。您正在寻找除 int 以外的任何内容，您应该使用 sm_12 或更高版本。

正如 Robert Crovella 所说，支持 unsigned long long int 变量，但不支持 long long int。

使用的代码来自：Beginner CUDA - Simple var increment not working

#include <iostream>

using namespace std;

__global__ void inc(unsigned long long int *foo) {
  atomicAdd(foo, 1);
}

int main() {
  unsigned long long int count = 0, *cuda_count;
  cudaMalloc((void**)&cuda_count, sizeof(unsigned long long int));
  cudaMemcpy(cuda_count, &count, sizeof(unsigned long long int), cudaMemcpyHostToDevice);
  cout << "count: " << count << '\n';
  inc <<< 100, 25 >>> (cuda_count);
  cudaMemcpy(&count, cuda_count, sizeof(unsigned long long int), cudaMemcpyDeviceToHost);
  cudaFree(cuda_count);
  cout << "count: " << count << '\n';
  return 0;
}

从 Linux 编译：nvcc -gencode arch=compute_12,code=sm_12 -o add add.cu

结果：

count: 0
count: 2500

【讨论】：

为什么unsigned long long int 在文档中显示为atomicAdd() 的重载，而不是unsigned long int？
@Adam27X 我的猜测是在NVidia架构上，long的大小和int的大小是一样的，但是long long int更大。