Параметры функции загрузки в встроенном ptx

Question

У меня есть следующая функция со встроенной сборкой, которая отлично работает в режиме отладки в 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];"::);Как правильно загрузить параметры функции с помощью встроенного ptx в режиме разблокировки?

PS построение режима выпуска с флагом -G (генерирует отладочную информацию GPU) заставляет код правильно работать в режиме выпуска. Спасибо,

Answer 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