CUDA 无效设备符号错误

Posted 2023-02-16

【中文标题】CUDA 无效设备符号错误

【英文标题】：CUDA invalid device symbol error

【发布时间】：2014-03-12 22:05:56

【问题描述】：

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

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));

【问题讨论】：

这可能与错误无关，但这行是什么意思：if(0

感谢您的评论，这是一个错误，但它并没有解决问题

我无法重现该错误。它编译和运行对我来说没有错误。什么CUDA版本？什么操作系统？什么 GPU？

设备 0：“GeForce 9800 GT”CUDA 驱动程序版本/运行时版本 5.5 / 5.5 Kubuntu 12.04

【参考方案1】：

一个相当不寻常的错误。

发生故障是因为：

通过仅指定虚拟架构 (-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 devices are limited to 16KB local memory per thread，而您的 find 内核请求的远不止这些（超过 64KB）。

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

一般来说，我建议同时指定一个虚拟架构和一个机器架构，这样做的简写方式是：

nvcc -arch=sm_11 ....

（对于 cc1.1 设备）

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

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

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

【参考方案2】：

这个错误通常意味着内核被编译为错误的架构。您需要find out what the compute capability of your GPU is，然后为该架构编译它。例如。如果您的 GPU 具有计算能力 1.1，请使用 -arch=sm_11 对其进行编译。您还可以为多个架构构建可执行文件。