CUDA无效的设备符号错误

6
以下代码可以编译通过,但当我尝试运行它时,出现了问题。
GPUassert: invalid device symbol file.cu 114

当我注释掉标有“!!!”的行时,错误就不会显示出来。我的问题是什么导致了这个错误,因为它对我毫无意义。
使用nvcc编译文件.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));
}

这可能与错误无关,但是您在这行中的意思是什么:if(0<=hash_num<HASH_ROW) { - Engin Kayraklioglu
谢谢您的评论,那是一个错误,但它并没有解决问题。 - user3390078
我无法重现这个错误。它在我的电脑上编译和运行时没有出现错误。CUDA 版本是什么?操作系统是什么?使用了哪款 GPU? - Robert Crovella
设备0:"GeForce 9800 GT" CUDA驱动程序版本/运行时版本 5.5 / 5.5 Kubuntu 12.04 - user3390078
2个回答

9

一个相当不寻常的错误。

故障发生原因如下:

  • 只指定了虚拟架构(-arch compute_11),导致PTX编译步骤被推迟到运行时(即强制JIT编译)
  • JIT编译在运行时失败了
  • JIT编译(和链接)的失败意味着设备符号无法正常建立
  • 由于设备符号出现问题,对设备符号dev_base执行cudaGetSymbolAddress操作失败,并抛出错误。

为什么JIT编译会失败?您可以通过指定-arch=sm_11而不是-arch compute_11来触发机器代码编译(运行ptxas汇编器)来自行查找答案。如果您这样做,您将得到以下结果:

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

即使您的代码没有调用find内核,它也必须成功编译才能为符号提供正常的设备环境。
为什么会出现这个编译错误?因为您正在请求每个线程太多的本地内存。 cc 1.x设备每个线程的本地内存限制为16KB,而您的find内核请求的要比那多得多(超过64KB)。
当我最初在我的设备上尝试时,我使用的是cc2.0设备,其限制更高(每个线程512KB),因此JIT编译步骤成功了。
一般来说,我建议同时指定虚拟架构和机器架构,简写方式如下:
nvcc -arch=sm_11 ....

(针对cc1.1设备)

这个问题/答案也可能会引起您的兴趣,nvcc手册有关于虚拟和机器架构的更多细节,以及如何为每个阶段指定编译过程。

我认为当您在内核中注释掉那些特定行时,错误消失的原因是,当这些被注释掉时,编译器能够优化访问那些本地内存区域,并优化本地内存的实例化。这使得JIT编译步骤能够成功完成,您的代码可以“无运行时错误”运行。

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


2

这个错误通常意味着内核已经为错误的架构编译。你需要找出你的GPU的计算能力,然后为该架构进行编译。例如,如果你的GPU具有计算能力1.1,则使用-arch=sm_11进行编译。你也可以为多个架构构建可执行文件。


网页内容由stack overflow 提供, 点击上面的
可以查看英文原文,
原文链接