Struct as template argument problem

Accelerated Computing CUDA CUDA Programming and Performance
1

Look at this code, please:

File test.cu:

[codebox]

struct S{};

template global void template_test() {

float4 f = make_float4( 0, 0, 0, 0 );

//int a;

}

void test() {

dim3 g( 1, 1, 1 );

dim3 b(1, 1, 1 );

template_test<S><<<g,b>>>();

}

[/codebox]

It compiles without errors.

Now uncomment //int a and comment float4 :

[codebox]

struct S{};

template global void template_test() {

//float4 f = make_float4( 0, 0, 0, 0 );

int a;

}

void test() {

dim3 g( 1, 1, 1 );

dim3 b(1, 1, 1 );

template_test<S><<<g,b>>>();

}

[/codebox]

Nvcc output:

[codebox]

1>Executing tool NVCC_CUDA on file d:\SciArtSVN\Sciart\PluginSystem\GpuBoard\test.cu, result Tmp\Debug\test.obj

1>test.cu

1>d:/SciArtSVN/Sciart/PluginSystem/GpuBoard/test.cu(5): warning: variable “a” was declared but never referenced

1> detected during instantiation of “void template_test() [with T=S]”

1>(11): here

1>tmpxft_00000d30_00000000-3_test.cudafe1.gpu

1>tmpxft_00000d30_00000000-8_test.cudafe2.gpu

1>d:/SciArtSVN/Sciart/PluginSystem/GpuBoard/test.cu(5): warning: variable “a” was declared but never referenced

1>tmpxft_00000d30_00000000-3_test.cudafe1.cpp

1>tmpxft_00000d30_00000000-13_test.ii

1>d:/SciArtSVN/Sciart/PluginSystem/GpuBoard/test.cu(11) : error C2065: ‘template_test__entry’ : undeclared identifier

1>d:/SciArtSVN/Sciart/PluginSystem/GpuBoard/test.cu(11) : error C2275: ‘S’ : illegal use of this type as an expression

1> d:/SciArtSVN/Sciart/PluginSystem/GpuBoard/test.cu(1) : see declaration of ‘S’

1>d:/SciArtSVN/Sciart/PluginSystem/GpuBoard/test.cu(11) : error C2059: syntax error : ‘)’

1>C:\DOCUME~1\Sieciech\LOCALS~1\Temp/tmpxft_00000d30_00000000-3_test.cudafe1.stub.c(10) : error C2275: ‘S’ : illegal use of this type as an expression

1> d:/SciArtSVN/Sciart/PluginSystem/GpuBoard/test.cu(1) : see declaration of ‘S’

1>C:\DOCUME~1\Sieciech\LOCALS~1\Temp/tmpxft_00000d30_00000000-3_test.cudafe1.stub.c(10) : error C2059: syntax error : ‘)’

1>C:\DOCUME~1\Sieciech\LOCALS~1\Temp/tmpxft_00000d30_00000000-3_test.cudafe1.stub.c(10) : error C2275: ‘S’ : illegal use of this type as an expression

1> d:/SciArtSVN/Sciart/PluginSystem/GpuBoard/test.cu(1) : see declaration of ‘S’

1>C:\DOCUME~1\Sieciech\LOCALS~1\Temp/tmpxft_00000d30_00000000-3_test.cudafe1.stub.c(10) : error C2059: syntax error : ‘)’

1>C:\DOCUME~1\Sieciech\LOCALS~1\Temp/tmpxft_00000d30_00000000-3_test.cudafe1.stub.c(23) : error C2143: syntax error : missing ‘;’ before ‘<’

1>C:\DOCUME~1\Sieciech\LOCALS~1\Temp/tmpxft_00000d30_00000000-3_test.cudafe1.stub.c(23) : error C2182: ‘__wrapper__device_stub_template_test’ : illegal use of type ‘void’

1>C:\DOCUME~1\Sieciech\LOCALS~1\Temp/tmpxft_00000d30_00000000-3_test.cudafe1.stub.c(23) : error C2988: unrecognizable template declaration/definition

1>C:\DOCUME~1\Sieciech\LOCALS~1\Temp/tmpxft_00000d30_00000000-3_test.cudafe1.stub.c(23) : error C2059: syntax error : ‘<’

1>C:\DOCUME~1\Sieciech\LOCALS~1\Temp/tmpxft_00000d30_00000000-3_test.cudafe1.stub.c(25) : error C2143: syntax error : missing ‘;’ before ‘{’

1>C:\DOCUME~1\Sieciech\LOCALS~1\Temp/tmpxft_00000d30_00000000-3_test.cudafe1.stub.c(25) : error C2447: ‘{’ : missing function header (old-style formal list?)

[/codebox]

What the hell is going on??

I coudn’t use Struct as template arguments in my kernels,

any help will be appreciated

2

After renaming test() to main() and it compiles fine on my system:

[codebox]$ cat test.cu

include <cuda.h>

struct S{};

template

global void template_test()

{

//float4 f = make_float4( 0, 0, 0, 0 );

int a;

}

int main(void)

{

dim3 g( 1, 1, 1 );

dim3 b(1, 1, 1 );

template_test<S><<<g,b>>>();

}

$ nvcc test.cu

test.cu(9): warning: variable “a” was declared but never referenced

test.cu(9): warning: variable “a” was declared but never referenced

$ nvcc --version

nvcc: NVIDIA ® Cuda compiler driver

Copyright © 2005-2009 NVIDIA Corporation

Built on Thu_Apr__9_05:05:52_PDT_2009

Cuda compilation tools, release 2.2, V0.2.1221[/codebox]

CUDA 2.2 has much better template support than 2.1, so make sure that you have it installed (w/ nvcc --version)

3

Thanks for the reply.

the same code on my computer generates following error:

[codebox]

test.cu

test.cu(9): warning: variable “a” was declared but never referenced

      detected during instantiation of "void template_test<T>() [with T=S]"

(16): here

tmpxft_00000ef4_00000000-3_test.cudafe1.gpu

tmpxft_00000ef4_00000000-8_test.cudafe2.gpu

test.cu(9): warning: variable “a” was declared but never referenced

tmpxft_00000ef4_00000000-3_test.cudafe1.cpp

tmpxft_00000ef4_00000000-13_test.ii

test.cu(16) : error C2065: ‘template_test__entry’ : undeclared identifier

test.cu(16) : error C2275: ‘S’ : illegal use of this type as an expression

    test.cu(3) : see declaration of 'S'

test.cu(16) : error C2059: syntax error : ‘)’

C:\DOCUME~1\Sieciech\LOCALS~1\Temp/tmpxft_00000ef4_00000000-3_test.cudafe1.stub.

c(10) : error C2275: ‘S’ : illegal use of this type as an expression

    test.cu(3) : see declaration of 'S'

C:\DOCUME~1\Sieciech\LOCALS~1\Temp/tmpxft_00000ef4_00000000-3_test.cudafe1.stub.

c(10) : error C2059: syntax error : ‘)’

C:\DOCUME~1\Sieciech\LOCALS~1\Temp/tmpxft_00000ef4_00000000-3_test.cudafe1.stub.

c(10) : error C2275: ‘S’ : illegal use of this type as an expression

    test.cu(3) : see declaration of 'S'

C:\DOCUME~1\Sieciech\LOCALS~1\Temp/tmpxft_00000ef4_00000000-3_test.cudafe1.stub.

c(10) : error C2059: syntax error : ‘)’

C:\DOCUME~1\Sieciech\LOCALS~1\Temp/tmpxft_00000ef4_00000000-3_test.cudafe1.stub.

c(23) : error C2143: syntax error : missing ‘;’ before ‘<’

C:\DOCUME~1\Sieciech\LOCALS~1\Temp/tmpxft_00000ef4_00000000-3_test.cudafe1.stub.

c(23) : error C2182: ‘__wrapper__device_stub_template_test’ : illegal use of typ

e ‘void’

C:\DOCUME~1\Sieciech\LOCALS~1\Temp/tmpxft_00000ef4_00000000-3_test.cudafe1.stub.

c(23) : error C2988: unrecognizable template declaration/definition

C:\DOCUME~1\Sieciech\LOCALS~1\Temp/tmpxft_00000ef4_00000000-3_test.cudafe1.stub.

c(23) : error C2059: syntax error : ‘<’

C:\DOCUME~1\Sieciech\LOCALS~1\Temp/tmpxft_00000ef4_00000000-3_test.cudafe1.stub.

c(25) : error C2143: syntax error : missing ‘;’ before ‘{’

C:\DOCUME~1\Sieciech\LOCALS~1\Temp/tmpxft_00000ef4_00000000-3_test.cudafe1.stub.

c(25) : error C2447: ‘{’ : missing function header (old-style formal list?)

[/codebox]

[indent]

–version

nvcc: NVIDIA ® Cuda compiler driver

Copyright © 2005-2009 NVIDIA Corporation

Built on Sat_May__2_13:16:30_PDT_2009

Cuda compilation tools, release 2.2, V0.2.1221[/indent]

version is the same as yours. I don’t understand this error messages.

4

I’ve been using template kernels with CUDA 2.1 and they were working fine. Now with version 2.2 I have these errors. I made another test script

template <int a> 

__global__ void kernelA(int * r)//error

{

	*r = a;

	int c;

};

template <int a> 

__global__ void kernelB(int * r)//OK

{

	*r = 0;

};

template <int a> 

__global__ void kernelC(int * r)//error

{

	int c;

	*r = 0;

};

template <class A> 

__global__ void kernelD(int * r)//OK

{

	A a=0;

	*r = 0;

};

void f(){

	kernelA<10><<<1,1>>>(0);//error C2065: kernelA__entry undeclared identifier

	kernelB<10><<<1,1>>>(0);//OK

	kernelC<10><<<1,1>>>(0);//error C2065: kernelC__entry undeclared identifier

	kernelD<int><<<1,1>>>(0);//OK

};

It is WIN32 Cuda compilation tools, release 2.2, V0.2.1221.

I found that compiling with switch

–host-compilation C

actually makes it work, but some C++ features like exceptions will be disabled.

5

Thank you very much. It helps.

Maybe nVidia fix this bug in new release

6

So the problem occurs when there is a warning reported in a template kernel.
After fixing warnings I was able to compile my kernels with --host-compilation C++

(To launch my kernels I need all the data structures which means lots of C++, and I can’t separate the code because textures can only be global file-scope).

7

Thanks for reporting this bug.
It will be fixed in the 2.3 release.