Matrizes dinâmicas globais CUDA (como em C) alocadas na memória do dispositivo

Question

Então, estou tentando escrever algum código que utilize a arquitetura CUDA da Nvidia.Percebi que copiar de e para o dispositivo estava prejudicando muito meu desempenho geral, então agora estou tentando mover uma grande quantidade de dados para o dispositivo.

Como esses dados são usados ​​em inúmeras funções, gostaria que fossem globais.Sim, posso passar dicas, mas gostaria muito de saber como trabalhar com globais neste caso.

Portanto, tenho funções de dispositivo que desejam acessar uma matriz alocada de dispositivo.

Idealmente, eu poderia fazer algo como:

__device__ float* global_data;

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

No entanto, ainda não descobri como criar uma matriz dinâmica.Eu descobri uma solução declarando o array da seguinte maneira:

__device__ float global_data[REALLY_LARGE_NUMBER];

E embora isso não exija uma chamada cudaMalloc, eu preferiria a abordagem de alocação dinâmica.

Answer 1

Algo assim provavelmente deveria 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;
}

Dê uma olhada.

Answer 2

Passe algum tempo concentrando-se na abundante documentação oferecida pela NVIDIA.

Do Guia de Programação:

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

Esse é um exemplo simples de como alocar memória.Agora, em seus kernels, você deve aceitar um ponteiro para um float assim:

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

Então agora você pode invocá-los assim:

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

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

Como esses dados são usados ​​em inúmeras funções, gostaria que fossem globais.

Existem poucos bons motivos para usar globais.Este definitivamente não é um deles.Deixarei como exercício expandir este exemplo para incluir a mudança de "devPtr" para um escopo global.

EDITAR:

Ok, o problema fundamental é este:seus kernels só podem acessar a memória do dispositivo e os únicos ponteiros de escopo global que eles podem usar são os da GPU.Ao chamar um kernel de sua CPU, nos bastidores o que acontece é que os ponteiros e primitivos são copiados para os registros da GPU e/ou memória compartilhada antes que o kernel seja executado.

Então o mais próximo que posso sugerir é o seguinte:use cudaMemcpyToSymbol() para atingir seus objetivos.Mas, no fundo, considere que uma abordagem diferente pode ser a coisa certa.

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

Não se esqueça de '--host-compilation=c++' para este exemplo.

Answer 3

Fui em frente e tentei a solução de alocar um ponteiro temporário e passá-lo para uma função global simples semelhante ao kernel1.

A boa notícia é que funciona :)

No entanto, acho que confunde o compilador, pois agora recebo "Aviso:Não consigo dizer para onde o ponteiro aponta, assumindo espaço de memória global" sempre que tento acessar os dados globais.Felizmente, a suposição está correta, mas os avisos são irritantes.

De qualquer forma, para que conste - eu olhei muitos dos exemplos e fiz os exercícios da nvidia onde o objetivo é fazer com que a saída diga "Correto!".No entanto, eu não olhei todos deles.Se alguém souber de um exemplo de SDK em que alocação dinâmica de memória de dispositivo global, eu ainda gostaria de saber.

Answer 4

Erm, era exatamente esse problema de mover o devPtr para o escopo global que era o meu problema.

Eu tenho uma implementação que faz exatamente isso, com os dois kernels tendo um ponteiro para os dados passados.Eu explicitamente não quero passar essas dicas.

Eu li a documentação com bastante atenção e acessei os fóruns da nvidia (e o Google pesquisou por cerca de uma hora), mas não encontrei uma implementação de uma matriz de dispositivo dinâmico global que realmente seja executada (tentei vários que compilam e então fracasse de maneiras novas e interessantes).

Answer 5

confira os exemplos incluídos no SDK.Muitos desses exemplos de projetos são uma maneira decente de aprender pelo exemplo.

Answer 6

Como esses dados são usados ​​em inúmeras funções, gostaria que fossem globais.

-

Existem poucos bons motivos para usar globais.Este definitivamente não é um deles.Deixarei como exercício expandir este exemplo para incluir a movimentação de "devptr" para um escopo global.

E se o kernel operar em uma grande estrutura const composta por arrays?Usar a chamada memória constante não é uma opção, porque seu tamanho é muito limitado.então você tem que colocá-lo na memória global ..?