I want to prefetch in cuda But Error occur

Question

I am a new comer to cuda. And I want to add prefetch to cuda.

I add this to code.

_device_ void a(...)..{
   ....
   double value;
   ...
   asm("prefetch.global.L2 [&value];");
...}

but, I get this error.

ptxas /tmp/tmpxft_00000259_00000000-6_Simulation.ptx, line 279; fatal   : Parsing error near '&': syntax error
ptxas <macro util>, line 12; error   : Illegal modifier '.L2' for instruction 'mov'
ptxas <macro util>, line 12; error   : Illegal modifier '.L2' for instruction 'mov'
ptxas fatal   : Ptx assembly aborted due to errors

Can anybody help me T^T

Answer 1

First of all this:

double value;

is not in the global space in the GPU. It is in the logical local space. You can find info about logical spaces in the PTX documentation

Next, when using PTX instructions, you are generally working with registers. You should probably also familiarize yourself with the inline PTX documentation.

This seems to compile for me:

#include <cstdio>
__device__ void a(){
   double value = 1.0;
   asm("prefetch.local.L2 [%0];" :: "l"(&value));
   printf("%f\n", value);
}

this may also be of interest.