CUDA toolkit version | using constant memory

Hi

I've installed CUDA toolkit 4.1 instead of 4.0. And now matlab reports some errors depending on kernel that Im trying to create:

Error using handleKernelArgs (line 48)
The first input to parallel.gpu.CUDAKernel must be a file that contains PTX code or a string that is PTX code.

Error using handleKernelArgs (line 61)
No entries were found in the PTX file

NVIDIA is using new compiler in 4.1, and it seems to me that pct does not support it. Am I right? It is unclear for user which version of toolkit should he use. Maybe I missed something, but I could not find toolkit version requirements.

The second question is about using constant memory. Help says that user should define constant variables in CU file

__constant__ int N;
__constant__ double CONST_DATA[256];

Placing such definition in kernel does not work:

error: a "__constant__" variable declaration is not allowed inside a function body

Answer 1

[Ben Tordoff]

Hi, I will tackle the second question first as I am pretty certain of the problem. _constants_ must be defined at the global scope, i.e. not inside functions, classes etc. So, for example,

__constant__ double CONST_DATA[256];
__global__ void myKernel(double * out, double * in) {
   // Some code
}

is fine, whilst

__global__ void myKernel(double * out, double * in) {
   __constant__ double CONST_DATA[256];
   // Some code
}

is not. If you already have your constants at global scope then there must be something else amiss. In that case, perhaps you could post some short CU code that shows the problem?

Regarding your first question, I'm afraid I don't have a definite answer. I've just tried compiling a simple kernel to PTX using NVCC from CUDA 4.1 with MATLAB R2012a and it worked fine. However this may depend on the compile options specified. Perhaps you could provide some details about which MATLAB release you are using, what compile command you used, and maybe even a short snippet of code that will help me to understand the problem?

Thanks.

Ben

[m4 Chrennikov]

Ben, thanks for reply!

Second problem is solved :-)

About compilers:

Trying to create test kernel from help works fine:

/*

* Add a constant to a vector.

*/

__global__ void addToVector(float * pi, float c, int vecLen) {

int idx = blockIdx.x * blockDim.x + threadIdx.x;

if (idx < vecLen) {

pi[idx] += c;

}

}

nvcc tst.cu -ptx

>>k = parallel.gpu.CUDAKernel( 'tst.ptx', 'tst.cu' );

BUT, when you specify arch:

nvcc tst1.cu -ptx -arch=sm_20

>> k = parallel.gpu.CUDAKernel( 'tst.ptx', 'tst.cu' );

Error using handleKernelArgs (line 61)

No entries were found in the PTX file

MATLAB fails to create kernel.

PTX files with and w/o -arch=sm_20 are very different.

Please try this key for compilation.

[Ben Tordoff]

OK, you are right. The PTX that the new version of NVCC spits out is sufficiently different that MATLAB fails to parse it. In CUDA 4.1 NVCC actually switches between two different compilers according to the "arch" setting, and the newer compiler produces pretty different output. For now I believe that your only two options are to:

1) Compile using -arch=sm_13 (or below)

2) Switch back to the CUDA 4.0 SDK

I will look into whether we can send out a patch for MATLAB R2012a and will update this post if one becomes available.

Sorry about that.

Ben

[m4 Chrennikov]

Since I use atomics second option is the best.

If you know any simple way to parse new version of PTX in MATLAB please let me know. Maybe the user only need to modify new PTXes a bit :)

NV says that new compiler is about 10% fast!

[Ben Tordoff]

A fix for the "no entries" problem is now available:

http://www.mathworks.com/support/bugreports/814923

Can you try it and let me know if you still have problems?

[m4 Chrennikov]

Ben everything seems to be fine. Ill post additionally if something happens.

Thanks!

Answer 2

[Matthew]

the problem is still there .

nvcc -ptx -arch=sm_20 xxx.cu with cuda4.1 & 4.2 it run on ubuntu 11.10

matlab 2011b