Usando a variável de dispositivo por vários threads no CUDA

Question

Estou brincando com Cuda.

No momento, tenho um problema. Estou testando uma grande matriz para respostas específicas e, quando recebo a resposta, tenho que copiar os dados para outra matriz.

Por exemplo, minha matriz de teste de 5 elementos se parece com o seguinte:
] [] [v1] [] [] [v2

O resultado deve ser assim:
v1] [v2

O problema é como calcular o endereço da segunda matriz para armazenar o resultado? Todos os elementos da primeira matriz são verificados em paralelo.

Estou pensando em declarar uma variável de dispositivo int addr = 0. Toda vez que encontro uma resposta, incrementarei o addr. Mas não tenho certeza disso porque significa que addr pode ser acessado por vários threads ao mesmo tempo. Isso causará problemas? Ou o thread aguardará até que outro encadeamento termine usando essa variável?

Answer 1

Não é tão trivial quanto parece. Acabei de terminar de implementar um e posso dizer o que você precisa ler o Scan GPU Gems 3 Artigo em particular capítulo 39.3.1 Compactação de fluxo.

Para implementar seu próprio começo a partir do exemplo de LargearRaysCan no SDK, isso lhe dará apenas o Prescan. Supondo que você tenha a matriz de seleção na memória do dispositivo (uma matriz de 1 e 0 que significa 1- Selecione 0 - descarte), dev_selection_array uma dev_elements_array elementos a serem selecionados um dev_prescan_array e a dev_result_array Todo o tamanho N Então você faz

prescan(dev_prescan_array,dev_selection_array, N);
scatter(dev_result_array, dev_prescan_array,
         dev_selection_array, dev_elements_array, N);

onde está a dispersão

 __global__ void scatter_kernel( T*dev_result_array, 
                   const T* dev_prescan_array, 
                   const T* dev_selection_array,
                   const T* dev_elements_array, std::size_t size){

unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx >= size) return;
if (dev_selection_array[idx] == 1){
    dev_result_array[dev_prescan_array[idx]] = dev_elements_array[idx];
}
}

Para outra boa aplicação do Prescan, consulte o artigo Ble93

Divirta-se!

Answer 2

Você está falando sobre compactação clássica de fluxo. Geralmente eu recomendaria olhar para Impulso ou CUDPP (Esses links vão para a documentação da compactação). Ambos são de código aberto, se você quiser rolar o seu próprio, eu também sugiro olhar para a amostra SDK 'Scan'.