¿Cómo usar la memoria constante CUDA de forma agradable?

Question

Estoy trabajando en una aplicación de cálculo numérico utilizando el marco CUDA. Tengo algunos datos estáticos que deben ser accesibles a todas las roscas, por lo que lo he puesto en la memoria constante de esta manera:

__device__ __constant__ CaseParams deviceCaseParams; 

utilizo el cudaMemcpyToSymbol llamada a transferir estos parametros desde el host al dispositivo:

void copyMetaData(CaseParams* caseParams) 
{ 
    cudaMemcpyToSymbol("deviceCaseParams", caseParams, sizeof(CaseParams)); 
} 

que funciona.

De todos modos, parece (por ensayo y error, y también por mensajes de lectura en la red) que por alguna razón enferma, la declaración de deviceCaseParams y la operación de copia de la misma (la llamada a cudaMemcpyToSymbol) deben estar en el mismo archivo. Por el momento tengo estos dos en un archivo .cu, pero realmente quiero tener el parámetro struct en un archivo .cuh para que cualquier implementación pueda verlo si lo desea. Eso significa que también tengo que tener la función copyMetaData en el archivo de encabezado a, pero esto arruina el enlace (el símbolo ya está definido) ya que ambos archivos .cpp y .cu incluyen este encabezado (y así el compilador MS C++ y nvcc lo compilan)

¿Alguien tiene algún consejo sobre diseño aquí?

Actualización: Ver los comentarios

Answer 1

Con una CUDA actualizada (por ejemplo, 3.2) que debe ser capaz de hacer el establecimiento de memoria desde el interior de una unidad de traducción diferente si usted está mirando hacia arriba en el símbolo tiempo de ejecución (es decir, pasando una cadena como la primera arg a cudaMemcpyToSymbol como lo hace en su ejemplo).

Además, con los dispositivos de la clase Fermi, puede malloc la memoria (cudaMalloc), copiarla en la memoria del dispositivo y luego pasar el argumento como un puntero const. El compilador reconocerá si está accediendo a los datos uniformemente a través de los warps y, de ser así, usará el caché constante. Consulte la Guía de programación de CUDA para obtener más información. Nota: necesitaría compilar con -arch=sm_20.

Answer 2

Si está utilizando el CUDA pre-Fermi, ya habrá descubierto que este problema no solo se aplica a la memoria constante, se aplica a cualquier cosa que desee del lado CUDA. Las dos únicas maneras que he encontrado en torno a esta son ya sea a:

1. Escribir todo CUDA en un solo archivo (.cu), o
2. Si necesita romper el código en archivos separados, se limite a las cabeceras que su archivo .cu solo incluye

Si necesita compartir el código entre CUDA y C/C++, o tiene algún código común que comparte entre proyectos, la opción 2 es la única opción. Parece muy antinatural para empezar, pero resuelve el problema. Aún puedes estructurar tu código, pero no de una manera típica de C. La sobrecarga principal es que cada vez que haces una compilación compilas todo. El lado positivo de esto (que creo que es posiblemente la razón por la que funciona de esta manera) es que el compilador CUDA tiene acceso a todo el código fuente en un solo hit, lo cual es bueno para la optimización.