nvcc internal error

I’m getting the following error while trying to compile some code with CUDA 2.3 on OS X:

Assertion failure at line 123 of …/…/be/cg/NVISA/expand.cxx:

Compiler Error in file /tmp/tmpxft_000044d8_00000000-7_kCalculateCustomNonbondedForces.cpp3.i during Code_Expansion phase:

unexpected mtype

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

The error seems to be triggered by the following piece of code:

[codebox]float value = 0.0f;

switch (expression->op[i])

{

case CONSTANT:

    value = expression->arg[i];

    break;

case VARIABLE1:

    value = vars1.x;

    break;

case VARIABLE2:

    value = vars1.y;

    break;

… lots more cases …

}

[/codebox]

If I comment out either the VARIABLE1 or VARIABLE2 case, it compiles correctly. Having two different case statements that refer to different elements of vars1 (which is a float4) produces the error. Unfortunately, that’s essential for me. Any suggestions for workarounds?

Peter

According to my experience, such error occurs when you attempt to assign the value from uninitialized variable to a variable in global memory. I.e.

device some_type Arr[5];

global void t(){
some_type a;
Arr[2] = a;
}

may cause this error.

Unfortunately the error message is useless in this case; I’ve spent several days to understand what’s going on.

Mikle.

You were absolutely right! I was assigning the value to shared memory, not global memory, but initializing the source variable made the error go away.

Perhaps the compiler could be made to give a slightly more useful error message… :)

Peter

TOTALLY WEIRD with nvcc 2.3!

__device__ void d_func(float4& descr)

{

   descr.x = 1.0f; // write data to reference, but no read!

}

__global__ void kernel1()

{

   float4 descr;

   d_func1(descr);

}

__global__ void kernel2()

{

   float4 descr;

   d_func1(descr);

}

this code compiles with nvcc 2.2 but NOT with 2.3 giving the error mentioned here.

following code compiles (only altering kernel2, so weird):

__device__ void d_func(float4& descr)

{

   descr.x = 1.0f; // write data to reference, but no read!

}

__global__ void kernel1()

{

   float4 descr;

   d_func1(descr);

}

__global__ void kernel2()

{

   float4 descr;

   descr.y = 0.0f; // .x still remains unset

   d_func1(descr);

}

unusable!

NOT ONLY OS X specific! also linux …
… so maybe move this thread!

I just encountered the same thing on a linux system with only 2.3.
It may be fixed in 3.0 though since the same code had compiled fine on a system with the latest compiler