__constant__ stopped working

Accelerated Computing CUDA CUDA Programming and Performance
1

I have the following stuff in the .cu file:

device constant int target_hash[4];//also tried without device
device constant char charset_c[128];

When I try to do any of these:

cudaMemcpyToSymbol(“target_hash”, hash_i, 44, 0, cudaMemcpyHostToDevice);//Also tried without cudaMemcpyHostToDevice
cudaMemcpyToSymbol(target_hash, hash_i, 44, 0, cudaMemcpyHostToDevice);
cudaMemcpyToSymbol(“charset_c”, charset, charset_len+1);
cudaMemcpyToSymbol(charset_c, charset, charset_len+1);//charset_len<256

I am getting “invalid device symbol”.

That is pretty frustrating, because this code was perfectly working in pre-2.0 CUDA SDK.
My old binaries are still fine, but now I am getting this. Not sure what to do with it :-(

Any clues?

2

Still not solved :-(
Tried both Intel C++ and MSVS 2008, same behavior.
I am not even calling CUDA kernel, it’s failing on the initialization stage :-(

3

Here’s a simple code which I wrote to test this:

__device__ __constant__ char datad[128];

int32 main(int32 argc, char** argv) {

  char data[128];

  cudaMemcpyToSymbol(datad, data, 128, 0, cudaMemcpyHostToDevice);

  cudaError_t check123 = cudaGetLastError();

  if(check123 != cudaSuccess) {

	  std::cerr << "Cuda-runtime failure! " << cudaGetErrorString(check123) << std::endl;

	  exit(1);

	}

  }

  return 0;

}

I didn’t face any issue. The code ran successfully in my case… :mellow:

PS:

I’m using MSVC 2005 express edition + CUDA ver 2.3 + quadro NVS 295

4

When I removed constant and used cudaMemcpy, I’ve got “invalid device pointer” :-S
So it appears that first cuda-related operation is failing.

5

Post a little more code, like a self-contained test case.

then we can see if we can reproduce this problem.

Christian

6

This works for me too.

The problem is that I need to initialize a constant from .cpp file, while it is defined in .cu file. This causes problems.

Could you try to move device constant char datad[128]; to .cu file?

I tried both:

  1. Call cudaMemcpyToSymbol from cpp, with “datad” - it says “invalid device symbol”

  2. Call a host function in .cu file, and let it initialize datad, same “invalid device symbol”.

Looks like the issue is with separated cpp and cu files.

7

create an API function inside the .cu file that you call from the .cpp file. You pass the data pointer to the API call and

this in turn copies it into constant memory.

8

This is exactly what I do in 2nd variant:

__constant__ int target_hash[4];// = {0xb182b498, 0xf4d2ac41, 0x1f636569, 0xaf4caf00};

extern "C" void init_md5_cuda(gpu_data_xyz* data, int* hash_i, unsigned char *charset, int charset_len)

{

	int test[4];

	cudaMemcpyToSymbol(target_hash, test/*hash_i*/, 4*4, 0, cudaMemcpyHostToDevice);

//	cudaMemcpy(target_hash, hash_i, sizeof(int)*4, cudaMemcpyHostToDevice);

	CCU2();//this checks for lasterror

}
9

I was able to trigger “invalid device symbol” in single .cu program by specifying non-existing constant name.
So, that means that for some reason, CUDA does not see my constant target_hash.

10

Here goes minimal code to reproduce the problem:

md5.cpp:

#include "md5_kernel.h"

int main(int argc, char **argv)

{	

		init_md5_cuda(0/*&g->gpu[0].gpu*/, 0, 0, 0);

}

md5_kernel.h

extern "C" {

	void init_md5_cuda(void* data, int* hash_i, unsigned char *charset, int charset_len);

}

md5_kernel.cu

#include <stdio.h>

#include "md5_kernel.h"

__constant__ int targethash[4];// = {0xb182b498, 0xf4d2ac41, 0x1f636569, 0xaf4caf00};

void init_md5_cuda(void* data, int* hash_i, unsigned char *charset, int charset_len)

{

	int test[4];

	cudaMemcpyToSymbol("targethash", test/*hash_i*/, 1, 0, cudaMemcpyHostToDevice);

	cudaError_t result = cudaGetLastError();

	if(result != cudaSuccess)

	{

		char msg[512];msg[0]=0;

		strcpy_s(msg, 512, "CUDA kernel error: ");

		strcat_s(msg, 512, cudaGetErrorString(result));

		printf(msg);

		return;

	}

}

That’s it. I cannot tell how frustrating to spend 16 hours to make working code work again :-S

11

Both cubin & ptx have targethash inside.
Works in emulation.
:-S

12

Any ideas?

13

Sorry for the delayed response… Can you declare an ‘extern’ in the header file and try?
extern constant int targethash[4];

14

1>md5_kernel.cu(4) : error C2086: ‘int targethash__cuda_shadow_variable__[4]’ : redefinition

1> md5_kernel.h(1) : see declaration of ‘targethash__cuda_shadow_variable__’

15

Ah… sorry… my bad… that was not the correct way of declaring the global constant variables (silly me) :(
Please try the following instead…

In the .h file:
extern constant int targethash[4];

In the .cu file:
int targethash[4];

16

Same result:

1>Compiling...

1>md5_kernel.cu

1>md5_kernel.cudafe1.gpu

1>md5_kernel.cudafe2.gpu

1>md5_kernel.cudafe1.cpp

1>md5_kernel.cu.cpp

1>c:/my/work/BarsWF/algo/md5/md5_kernel.cu(4) : error C2086: 'int targethash__cuda_shadow_variable__[4]' : redefinition

1>		c:\my\work\barswf\algo\md5\md5_kernel.h(1) : see declaration of 'targethash__cuda_shadow_variable__'
17

Your point about this being particularly an issue with the separated cu and cpp files, seems to be right.
Anyone from nvidia?

18

Summary: Unable to set constant when working with separate .cu and .cpp files. When using single .cu program everything is fine. In emulation everything is fine.

Sourcecode to reproduce an issue at [url=“The Official NVIDIA Forums | NVIDIA”]The Official NVIDIA Forums | NVIDIA

19

It appeared CUDA Wizard thing was the cause. When I switched to SDK .rules file, everything started to work fine automagically :-D

20

Now all we need to do is find out what exactly the Wizard does wrong. After all a lot of people are using it.