tutorial - Conjuntos dinámicos CUDA globales(como en C) asignados a la memoria del dispositivo

programar gpu

Question

Por lo tanto, estoy tratando de escribir algún código que utiliza la arquitectura CUDA de Nvidia. Noté que copiar desde y hacia el dispositivo realmente estaba perjudicando mi rendimiento general, así que ahora estoy tratando de mover una gran cantidad de datos al dispositivo.

Como esta información se usa en numerosas funciones, me gustaría que sea global. Sí, puedo pasar punteros, pero realmente me gustaría saber cómo trabajar con globales en esta instancia.

Entonces, tengo funciones de dispositivo que quieren 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. Me di cuenta de un trabajo al declarar 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 así 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 concentrarse en la copiosa documentación ofrecida por 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 simple ejemplo de cómo asignar memoria. Ahora, en sus núcleos, debe aceptar un puntero a un flotador como ese:

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

Entonces 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 esta información se usa en numerosas funciones, me gustaría que sea global.

Hay pocas buenas razones para usar globales. Esto definitivamente no es uno. Lo dejo como ejercicio para expandir este ejemplo e incluir el movimiento de "devPtr" a un alcance global.

EDITAR:

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

Por lo tanto, lo más cerca que puedo sugerir es esto: use 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 olvides ''--host-compilation = c ++'' para este ejemplo.

Answer 3

Eh, ese era exactamente el problema de mover DevPtr al alcance global que era mi problema.

Tengo una implementación que hace exactamente eso, con los dos núcleos teniendo un puntero a los datos pasados. De manera explícita, no quiero pasar esos punteros.

He leído la documentación con bastante atención y he visitado los foros de nvidia (y google ha buscado durante aproximadamente una hora), pero no he encontrado una implementación de una matriz de dispositivos dinámicos globales que realmente se ejecute (he intentado varios que compilan y luego fracasan en formas nuevas e interesantes).

Answer 4

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 recibo "Aviso: no puedo decir a qué apunta el puntero, asumiendo el espacio de la memoria global" cada vez que intento acceder a los datos globales. Afortunadamente, la suposición es correcta, pero las advertencias son molestas.

De todos modos, para el registro, he analizado muchos de los ejemplos y repasé los ejercicios de nvidia, donde el objetivo es lograr que la salida diga "¡Correcto!". Sin embargo, no los he visto a todos . Si alguien sabe de un ejemplo de SDK donde hacen una asignación dinámica de la memoria global del dispositivo, me gustaría saber.

Answer 5

echa un vistazo a las muestras incluidas con el SDK. Muchos de esos proyectos de muestra son una forma decente de aprender con el ejemplo.

Answer 6

Como esta información se usa en numerosas funciones, me gustaría que sea global.

-

Hay pocas buenas razones para usar globales. Esto definitivamente no es uno. Lo dejo como ejercicio para expandir este ejemplo e incluir el movimiento de "devPtr" a un alcance global.

¿Qué pasa si el kernel opera en una gran estructura const compuesta de matrices? El uso de la llamada memoria constante no es una opción, porque tiene un tamaño muy limitado ... ¿entonces tienes que ponerlo en la memoria global ...?