在内联ptx中加载函数参数答案

【问题标题】：load function parameters in inlined ptx在内联ptx中加载函数参数
【发布时间】：2013-04-16 16:23:04
【问题描述】：

我在 32 位 Visual Studio 2008 的调试模式下具有以下内联汇编功能：

__device__ void add(int* pa, int* pb)
{
  asm(".reg .u32   s<3>;"::);
  asm(".reg .u32   r<14>;"::);

  asm("ld.global.b32    s0, [%0];"::"r"(&pa));      //load addresses of pa, pb
  printf(...);
  asm("ld.global.b32    s1, [%0];"::"r"(&pb));
  printf(...);
  asm("ld.global.b32    r1, [s0+8];"::);
  printf(...);  
  asm("ld.global.b32    r2, [s1+8];"::);
  printf(...);

  ...// perform some operations
}

pa和pb在设备上全局分配如

__device__ int pa[3] = {0, 0x927c0000, 0x20000011};  
__device__ int pb[3] = {0, 0xbb900000, 0x2000000b};

但是，此代码在发布模式下失败，在线 asm("ld.global.b32 r1, [s0+8];"::); 如何在发布模式下使用 inline ptx 正确加载函数参数？

附：使用 -G 标志（生成 GPU 调试信息）构建发布模式会导致代码在发布模式下正确运行。谢谢，

【问题讨论】：

您的发行版也是 32 位版本吗？否则这条线显然会失败。如果不是：如果您告诉我们错误消息是什么，则帮助处理错误消息会容易得多。
实体pa可能在release模式下被优化成寄存器，不能取寄存器地址。这只是一个猜测。如果您定义了您观察到的故障是在编译时还是在运行时，这将很有帮助。如果在编译时，你到底得到了什么错误？如果在运行时，您是如何本地化到这行代码的？为什么要取pa的地址呢？获取函数参数的地址似乎很奇怪。要在调用上下文中修改pa，你需要传递一个指向它的指针，即**pa 你的意思是*pa or &(pa[0]) or pa?
@tera：发布版本也是 32 位的。发布版本编译没有任何错误。当我跑步时，我得到不同的输出。此外，我在每个 ld.global.b32 指令之后都有 printf 语句。它在 asm("ld.global.b32 r1, [s0+8];"::); 之后失败这样它就不会打印下一个 printf 语句。
正如罗伯特所说，获取函数指针的地址很奇怪。为什么不直接移动参数本身：asm("mov.b32 s0, %0;"::"r"(pa));？
@RobertCrovella：错误出现在运行时。我得到完全错误的输出； asm("ld.global.b32 r1, [s0+8];"::); 之后的 printf 和其他指令没有执行，它们被跳过，然后程序以错误的输出退出。

标签： visual-studio-2008 cuda ptx

【解决方案1】：

希望这段代码会有所帮助。我仍在猜测您到底要做什么，但我从您的代码开始，并决定在 pa 和 pb 数组中添加一些值并将它们存储回 pa[0] 和 pb[0]。

此代码是为 64 位机器编写的，但将其转换为 32 位指针应该不难。我已经用注释标记了需要为 32 位指针更改的行。希望这将回答您关于如何使用作为设备内存指针的函数参数的问题：

#include <stdio.h>

__device__ int pa[3] = {0, 0x927c0000, 0x20000011};
__device__ int pb[3] = {0, 0xbb900000, 0x2000000b};

__device__ void add(int* mpa, int* mpb)
{
  asm(".reg .u64   s<2>;"::);  // change to .u32 for 32 bit pointers
  asm(".reg .u32   r<6>;"::);

  asm("mov.u64    s0, %0;"::"l"(mpa));      //change to .u32 and "r" for 32 bit
  asm("mov.u64    s1, %0;"::"l"(mpb));      //change to .u32 and "r" for 32 bit
  asm("ld.global.u32    r0, [s0+4];"::);
  asm("ld.global.u32    r1, [s1+4];"::);
  asm("ld.global.u32    r2, [s0+8];"::);
  asm("ld.global.u32    r3, [s1+8];"::);
  asm("add.u32    r4, r0, r2;"::);
  asm("add.u32    r5, r1, r3;"::);
  asm("st.global.u32    [s0], r4;"::);
  asm("st.global.u32   [s1], r5;"::);
}

__global__ void mykernel(){
  printf("pa[0] = %x, pb[0] = %x\n", pa[0], pb[0]);
  add(pa, pb);
  printf("pa[0] = %x, pb[0] = %x\n", pa[0], pb[0]);
}

int  main() {
  mykernel<<<1,1>>>();
  cudaDeviceSynchronize();
  return 0;
}

当我运行这段代码时，我得到：

$ ./t128
pa[0] = 0, pb[0] = 0
pa[0] = b27c0011, pb[0] = db90000b
$

我认为这是正确的输出。

我编译它：

nvcc -O3 -arch=sm_20 -o t128 t128.cu

【讨论】：