Matrices dinámicas CUDA globales (como en C) asignadas a la memoria del dispositivo

Question

Entonces, estoy tratando de escribir un código que utilice la arquitectura CUDA de Nvidia.Me di cuenta de que copiar hacia y desde el dispositivo realmente perjudicaba mi rendimiento general, por lo que ahora estoy intentando mover una gran cantidad de datos al dispositivo.

Como estos datos se utilizan en numerosas funciones, me gustaría que fueran globales.Sí, puedo transmitir sugerencias, pero realmente me gustaría saber cómo trabajar con valores globales en este caso.

Entonces, tengo funciones de dispositivo que desean acceder a una matriz asignada por dispositivo.

Idealmente, podría hacer algo como:

__device__ float* global_data;

main()
{
  cudaMalloc(global_data);
  kernel1<<<blah>>>(blah); //access global data
  kernel2<<<blah>>>(blah); //access global data again
}

Sin embargo, no he descubierto cómo crear una matriz dinámica.Descubrí una solución declarando la matriz de la siguiente manera:

__device__ float global_data[REALLY_LARGE_NUMBER];

Y aunque eso no requiere una llamada a cudaMalloc, preferiría el enfoque de asignación dinámica.

Answer 1

Algo como esto probablemente debería funcionar.

#include <algorithm>

#define NDEBUG
#define CUT_CHECK_ERROR(errorMessage) do {                                 \
        cudaThreadSynchronize();                                           \
         cudaError_t err = cudaGetLastError();                             \
         if( cudaSuccess != err) {                                         \
                     fprintf(stderr, "Cuda error: %s in file '%s' in line %i : %s.\n",    \
                                             errorMessage, __FILE__, __LINE__, cudaGetErrorString( err) );\
                     exit(EXIT_FAILURE);                                                  \
                 } } while (0)


__device__ float *devPtr;

__global__
void kernel1(float *some_neat_data)
{
    devPtr = some_neat_data;
}

__global__
void kernel2(void)
{
    devPtr[threadIdx.x] *= .3f;
}


int main(int argc, char *argv[])
{
    float* otherDevPtr;
    cudaMalloc((void**)&otherDevPtr, 256 * sizeof(*otherDevPtr));
    cudaMemset(otherDevPtr, 0, 256 * sizeof(*otherDevPtr));

    kernel1<<<1,128>>>(otherDevPtr);
    CUT_CHECK_ERROR("kernel1");

    kernel2<<<1,128>>>();

    CUT_CHECK_ERROR("kernel2");

    return 0;
}

Darle un giro.

Answer 2

Dedique un tiempo a centrarse en la abundante documentación que ofrece NVIDIA.

De la Guía de programación:

float* devPtr;
cudaMalloc((void**)&devPtr, 256 * sizeof(*devPtr));
cudaMemset(devPtr, 0, 256 * sizeof(*devPtr));

Ese es un ejemplo simple de cómo asignar memoria.Ahora, en tus núcleos, deberías aceptar un puntero a un flotante como este:

__global__
void kernel1(float *some_neat_data)
{
    some_neat_data[threadIdx.x]++;
}

__global__
void kernel2(float *potentially_that_same_neat_data)
{
    potentially_that_same_neat_data[threadIdx.x] *= 0.3f;
}

Ahora puedes invocarlos así:

float* devPtr;
cudaMalloc((void**)&devPtr, 256 * sizeof(*devPtr));
cudaMemset(devPtr, 0, 256 * sizeof(*devPtr));

kernel1<<<1,128>>>(devPtr);
kernel2<<<1,128>>>(devPtr);

Como estos datos se utilizan en numerosas funciones, me gustaría que sea global.

Hay pocas buenas razones para utilizar globales.Este definitivamente no es uno.Lo dejaré como ejercicio para ampliar este ejemplo e incluir mover "devPtr" a un ámbito global.

EDITAR:

Ok, el problema fundamental es este:sus núcleos solo pueden acceder a la memoria del dispositivo y los únicos punteros de alcance global que pueden usar son los de GPU.Cuando se llama a un kernel desde su CPU, lo que sucede detrás de escena es que los punteros y las primitivas se copian en los registros de la GPU y/o en la memoria compartida antes de que se ejecute el kernel.

Entonces lo más cercano que puedo sugerir es esto:utilice cudaMemcpyToSymbol() para lograr sus objetivos.Pero, en el fondo, considere que un enfoque diferente podría ser lo correcto.

#include <algorithm>

__constant__ float devPtr[1024];

__global__
void kernel1(float *some_neat_data)
{
    some_neat_data[threadIdx.x] = devPtr[0] * devPtr[1];
}

__global__
void kernel2(float *potentially_that_same_neat_data)
{
    potentially_that_same_neat_data[threadIdx.x] *= devPtr[2];
}


int main(int argc, char *argv[])
{
    float some_data[256];
    for (int i = 0; i < sizeof(some_data) / sizeof(some_data[0]); i++)
    {
        some_data[i] = i * 2;
    }
    cudaMemcpyToSymbol(devPtr, some_data, std::min(sizeof(some_data), sizeof(devPtr) ));
    float* otherDevPtr;
    cudaMalloc((void**)&otherDevPtr, 256 * sizeof(*otherDevPtr));
    cudaMemset(otherDevPtr, 0, 256 * sizeof(*otherDevPtr));

    kernel1<<<1,128>>>(otherDevPtr);
    kernel2<<<1,128>>>(otherDevPtr);

    return 0;
}

No olvide '--host-compilation=c++' para este ejemplo.

Answer 3

Seguí adelante y probé la solución de asignar un puntero temporal y pasarlo a una función global simple similar a kernel1.

La buena noticia es que funciona :)

Sin embargo, creo que confunde al compilador ya que ahora aparece "Aviso:No puedo saber a qué apunta el puntero, asumiendo espacio de memoria global" cada vez que intento acceder a los datos globales.Afortunadamente, la suposición resulta ser correcta, pero las advertencias son molestas.

De todos modos, para que conste, he visto muchos de los ejemplos y realicé los ejercicios de nvidia donde el objetivo es conseguir que el resultado diga "¡Correcto!".Sin embargo, no he mirado todo de ellos.Si alguien conoce un ejemplo de SDK en el que realizan una asignación dinámica de memoria global del dispositivo, todavía me gustaría saberlo.

Answer 4

Erm, mi problema era exactamente ese problema de mover devPtr al alcance global.

Tengo una implementación que hace exactamente eso, con los dos núcleos teniendo un puntero a los datos pasados.Explícitamente no quiero pasar esos consejos.

Leí la documentación con bastante atención y visité los foros de nvidia (y busqué en Google durante aproximadamente una hora), pero no encontré una implementación de una matriz de dispositivos dinámicos globales que realmente se ejecute (he probado varios que compilan y luego falla de maneras nuevas e interesantes).

Answer 5

consulte los ejemplos incluidos con el SDK.Muchos de esos proyectos de muestra son una forma decente de aprender con el ejemplo.

Answer 6

Como estos datos se utilizan en numerosas funciones, me gustaría que fueran globales.

-

Hay pocas buenas razones para utilizar globales.Este definitivamente no es uno.Dejaré como un ejercicio para expandir este ejemplo para incluir el movimiento "DevPTR" a un alcance global.

¿Qué pasa si el kernel opera en una estructura constante grande que consta de matrices?Usar la llamada memoria constante no es una opción, porque tiene un tamaño muy limitado.¿Entonces tienes que ponerlo en la memoria global?