PTX code "st"

Hi everyone,

I want to use ‘st’ to store a register t1 into a float variable, so I use the following PTX code in my cuda code:

asm(“st.global.f32\t[%0], t1;\n\t” : “=r” (&C[n]) :);

compiler tells me something was wrong: st.cu(22): error: expression must be a modifiable lvalue; st.cu(22): error: an asm operand must have scalar type

if I change the code like this:

asm(“st.global.f32\t[%0], t1;\n\t” : “=f” (C[n]) :);

compiling is OK. However, the result is wrong!! Why??

Thank you!

Use “r” instead of “=r” (after the correct number of colons). The address of the variable is input to the PTX instruction, not output from it.

Or you can use “=f” in the place where t1 gets set, replacing t1 with %0 and dropping the cited asm statement altogether.

Hi, I tried your method but got 2 errors. The code is " asm(“st.global.f32\t[%0], t1;\n\t” : “r” (&C[n]) :);".

The errors are:
st.cu(22): error: expression must be a modifiable lvalue

st.cu(22): error: an asm output operand must have one of the ‘=’ or ‘+’ modifiers

Input constraints are listed after the second colon, so the correct syntax is

asm("st.global.f32\t[%0], t1;\n\t" :: "r" (&C[n]));