visual-studio-2008 - tag - visual studio code shortcuts html

cargar los parámetros de la función en ptx en línea

Question

Tengo la siguiente función con ensamblado en línea que funciona bien en el modo de depuración en Visual Studio 2008 de 32 bits:

__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 y pb se asignan globalmente en el dispositivo, como

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

Sin embargo, este código falla en el modo de lanzamiento, en línea asm("ld.global.b32 r1, [s0+8];"::); ¿Cómo puedo cargar los parámetros de la función correctamente con ptx en línea en el modo de lanzamiento?

PS construye el modo de lanzamiento con el indicador -G (Genera información de depuración de GPU) hace que el código se ejecute correctamente en el modo de lanzamiento. Gracias,

Answer 1

Espero que este código ayude. Todavía estoy adivinando lo que estás tratando de hacer exactamente, pero comencé con tu código y decidí agregar algunos valores en las matrices pa y pb y almacenarlos de nuevo en pa[0] y pb[0] .

Este código está escrito para una máquina de 64 bits, pero convertirlo en punteros de 32 bits no debería ser difícil. He marcado las líneas que deben cambiarse para punteros de 32 bits con un comentario. Esperemos que esto responda a su pregunta sobre cómo usar los parámetros de función que son punteros a la memoria del dispositivo:

#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;
}

Cuando ejecuto este código obtengo:

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

que creo que es la salida correcta.

Lo compilé con:

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