1

我正在使用具有 1.3 计算能力和 nvcc 编译器驱动程序 4.0 的 Tesla C1060。我正在尝试对线程块进行一些本地计算。每个线程块都有一个共享数组，该数组首先初始化为零值。为了通过线程块的线程同步并发更新（添加）到共享数据，我使用 CUDA atomicAdd 原语。

一旦每个线程块准备好其共享数据数组中的结果，共享数据数组中的每个条目就会迭代地合并（使用 atomicAdd）到全局数据数组中的相应条目。

下面是一个与我基本上试图做的非常相似的代码。

#define DATA_SZ 16
typedef unsigned long long int ULLInt;

__global__ void kernel( ULLInt* data, ULLInt ThreadCount )
{
  ULLInt thid = threadIdx.x + blockIdx.x * blockDim.x;
  __shared__ ULLInt sharedData[DATA_SZ];

  // Initialize the shared data
  if( threadIdx.x == 0 )
  {
    for( int i = 0; i < DATA_SZ; i++ ) { sharedData[i] = 0; }
  }
  __syncthreads();

  //..some code here

  if( thid < ThreadCount )
  {
    //..some code here

    atomicAdd( &sharedData[getIndex(thid), thid );

    //..some code here        

    for(..a loop...)
    { 
      //..some code here

      if(thid % 2 == 0)
      {           
        // getIndex() returns a value in [0, DATA_SZ )
        atomicAdd( &sharedData[getIndex(thid)], thid * thid );
      }
    }
  }
  __syncthreads();

  if( threadIdx.x == 0 )
  {
    // ...
    for( int i = 0; i < DATA_SZ; i++ ) { atomicAdd( &Data[i], sharedData[i] ); }
    //...
  }
}

如果我使用 -arch=sm_20 进行编译，我不会收到任何错误。但是，当我使用 -arch=sm_13 选项编译内核时，出现以下错误：

ptxas /tmp/tmpxft_00004dcf_00000000-2_mycode.ptx, line error   : Global state space expected for instruction 'atom'
ptxas /tmp/tmpxft_00004dcf_00000000-2_mycode.ptx, line error   : Global state space expected for instruction 'atom'
ptxas fatal   : Ptx assembly aborted due to errors

如果我评论以下两行，我不会在 -arch=sm_13 中遇到任何错误：

atomicAdd( &sharedData[getIndex(thid), thid );
atomicAdd( &sharedData[getIndex(thid)], thid * thid );

有人可以建议我可能做错了什么吗？

4

1 回答 1

1

在 CUDA C 编程指南中找到了解决方案：在共享内存上运行的原子函数和在 64 位字上运行的原子函数仅适用于计算能力 1.2 及以上的设备。在共享内存中的 64 位字上运行的原子函数仅适用于计算能力为 2.x 及更高版本的设备。

所以基本上我不能在这里使用 ULLInt 来共享内存，不知何故我需要使用 unsigned int

于 2012-08-27T23:09:49.177 回答