Warning: Expression has no effect What does it mean?

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi everyone,

I am compiling my code with CUDA 2.2 and this error keeps popping up, although program execution is still the same (no bugs appearing):

Material.cuh(160): warning: expression has no effect

It is referring to the line: sr[zero] = originalSR;

Does anyone know what it’s saying that? There are two more warnings which seem to be when a ShadeRec is assigned another. Is it because it can’t find a copy constructor or something? How do I get rid of this warning.

[codebox]device RGBColour shade2(ShadeRec& originalSR, const RGBColour& backgroundColour)

{

float3 wo, wi, wt;	

RGBColour fr, ft; 

ShadeRec sr[16];

Ray newRays[8];

Ray currRays[8];

RGBColour srC[16];

volatile int idx = threadIdx.y*blockDim.x + threadIdx.x;

volatile int zero = 0;

volatile int one = 1;

volatile int three = 3;

sr[zero] = originalSR;[/codebox]
2

that means compiler is going to throw these calculations out, because it doesn’t affect the output anyhow.

3

As far as I can tell it makes no sense declaring idx, zero, one and three as volatile, unless that from some reason you are trying to take them away from the registers (and that makes no sense too).

4

Declaring those variables as volatile reduces register usage from 45 registers to 35

5

A bit of explanation for Noel Lopez: the multiplication to compute idx would otherwise get inlined by the compiler at several places, bringing register usage up needlessly (and total instruction count as well). The volatile trick forces the result of the computation into a register early on, so no inlining will occur later. This also works with constants (like the one, two, three) which would otherwise get loaded into registers over and over when accessed at several places.

The trick works, but is a little unintuitive. Usually the volatile keyword is a performance killer, but in this context (local variables in CUDA kernels) this helps greatly.

6

Sorry for the delay on the reply. Still this makes no sense to me:

Volatile basically tells the compiler that the variable value may change outside the function code. So it basically tells the compiler “do not to thrust the value of the variable that is on the register”.

This forces the compiler not to optimize code (hence the value of the variable is always read from memory).

So yes the result of the computation is calculated early and placed on the memory. Then each time its needed is taken from memory (This will for sure kill performance). But you gain registers. I still think its not a good idea - There are other ways to do things (You can limit the visibility of the variable).

I’m I wrong? Does CUDA compiler work diferently? volatile is C standard so it shouldn’t be open to different interpretations.

7

In a vast majority (90% of all cases), local automatic variables are stored in registers only and never written to memory.

Edit: nevermind, I think what originally had here was wrong. We need a PTX/decuda expert to explain the details of what nvcc does with volatile.