Compiler Bug - Protection Classes There is a potential bug regarding inline functions utilizing memb

I’ve been doing some work with structs that contain private or protected variables. They work fine when generated for the device, but when I’m running in device-emulation mode, the compiler seems confused by the use of private or protected variables within member functions. It looks like when the compiler is in device-emulation mode that the function bodies are inlined, ignoring the scope of the original body. The problem seems to occur on both OS X and Linux, but with a slightly different error message. Here’s a sample case:

struct type

{

private:

		int a;

public:

		__device__ void set(int a_)

		{ a = a_; }

};

__global__ void kernel(type arg)

{

		arg.set(3);

}

int main()

{ }

Linux:


When I compile this on 64-bit Linux (Ubuntu 9.04.1, CUDA 2.1), I get an error saying that the access of a is disallowed since a is private:

$ nvcc -deviceemu test.cu -o test

In file included from /tmp/tmpxft_00003cbe_00000000-1_test.cudafe1.stub.c:5,

				 from test.cu:20:

test.cu: In function ‘void _Z6kernel4type(type)’:

test.cu:5: error: ‘int type::a’ is private

test.cu:13: error: within this context

make: *** [debug] Error 255

But when I compile this without device emulation, it works fine.

$ nvcc test.cu -o test

test.cu(11): warning: parameter "arg" was set but never used
$ nvcc --version

nvcc: NVIDIA (R) Cuda compiler driver

Copyright (c) 2005-2007 NVIDIA Corporation

Built on Wed_Dec__3_16:25:17_PST_2008

Cuda compilation tools, release 2.1, V0.2.1221
$ gcc --version

gcc (GCC) 4.2.4 (Ubuntu 4.2.4-1ubuntu3)

Copyright (C) 2007 Free Software Foundation, Inc.

This is free software; see the source for copying conditions.  There is NO

warranty; not even for MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.

OS X:


I tried it on OS X (10.5, where I have Cuda 2.0), but got a slightly different error with the same code. This is probably because of the cuda version difference (2.1 on Linux, 2.0 on OS X):

$ nvcc -deviceemu test.cu -o test

test.cu: In function 'void __globfunc__Z6kernel4type(type)':

test.cu:5: error: 'int type::a' is private

test.cu:15: error: within this context

make: *** [debug] Error 255
$ nvcc --version

nvcc: NVIDIA (R) Cuda compiler driver

Copyright (c) 2005-2007 NVIDIA Corporation

Built on Wed_Jul_16_12:14:50_PDT_2008

Cuda compilation tools, release 2.0, V0.2.1221