在CUDA程序中嵌套PTX遇到了一些问题
现在要进行a+b=c 这个运算，
（1）将变量a 存入寄存器f1， asm[a，f1]
（2）将变量b存入寄存器f2， asm[a，f2]
（3）将寄存器f1和寄存器f2加起来，得到的结果存入寄存器f3， f3=f1+f2
（4）然后将结果返回给c。 c=f3
整个过程是这样，现在我需要实现的是将（3）这个步骤重复多次计算，这里我了一个for循环语句，但是因为前面是两个asm
这样的话，执行到f3=f2+f1的时候，此处无法识别f1和f2这两个寄存器。我试着把这四步写到一个asm里面，但是这个循环语句在asm里面我又不知道该如何去写。请问是否有方法解决，还有就是关于asm这个ptx文档讲得十分粗略，有没有什么地方可以看到比较详细的介绍。非常感谢！
楼主您好，您的asm[a,f1]这种格式我没有看懂。
按照您的要求，我为您重写了你的这个步骤，您看看是否符合：
（假设您的数据类型是int)
asm (“add.s32 %0, %1, %2; "
: “=r”(c)
: “r”(a), “r”(b)
);
%1和%2是分别对应a,b(如果他们不是寄存器，则编译器自动生成代码给您装入寄存器)。%0是结果的c。注意c使用了”=r"，作为输出并被破坏掉。
如果您是一个循环，例如您有int alex[100]和int xu[100]两个数组，并将结果写入int ice[100], 您需要在kernel里循环（假设您真的需要这样), 那么您可以这样：
for (int i = 0; i < 100; i++)
asm (“add.s32 %0, %1, %2; " :”=r"(ice[i]) : “r”(alex[i]), “r”(xu[i]));
如果您不满意，请继续回帖指出您的需要。我将继续为您服务。
楼主已经15分钟没有回应了，那我继续猜测您的意思。
您想体验一下控制的感觉，所以您想要在嵌入的ptx中，自己完成整个循环（而不是依赖外层的C)? 那也可以，您可以这样：
(以上是上文三个数组int alex[100]; int xu[100];和int ice[100];为例)
asm volatile
(“”
“{ .reg .u32 p_ice;” //用{}是为防止名字冲突。定义3个32位指针（假设您是32位环境）
" .reg .u32 p_alex;"
" .reg .u32 p_xu;"
" mov.u32 p_ice, %0;" //初始化p_ice, p_alex和p_xu
" mov.u32 p_alex, %1;"
" mov.u32 p_xu, %2;"
“”
" .reg .s32 count, t<3>;"
" .reg .pred p;"
" mov.s32 count, 0;" //循环变量初始化
“”
“alex_xu: ld.s32 t0, [p_alex];” //读取2个源
" ld.s32 t1, [p_xu];"
" add.s32 t2,t0,t1;" //加法
" st.s32 [p_ice], t2;" //回写1个结果
" add.s32 p_ice, p_ice, 4;" //修改指针（实际jit出来的sass可能无这几条）
" add.s32 p_alex, p_alex, 4;"
" add.s32 p_xu, p_xu, 4;"
" add.s32 count, count, 1;" //修改循环控制变量
" setp.ne.s32 p, count, 100;" //到100次就结束
" @p bra alex_xu;" //不到就继续循环
" }" :“=r”(c) : “r”(a), “r”(b) : “memory”
);
这样有快感么？您可能有，当然也可能没有。
这只是用于教学用途的，实际上，这个的结果和您直接写C, 并无区别。
我不建议您写ptx. 除非您觉得CUDA C的表达能力不够，无法满足您。
额。上个帖子不小心写顺手了。。
这里 :“=r”(c) : “r”(a), “r”(b) : “memory”
改成： :: “r”(ice), “r”(alex), “r”(xu) : “memory”
（我老是想着你1楼的a,b,c三个变量去了。忘记我举的例子不同。。）
非常感谢您的回复 只是我是想把将变量存入寄存器和执行相加这两个分开，然后单独地循环执行相加这个动作 用多个asm去定义的话寄存器识别不了 用一个asm的话 我又不知道在asm内部如何写循环语句 所以想问问你们有没有什么办法
抱歉，我完全不懂PTX，请延教高明。
算了，刚才是我生气你丝毫不看我的回复。
如果在1个asm里面写循环我已经在3#示范你了，麻烦看一下吧。否则很伤人的。
此外，如何分开这个简单。你需要定义几个你自己用到的名字。然后将a,b的值给它们即可。
例如：asm volatile (“.reg .s32 alex_f, alex_k; mov.s32 alex_f, %0; mov.s32 alex_k, %1;” :: “r”(a), “r”(b)); 这样即可。注意这个实际上不会产生额外的实质指令的，如果a,b已经是在寄存器中了。
然后你下文可以过一段再用他们(alex_f, alex_k), 直接下下文使用这个2个名字参与运算即可。例如：asm volatile (“add.s32 %0, alex_f, alex_k;”:“=r”(你的结果变量)); 即可。
alex_f, alex_k只要你不是在{}中用.reg定义的，那么一直有效到函数结束。
这个实际上在前文2#也给你示范过了。麻烦看一下。
有点心。给你回复打字也是不容易的。给你写了2楼，三个示范，你一点都不看，直接上去继续问，真的伤人。
好吧 确实不好意思 我是刚开始学的 不是很熟析 因为没学汇编语言 非常感谢您的回复 为我的失礼举动向你道歉
嘟嘟。