cudaGetSymbolAddress() Fail in Cuda 5

The following code snippet was working fine in cuda 3.xx and below but appears to fail for me in cuda 5.

#define NSTATES 512
device unsigned m_z[NSTATES] ;
device unsigned m_w[NSTATES];

// initialize independent seeds
extern “C” LIBSPEC void initseeds()
{
unsigned mz[NSTATES] ;
unsigned mw[NSTATES] ;
unsigned *devmz, devmw ;
int k ;
cudaError cuerr ;
for (k = 0 ; k < NSTATES ; k++) {
mz[k] = xor128() ;
mw[k] = xor128() ;
}
/
get symbol address for state */
cudaGetSymbolAddress((void **) &devmz, “m_z” );
// Error occurs here: CUDA error: invalid device symbol

}

Replacing device with constant doesn’t fix my problem either. What has changed in Cuda 5 that would cause this error?

For a number of overloaded API functions that used to accept string arguments for symbol names as one of the variants, the variants with string arguments were deprecated in CUDA 4.2 and removed in CUDA 5.0. This is presumably one of them (please check the release notes).

Try using the symbol name directly, e.g:

cudaGetSymbolAddress((void **) &devmz, mz );

Ahh thanks that did the trick. I must have missed that memo. Does the same result hold true then for cudaMemcpyToSymbol() ? Is the first argument no longer a string, but simply the device variable name?

I have a few of those lying around but that code has not yet been tested. Funny how it still compiles using a string as an input. It doesn’t even raise a warning flag.

CUDA APIs related to symbols previously were overloaded to accept either the symbol directly, or the stringified symbol name. For various reasons, functions with string arguments turned out to be a bad idea, and therefore the variants accepting string arguments were first deprecated (at minimum in 4.2, if not in 4.1) and ultimately removed in 5.0. If you change affected code to pass just the symbol, it should operate fine across all CUDA versions, as that variant has always existed.

Compile-time error checking for the now unsupported string variants of these APIs is desirable as you note. From my limited understanding, technical issues precluded compile-time detection in CUDA 5.0, so detection is via runtime error.