Compiler problem with pointer init

Hello,

The following code:

device
int test [3] = {1,2,3};

device
int *test_ptr = test;

Results in:

Assertion failure at line 913 of …/…/be/cg/NVISA/cgemit_targ.cxx:

Compiler Error in file /tmp/tmpxft_0000345f_00000000-7_first.cpp3.i during Assembly phase:

NYI initv kind 1

nvopencc INTERNAL ERROR: /usr/local/cuda/open64/lib//be returned non-zero status 1

The compiler is from SDK 2.1

wojtek@srebrny:~/NVIDIA_CUDA_SDK/projects/first$ nvcc --version
nvcc: NVIDIA ® Cuda compiler driver
Copyright © 2005-2007 NVIDIA Corporation
Built on Wed_Dec__3_18:29:25_PST_2008
Cuda compilation tools, release 2.1, V0.2.1221

Unimplemented is INITVKIND_SYMOFF…

–Regards,
Wojtek N.

Try

__device__

int test [3] = {1,2,3};

__device__

int *test_ptr = 0;

cudaGetSymbolAddress((void**)&test_ptr, test);

That won’t work either… the host cannot write to test_ptr. You’ll need a combination of cudatGetSymbolAddress and cudaMemcpyToSymbol in order to truly initialize that pointer. But ask yourself if you truly want to do this… There is almost never a good reason to want to do an uncoalesced read from global mem just to get a pointer to perform yet another read from global mem.

I am just reporting a compiler defect. Please don’t read my code literally.

As for workarounds, I guess the simplest one would be to call a host function to init the pointers. But in my real code the tables are supposed to reside constant memory, I haven’t found an acceptable workaround, so suggestions welcome.

Well, I do. I need starting addresses of some variable length tables. I don’t see why not if this happens outside any performance critical loop. Alternative would be a huge case statement.

–Regards,

Wojtek N.

well, I agree with MisterAnderson42 - this is indeed not implemented in nvcc but I guess this is not a compiler bug - global memory is just not supposed to be initialized in this way.

So you can either use constant memory properly initialized or, if your tables are too large, put them to texture memory and pass starting indices as parameters to your kernel

Actually, you don’t agree MisterAnderson42. He never said that this compiler behaviour was not a compiler bug. If a failed assertion is not a compiler defect for you then I am grateful that you don’t write compilers I am used to use.

MisterAnderson42 only warned me about the adverse performance effect of such statement, and I replied that in my code it is used so marginally, that there would not be any.

Yes, daddy. I will never take constant pointer to a constant again. This is sooooo improper. :">

Seriously, I don’t find your void patronizing statements appropriate.

–Regards,

Wojtek N.

Correct. Sorry for misinterpreting your post as a question of how to initialize such a pointer and not as just a bug report.

ok, perhaps I didn’t state it clear - I only agree with him that this way of initializing device pointers is inappropriate

as for compilers: you are probably not used to them at all if you believe that they are mature pieces of code with a solid fool-protection

compilers are constantly under development - just upload nvopencc sources and grep for ‘NYI’ - you will see dozern other places

actually your attempt to initialize pointer was a violation of this fool-protection…

or perhaps you would be more satisfied to have it compiled quietly to something erroneous which would then cause a runtime fault ? I doubt it

You can’t tell a legal C statement from an illegal C statement. You can’t tell a failed assertion from a compile time warning. And you even can’t even tell upload form download. It is therefore impossible for you to offend me. By calling me a fool, or in any other way.

You claim that nvcc is full of bugs. So why don’t you disclose the list of bugs you found?

Prior to answering me, ask yourself a question if you really have something to say.

nvcc is non-compliant in some well-defined ways (no function pointers, no recursion, double-precision behavior on 1.1 devices, “fast-math” etc.).

nvcc is also non-compliant in some loosely-defined ways (C++ support is “spotty”, your mileage may vary).

nvcc also has bugs. WojtekN’s example is a case of this.

ok, cool down please )) you really take it too much personal and this forum is clearly not the right place for such discussions…

I did not call you a fool personally

I told you - compilers are constantly under development.

If you want to disclose more bugs in nvcc you can do this on your own, here is a URL:

http://svn.open64.net/listing.php?repname=…targia32_nvisa_

check it out from svn, grep for NYI and put it on the forum, however I doubt that this will change anything…

These are not critical bugs - I rebuilt nvopencc compiler from these sources and I never got assertions like yours

thank you for valuable advice