cuda error message : invalid device symbol

Question

I'm seeing the message of "invalid device symbol" when I run the cuda binary. There was no error during compilation it. The message is in below.

Cuda error in file 'euler3d.cu' in line 416 : invalid device symbol.

And the related source code is as following.

CUDA_SAFE_CALL( cudaMemcpyToSymbol(ff_variable, h_ff_variable, NVAR*sizeof(float)) );

Is there any fault from source code? Actually this code is from Rodinia v2.1, cfd benchmark program. I'm using cuda version 3.1 and did compile with following options.

nvcc -Xptxas -v -O3 --gpu-architecture=compute_13 --gpu-code=compute_13 euler3d.cu -o euler3d -I$(CUDA_SDK_PATH)/common/inc  -L$(CUDA_SDK_PATH)/lib $(CUTIL_LIB)

The ff_variable related code is here.

#define NDIM 3
#define VAR_MOMENTUM  1
#define VAR_DENSITY_ENERGY (VAR_MOMENTUM+NDIM)
#define NVAR (VAR_DENSITY_ENERGY+1)
__constant__ float ff_variable[NVAR];

Rodinia v2.1 was released in mid-2012, when most folks were on CUDA 4 or CUDA 5 prerelease. Why are you using CUDA 3.1 which is pretty old? The syntax expected for cudaMemcpyToSymbol has changed over time, in particular with what can be passed as an acceptable device symbol. I think you're likely to have better luck with CUDA 4.1 or newer. If you post the declaration syntax of ff_variable from the source code, it may be possible to be more specific about what is going on. — Robert Crovella
Thank you Robert, but I checked cfd source code has not been changed from rodinia version 1.0 which was released March 2010. The cuda 3.1 released July 2010. I added the declaration of ff_variable to the text. — user1292251
I downloaded rodinia 2.1 and set up a machine with a cc1.3 device (QuadroFX 4800), RHEL 5.5, cuda 3.1 plus cuda 3.1 sdk. I was able to compile euler3d.cu successfully and run it as well (using the supplied run script). I also noticed that my line 416 in euler3d.cu appears to differ from your line 416. Have you modified the source in any way? Are you using both the cuda 3.1 toolkit and the cuda 3.1 sdk? Which driver is loaded on your machine? What GPU specifically are you using? Are you able to run other CUDA codes (such as deviceQuery from the SDK)? What is your OS and 32/64 bit? — Robert Crovella
My environment info: Ubuntu 10.04 64bit, GeForce GT 330, and I found one more commented line at the upper code, so you're right, the line I mentioned is actually 415th line. I think only graphic driver is differ from your environment. Do you think it can cause the runtime error (invalid device symbol)? BTW what is the cc1.3? — user1292251

P Marecki P Marecki · Accepted Answer · 2012-11-29T09:36:00

Code below compiler and runs OK, but reproduces your error on replacement of kk_d with "kk_d" (i.e. it compiles, but reports invalid device symbol in runtime). CUDA reference (v. 4.2) is a little misleading here, as it says first argument should indeed be a const literal.

#include <cstdio>
#include "XFC_cudaError.cuh"   //my error reporting

__device__ int kk_d;

__global__ void foo() {
  printf("%i ", kk_d);
}

int main() {
  int kk = 10;
  cudaMemcpyToSymbol(kk_d, &kk, 4);
  CUDA_CHK;
  foo<<<1,1>>>();
  CUDA_CHK;
  cudaDeviceSynchronize();
}

You should now be able to modify the code appropriately. Note, that printf will not compile with arch=sm_13; you would need at least sm_20, but this is irrelevant to your issue.

With your edit: Syntax for copying arrays is ... exactly as you have in your code. Specifically adding __constant__ float ff[2]; to global device variables, and float rr[] = {1,2}; to host code allows for a copy via cudaMemcpyToSymbol(ff, rr, 8);, which compiles, and runs OK, even on sm_13.

Maybe the error originates from an earlier code, or from your h_ff_variable?

cuda error message : invalid device symbol

1 Answers