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]));