Empurrar dentro de kernels escritos pelo usuário

Eu sou um novato em Thrust. Vejo que todas as apresentações e exemplos do Thrust mostram apenas o código do host.

Gostaria de saber se posso passar um device_vector para o meu próprio kernel? Como? Se sim, quais são as operações permitidas dentro do kernel / código do dispositivo?

Como foi originalmente escrito, Thrust é puramente uma abstração do lado do host. Não pode ser usado dentro de kernels. Você pode passar a memory do dispositivo encapsulada dentro de um thrust::device_vector para o seu próprio kernel assim:

 thrust::device_vector< Foo > fooVector; // Do something thrust-y with fooVector Foo* fooArray = thrust::raw_pointer_cast( &fooVector[0] ); // Pass raw array and its size to kernel someKernelCall<<< x, y >>>( fooArray, fooVector.size() ); 

e você também pode usar a memory do dispositivo não alocada por empuxo dentro de algoritmos de empuxo, instanciando um thrust :: device_ptr com o ponteiro da memory do dispositivo bare cuda.

Editado quatro anos e meio depois para acrescentar que, de acordo com a resposta do JackOLantern, o empuxo 1.8 adiciona uma política de execução sequencial, o que significa que você pode executar versões de encadeamento único de algoritmos de empuxo no dispositivo. Observe que ainda não é possível transmitir diretamente um vetor de dispositivo de empuxo para um kernel e os vetores de dispositivo não podem ser usados ​​diretamente no código do dispositivo.

Observe que também é possível usar a política de execução thrust::device em alguns casos para que a execução de thrust paralela seja iniciada por um kernel como uma grade filho. Isso requer uma compilation / conexão de dispositivo / hardware separada que suporte o paralelismo dynamic. Não tenho certeza se isso é realmente suportado em todos os algoritmos de empuxo ou não, mas certamente funciona com alguns.

Eu gostaria de fornecer uma resposta atualizada para esta pergunta.

A partir do Thrust 1.8, as primitivas de impulso CUDA podem ser combinadas com a política de execução thrust::seq para executar sequencialmente em um único thread CUDA (ou seqüencialmente em um único thread de CPU). Abaixo, um exemplo é relatado.

Se você deseja a execução paralela em um encadeamento, pode considerar o uso do CUB, que fornece rotinas de redução que podem ser chamadas de dentro de um encadeamento de encadeamentos, desde que sua placa permita o paralelismo dynamic.

Aqui está o exemplo com o Thrust

 #include  #include  #include  /********************/ /* CUDA ERROR CHECK */ /********************/ #define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); } inline void gpuAssert(cudaError_t code, char *file, int line, bool abort=true) { if (code != cudaSuccess) { fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line); if (abort) exit(code); } } __global__ void test(float *d_A, int N) { float sum = thrust::reduce(thrust::seq, d_A, d_A + N); printf("Device side result = %f\n", sum); } int main() { const int N = 16; float *h_A = (float*)malloc(N * sizeof(float)); float sum = 0.f; for (int i=0; i>>(d_A, N); } 

Esta é uma atualização da minha resposta anterior.

A partir do Thrust 1.8.1, as primitivas de impulso CUDA podem ser combinadas com a política de execução thrust::device para executar em paralelo dentro de um único paralelismo dynamic CUDA de exploração de thread CUDA. Abaixo, um exemplo é relatado.

 #include  #include  #include  #include "TimingGPU.cuh" #include "Utilities.cuh" #define BLOCKSIZE_1D 256 #define BLOCKSIZE_2D_X 32 #define BLOCKSIZE_2D_Y 32 /*************************/ /* TEST KERNEL FUNCTIONS */ /*************************/ __global__ void test1(const float * __restrict__ d_data, float * __restrict__ d_results, const int Nrows, const int Ncols) { const unsigned int tid = threadIdx.x + blockDim.x * blockIdx.x; if (tid < Nrows) d_results[tid] = thrust::reduce(thrust::seq, d_data + tid * Ncols, d_data + (tid + 1) * Ncols); } __global__ void test2(const float * __restrict__ d_data, float * __restrict__ d_results, const int Nrows, const int Ncols) { const unsigned int tid = threadIdx.x + blockDim.x * blockIdx.x; if (tid < Nrows) d_results[tid] = thrust::reduce(thrust::device, d_data + tid * Ncols, d_data + (tid + 1) * Ncols); } /********/ /* MAIN */ /********/ int main() { const int Nrows = 64; const int Ncols = 2048; gpuErrchk(cudaFree(0)); // size_t DevQueue; // gpuErrchk(cudaDeviceGetLimit(&DevQueue, cudaLimitDevRuntimePendingLaunchCount)); // DevQueue *= 128; // gpuErrchk(cudaDeviceSetLimit(cudaLimitDevRuntimePendingLaunchCount, DevQueue)); float *h_data = (float *)malloc(Nrows * Ncols * sizeof(float)); float *h_results = (float *)malloc(Nrows * sizeof(float)); float *h_results1 = (float *)malloc(Nrows * sizeof(float)); float *h_results2 = (float *)malloc(Nrows * sizeof(float)); float sum = 0.f; for (int i=0; i>>(d_data, d_results1, Nrows, Ncols); gpuErrchk(cudaPeekAtLastError()); gpuErrchk(cudaDeviceSynchronize()); printf("Timing approach nr. 1 = %f\n", timerGPU.GetCounter()); gpuErrchk(cudaMemcpy(h_results1, d_results1, Nrows * sizeof(float), cudaMemcpyDeviceToHost)); for (int i=0; i>>(d_data, d_results1, Nrows, Ncols); gpuErrchk(cudaPeekAtLastError()); gpuErrchk(cudaDeviceSynchronize()); printf("Timing approach nr. 2 = %f\n", timerGPU.GetCounter()); gpuErrchk(cudaMemcpy(h_results1, d_results1, Nrows * sizeof(float), cudaMemcpyDeviceToHost)); for (int i=0; i 

O exemplo acima realiza reduções das linhas de uma matriz no mesmo sentido que Reduzir linhas de matriz com CUDA , mas é feito de forma diferente do post acima, a saber, chamando primitivas de impulso de CUDA diretamente de kernels escritos pelo usuário. Além disso, o exemplo acima serve para comparar o desempenho das mesmas operações quando feito com duas políticas de execução, a saber, thrust::seq e thrust::device . Abaixo, alguns charts mostrando a diferença no desempenho.

Horários

Speedups

O desempenho foi avaliado em um Kepler K20c e em um Maxwell GeForce GTX 850M.

Se você pretende usar os dados alocados / processados ​​pelo empuxo sim, você pode, apenas obter o ponteiro bruto dos dados alocados.

 int * raw_ptr = thrust::raw_pointer_cast(dev_ptr); 

Se você quiser alocar vetores de empuxo no kernel, eu nunca tentei, mas acho que não funcionará e, se funcionar, não acho que isso trará nenhum benefício.