c++ - Clase virtual Cuda

virtual-functions

Question

Me gustaría ejecutar algunos métodos virtuales en un kernel cuda, pero en lugar de crear el objeto en el mismo kernel, me gustaría crearlo en el host y copiarlo a la memoria gpu.

Estoy creando con éxito objetos en un kernel y llamo a un método virtual. El problema surge al copiar el objeto. Esto tiene sentido porque obviamente el puntero de la función virtual es falso. Lo que sucede es simplemente "Falló el lanzamiento de la red Cuda", al menos esto es lo que dice Nsight. Pero al echarle un vistazo al SASS, falla al desreferenciar el puntero de función virtual, lo que tiene sentido.

Por supuesto, estoy usando Cuda 4.2 y compilando con "compute_30" en una tarjeta de adaptación.

Entonces, ¿cuál es el camino recomendado para seguir? ¿O esta característica simplemente no es compatible?

Tuve la idea de ejecutar un kernel diferente primero que crea objetos ficticios y extraer el puntero de la función virtual para "parchar" mis objetos antes de copiarlos. Lamentablemente, esto no funciona realmente (aún no lo hemos descubierto) y sería una solución desagradable.

PD. Esto es en realidad una repetición de esta pregunta, que lamentablemente nunca fue respondida completamente.

Editar:

Entonces encontré la manera de hacer lo que quería. Pero solo para ser claros: esto no es en absoluto una respuesta o solución, la respuesta ya fue provista, esto es solo un truco, solo por diversión.

Así que primero veamos qué hace Cuda cuando llamamos a un método virtual, a continuación se muestra la depuración SASS

//R0 is the address of our object
LD.CG R0, [R0];  
IADD R0, R0, 0x4;  
NOP;  
MOV R0, R0;  
LD.CG R0, [R0];
...
IADD R0, RZ, R9;  
MOV R0, R0;  
LDC R0, c[0x2][R0];
...
BRX R0 - 0x5478

Asumiendo que "c [0x2] [INDICE]" es constante para todos los kernels, podemos obtener el índice de una clase simplemente ejecutando un kernel y haciendo esto, donde obj es un objeto recién creado de la clase que mira:

unsigned int index = *(unsigned int*)(*(unsigned int*)obj + 4);

Luego usa algo como esto:

struct entry
{
    unsigned int vfptr;// := &vfref, thats our value to store in an object
    int dummy;// := 1234, great for debugging
    unsigned int vfref;// := &dummy
    unsigned int index;
    char ClassName[256];//use it as a key for a dict
};

Almacene esto en el host así como en la memoria del dispositivo (las ubicaciones de la memoria son las del dispositivo) y en el host puede usar ClassName como búsqueda de un objeto para "parchear".

Pero otra vez: no usaría esto en algo serio, porque las funciones virtuales de rendimiento no son para nada buenas.

Answer 1

Lo que estás tratando de hacer no es compatible, actualmente, con el compilador CUDA y el tiempo de ejecución (a partir de CUDA 5.0). La Sección D.2.6.3 de la Guía de programación de CUDA C v5.0 dice:

D.2.6.3 Funciones virtuales

Cuando una función en una clase derivada anula una función virtual en una clase base, los calificadores del espacio de ejecución (es decir, __host__ , __device__ ) en las funciones anuladas y anuladas deben coincidir.

No se permite pasar como argumento a una función __global__ un objeto de una clase con funciones virtuales.

La tabla de funciones virtuales se coloca en la memoria global o constante por el compilador.

Lo que recomiendo es que encapsules los datos de tu clase por separado de la funcionalidad de la clase. Por ejemplo, almacene los datos en una estructura. Si planea operar en matrices de estos objetos, almacene los datos en una estructura de matrices (para rendimiento, fuera del alcance de esta pregunta). Asigne las estructuras de datos en el host usando cudaMalloc , y luego pase los datos al kernel como argumentos, en lugar de pasar la clase con métodos virtuales.

Luego construye tus objetos con métodos virtuales en el dispositivo. El constructor de su clase con métodos virtuales tomaría los parámetros del núcleo del puntero del dispositivo como argumentos. Los métodos del dispositivo virtual podrían funcionar en los datos del dispositivo.

El mismo enfoque funcionaría para permitir asignar los datos en un kernel en el dispositivo y acceder a él en otro kernel del dispositivo (ya que de nuevo, las clases con funciones virtuales no pueden ser parámetros para los kernels).