CODEWITHSUNDEEP
c++cuda
shrimp
shrimp

Reputation: 1

invalid device symbol in cudaGetSymbolAddress(&cuset_addr, random_states_g)

#define cuset(symbol, T, val)
{
   void *cuset_addr;
   cucheck(cudaGetSymbolAddress(&cuset_addr, symbol));
   T cuset_val=(val);
   cucheck(cudaMemcpy(cuset_addr, &cuset_val, sizeof(cuset_val),cudaMemcpyHostToDevice))
}

when I compiler and run the project,it present a error:

invalid device symbol in cudaGetSymbolAddress(&cuset_addr, random_states_g).

I download the project from:https://github.com/canonizer/halloc

The readme in the project present this:

Note: libraries and tests are currently not compiled for compute_50/sm_50, i.e. Maxwell.

My environment is:Ubuntu14.04, cuda7.5, capability5.0。

The compile as this: 

nvcc -gencode arch=compute_20,code=sm_20 -gencode arch=compute_30,code=sm_30 -gencode arch=compute_35,code=sm_35 -lineinfo -O3 -lib -rdc=true -Xptxas -dlcm=cg -Xptxas -dscm=wb \
    -Xptxas -maxrregcount=64 -o bin/libhalloc.a src/*.cu

I don't know whether it is due to my cuda version, when I search the error, I found it seems there are some symbol that were remove from 5.0.

Upvotes: 0

Views: 486

Answers (1)

Robert Crovella
Robert Crovella

Reputation: 152259

In the makefile, change this:

ARCH= -gencode arch=compute_20,code=sm_20 \
    -gencode arch=compute_30,code=sm_30 \
    -gencode arch=compute_35,code=sm_35

to this:

ARCH= -gencode arch=compute_20,code=sm_20 \
    -gencode arch=compute_30,code=sm_30 \
    -gencode arch=compute_35,code=compute_35 \
    -gencode arch=compute_35,code=sm_35

and rebuild the project.

The underlying problem here is that the project compile options are set up to include device code (SASS) only, i.e. no options are included to generate PTX. With only device code for sm_20, sm_30 and sm_35, there is no option to produce device code for your sm_50 device, and therefore there is no loadable module for your device.

Normally, this sort of error might show up as an "invalid device function" error, when you try to run a kernel. But if the first activity is accessing device symbols, those symbols are also not valid because there is no proper image loaded on the device, so you will get this particular error in that case.

By including a compile option that ends in code=compute_35, we instruct the CUDA compiler to also include a PTX module for your project, and PTX can be forward-JIT-compiled at runtime to match whatever newer device you are running on.

This is not the only possible way to modify the compile operation to work with your device, but it is one possible approach that will generally be forward-compatible with a range of devices.

Upvotes: 2

Related Questions