CUDA 无效设备符号错误

Question

下面的代码编译得很好。但是当我尝试运行它时，我得到了

GPUassert: invalid device symbol file.cu 114

当我注释由 (!!!) 标记的行时，错误不会出现。我的问题是是什么导致了这个错误，因为它让我没有任何意义。

使用 nvcc file.cu -arch compute_11 编译

#include "stdio.h"
#include <algorithm>
#include <ctime>

#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
#define THREADS 64
#define BLOCKS 256
#define _dif (((1ll<<32)-121)/(THREADS*BLOCKS)+1)

#define HASH_SIZE 1024
#define ROUNDS 16
#define HASH_ROW (HASH_SIZE/ROUNDS)+(HASH_SIZE%ROUNDS==0?0:1)
#define HASH_COL 1000000000/HASH_SIZE


typedef unsigned long long ull;

inline void gpuAssert(cudaError_t code, char *file, int line, bool abort=true)
{
  if (code != cudaSuccess) 
  {
  //fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
  printf("GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
  if (abort) exit(code);
  }
}

__device__ unsigned int primes[1024]; 
//__device__ unsigned char primes[(1<<28)+1];
__device__ long long n = 1ll<<32; 
__device__ ull dev_base;
__device__ unsigned int dev_hash; 
__device__ unsigned int dev_index; 

time_t curtime;

__device__ int hashh(long long x) {
  return (x>>1)%1024;
}
// compute (x^e)%n
__device__ ull mulmod(ull x,ull e,ull n) {
ull ans = 1;
while(e>0) {
    if(e&1) ans = (ans*x)%n;
    x = (x*x)%n;
    e>>=1;
}
return ans;
}

// determine whether n is strong probable prime base a or not.
// n is ODD
__device__ int is_SPRP(ull a,ull n) {
  int d=0;
  ull t = n-1;
  while(t%2==0) {
      ++d;
      t>>=1;
  }
  ull x = mulmod(a,t,n);
  if(x==1) return 1; 
  for(int i=0;i<d;++i) {
      if(x==n-1) return 1;
      x=(x*x)%n;
  }
  return 0;
}


__device__ int prime(long long x) {
//unsigned long long b = 2;
//return is_SPRP(b,(unsigned long long)x);
return is_SPRP((unsigned long long)primes[(((long long)0xAFF7B4*x)>>7)%1024],(unsigned long long)x);
}

__global__ void find(unsigned int *out,unsigned int *c) {

unsigned int buff[HASH_ROW][256];
int local_c[HASH_ROW];
for(int i=0;i<HASH_ROW;++i) local_c[i]=0;

long long b = 121+(threadIdx.x+blockIdx.x*blockDim.x)*_dif;
long long e = b+_dif;
if(b%2==0) ++b;
for(long long i=b;i<e && i<n;i+=2) {
    if(i%3==0 || i%5==0 || i%7==0) continue;
    int hash_num = hashh(i)-(dev_hash*(HASH_ROW));
    if(0<=hash_num && hash_num<HASH_ROW) {
    if(prime(i)) continue;
    buff[hash_num][local_c[hash_num]++]=(unsigned int)i;
    if(local_c[hash_num]==256) {
        int start = atomicAdd(c+hash_num,local_c[hash_num]);
        if(start+local_c[hash_num]>=HASH_COL) return;

        unsigned int *out_offset = out+hash_num*(HASH_COL)*4;
        for(int i=0;i<local_c[hash_num];++i) out_offset[i+start]=buff[hash_num][i]; //(!!!)
        local_c[hash_num]=0;
    }
    }
}
for(int i=0;i<HASH_ROW;++i) {
  int start = atomicAdd(c+i,local_c[i]);
  if(start+local_c[i]>=HASH_COL) return;
  unsigned int *out_offset = out+i*(HASH_COL)*4;
  for(int j=0;j<local_c[i];++j) out_offset[j+start]=buff[i][j]; //(!!!)
}

}

int main(void) {
printf("HASH_ROW: %d\nHASH_COL: %d\nPRODUCT: %d\n",(int)HASH_ROW,(int)HASH_COL,(int)(HASH_ROW)*(HASH_COL));

ull *base_adr;
gpuErrchk(cudaGetSymbolAddress((void**)&base_adr,dev_base));
gpuErrchk(cudaMemset(base_adr,0,7));
gpuErrchk(cudaMemset(base_adr,0x02,1));
}

Answer 1

一个相当不寻常的错误。

发生故障的原因是:

通过仅指定虚拟架构( -arch compute_11 )，您将 PTX 编译步骤推迟到运行时(即您强制进行 JIT 编译)

JIT 编译失败(在运行时)

JIT 编译(和链接)失败意味着无法正确建立设备符号

由于设备符号问题，操作cudaGetSymbolAddress在设备符号上 dev_base失败，并抛出错误。

为什么 JIT 编译失败？您可以通过指定 ptxas 触发机器代码编译(运行 -arch=sm_11 汇编程序)来找出自己的身份。而不是 -arch compute_11 .如果你这样做，你会得到这样的结果:

ptxas error   : Entry function '_Z4findPjS_' uses too much local data (0x10100 bytes, 0x4000 max)

所以即使你的代码没有调用 find内核，它必须成功编译才能为符号提供一个健全的设备环境。

为什么会出现这个编译错误？因为您每个线程请求太多本地内存。 cc 1.x devices are limited to 16KB local memory per thread ，以及您的 find内核要求的远不止这些(超过 64KB)。

当我最初在我的设备上尝试它时，我使用的是具有更高限制(每个线程 512KB)的 cc2.0 设备，因此 JIT 编译步骤成功了。

一般来说，我建议同时指定虚拟架构和机器架构，简写方式是:

nvcc -arch=sm_11 ....

(对于 cc1.1 设备)

此 question/answer可能也有兴趣，还有 nvcc manual有更多关于虚拟与机器架构的细节，以及如何为每个架构指定编译阶段。

我相信当您注释掉内核中的那些特定行时错误消失的原因是，注释掉那些行后，编译器能够优化对这些本地内存区域的访问，并优化对本地内存。这允许 JIT 编译步骤成功完成，并且您的代码运行“没有运行时错误”。

您可以通过将这些行注释掉然后指定完整编译( nvcc -arch=sm_11 ... )来验证这一点，其中 -arch --gpu-architecture的缩写.