CUDA invalid device symbol error

Question

the code below compiles just fine. But when i try to run it, i got

GPUassert: invalid device symbol file.cu 114

When i comment lines marked by (!!!) the error wont show up. My question is what is causing this error because it gives me no sense.

Compiling with 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

A rather unusual error.

The failure is occurring because:

• By specifying a virtual architecture only (-arch compute_11) you defer the PTX compile step until runtime (i.e. you are forcing JIT-compile)
• The JIT-compile is failing (at runtime)
• The failure of the JIT-compile (and link) means device symbols cannot be properly established
• Due to the problem with device symbols, the operation cudaGetSymbolAddress on the device symbol dev_base fails, and throws an error.

Why is the JIT-compile failing? You can find out yourself by triggering the machine code compile (which runs the ptxas assembler) by specifying -arch=sm_11 instead of -arch compute_11. If you do that, you'll get this result:

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

So even though your code doesn't call the find kernel, it must compile successfully to have a sane device environment for symbols.

Why does this compile error occur? Because you are requesting too much local memory per thread. cc 1.x devices are limited to 16KB local memory per thread, and your find kernel is requesting quite a bit more than that (over 64KB).

When I initially tried it on my device, I was using a cc2.0 device which has a higher limit (512KB per thread) and so the JIT-compile step succeeded.

In general, I would recommend specifying both a virtual architecture and a machine architecture, and the shorthand way to do that is:

nvcc -arch=sm_11 ....

(for a cc1.1 device)

This question/answer may also be of interest, and the nvcc manual has more details about virtual vs. machine architecture, and how to specify the compilation phases for each.

I believe the reason the error goes away when you comment out those particular lines in the kernel, is that with those commented out, the compiler is able to optimize-out the accesses to those local memory areas, and optimize-out the instantiation of the local memory. This allows the JIT-compile step to complete successfully, and your code runs "without runtime error".

You can verify this by commenting those lines out and then specify a full compile (nvcc -arch=sm_11 ...), where -arch is short for --gpu-architecture.