¿Por qué no se pueden compartir las variables miembro?

Question

Me gustaría crear una instancia de una clase en código CUDA, que comparte algunos de sus miembros con otros hilos en el mismo bloque.

Sin embargo, cuando intento compilar el siguiente código, aparece el error: »el atributo" compartido "no se aplica aquí« (nvcc versión 4.2).

class SharedSomething { 

public: 
    __shared__ int i; // this is not allowed 
}; 

__global__ void run() { 

    SharedSomething something; 
} 

¿Cuál es la razón detrás de eso? ¿Hay una solución alternativa para lograr el comportamiento deseado (miembros compartidos de una clase en un bloque)?

Answer 1

Rost explicó la razón de ser de la limitación. Para responder a la segunda parte de la pregunta, una solución simple es hacer que el kernel declare la memoria compartida e inicialice un puntero al que pertenece la clase, p. en el constructor de la clase. Ejemplo.

class Foo 
{ 
public: 
    __device__ 
    Foo(int *sPtr) : sharedPointer(sPtr, gPtr) { 
    sharedPointer[threadIdx.x] = gPtr[blockIdx.x * blockDim.x + threadIdx.x]; 
    __syncthreads(); 
    } 

    __device__ 
    void useSharedData() { printf("my data: %f\n", sharedPointer[threadIdx.x]); } 

private: 
    int *sharedPointer; 
}; 

__global__ void example(int *gData) 
{ 
    __shared__ int sData[BLOCKDIM]; 

    Foo f(sData, gData); 

    f.useSharedData(); 
} 

Advertencia: código escrito en el navegador, no verificada, no probado (y es un ejemplo trivial, pero el concepto se extiende a código real — He utilizado esta técnica yo).

Answer 2

Los objetos marcados como __shared__ residen en la memoria compartida dedicada por bloque de subprocesos. Tiene un tamaño limitado y tiene la misma duración que el bloque de hilos.

Este es el motivo por el que no puede declarar miembros de clase como compartidos: su duración no está gestionada por instancia de clase, sino por bloque de subprocesos. Posiblemente los miembros de la clase static podrían ser compartidos, pero no lo verificaron.

Consulte CUDA Programming Guide para más detalles, sección B.2.3.