Упор внутри написанных пользователем ядер

Question

Я новичок в Thrust. Я вижу, что все презентации и примеры Thrust только показывают код хоста.

Я хотел бы знать, могу ли я передать device_vector в свое собственное ядро? Как? Если да, то какие операции разрешены на нем внутри кода ядра/устройства?

Answer 1

Если вы хотите использовать данные, выделенные/обработанные нажатием да, вы можете просто получить необработанный указатель выделенных данных.

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

если вы хотите выделить упорных векторы в ядре я никогда не пробовал, но я не думаю, что будет работать , а также, если это работает, я не думаю, что это даст какую-либо выгоду.

Answer 2

Я хотел бы дать обновленный ответ на этот вопрос.

Начиная с Thrust 1.8, примитивы CUDA Thrust могут быть объединены с политикой выполнения thrust::seq для последовательного запуска в пределах одного потока CUDA (или последовательно в пределах одного потока ЦП). Ниже приведен пример.

Если вы хотите выполнить параллельное выполнение в потоке, вы можете подумать об использовании CUB, который предлагает процедуры сокращения, которые могут быть вызваны из потокаблока при условии, что ваша карта обеспечивает динамический параллелизм.

Вот пример с Упорный

#include <stdio.h> 

#include <thrust/reduce.h> 
#include <thrust/execution_policy.h> 

/********************/ 
/* 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<N; i++) { 
     h_A[i] = i; 
     sum = sum + h_A[i]; 
    } 
    printf("Host side result = %f\n", sum); 

    float *d_A; gpuErrchk(cudaMalloc((void**)&d_A, N * sizeof(float))); 
    gpuErrchk(cudaMemcpy(d_A, h_A, N * sizeof(float), cudaMemcpyHostToDevice)); 

    test<<<1,1>>>(d_A, N); 

} 

Answer 3

Это обновление моего предыдущего ответа.

Начиная с Thrust 1.8.1, CUDA упорных примитивы могут быть объединены с политикой thrust::device выполнения для запуска параллельно в одном потоке CUDA эксплуатирующего CUDA динамический параллелизм. Ниже приведен пример.

#include <stdio.h> 

#include <thrust/reduce.h> 
#include <thrust/execution_policy.h> 

#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<Nrows; i++) { 
     h_results[i] = 0.f; 
     for (int j=0; j<Ncols; j++) { 
      h_data[i*Ncols+j] = i; 
      h_results[i] = h_results[i] + h_data[i*Ncols+j]; 
     } 
    } 

    TimingGPU timerGPU; 

    float *d_data;   gpuErrchk(cudaMalloc((void**)&d_data,  Nrows * Ncols * sizeof(float))); 
    float *d_results1;  gpuErrchk(cudaMalloc((void**)&d_results1, Nrows   * sizeof(float))); 
    float *d_results2;  gpuErrchk(cudaMalloc((void**)&d_results2, Nrows   * sizeof(float))); 
    gpuErrchk(cudaMemcpy(d_data, h_data, Nrows * Ncols * sizeof(float), cudaMemcpyHostToDevice)); 

    timerGPU.StartCounter(); 
    test1<<<iDivUp(Nrows, BLOCKSIZE_1D), BLOCKSIZE_1D>>>(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<Nrows; i++) { 
     if (h_results1[i] != h_results[i]) { 
      printf("Approach nr. 1; Error at i = %i; h_results1 = %f; h_results = %f", i, h_results1[i], h_results[i]); 
      return 0; 
     } 
    } 

    timerGPU.StartCounter(); 
    test2<<<iDivUp(Nrows, BLOCKSIZE_1D), BLOCKSIZE_1D>>>(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<Nrows; i++) { 
     if (h_results1[i] != h_results[i]) { 
      printf("Approach nr. 2; Error at i = %i; h_results1 = %f; h_results = %f", i, h_results1[i], h_results[i]); 
      return 0; 
     } 
    } 

    printf("Test passed!\n"); 

} 

Приведенный выше пример выполняет сокращения строк матрицы в том же смысле, как Reduce matrix rows with CUDA, но это делается в отличие от поста выше, а именно, с помощью вызова CUDA Упорный примитивы непосредственно из написанного пользователем ядер. Кроме того, приведенный выше пример служит для сравнения производительности тех же операций, когда они выполняются с двумя политиками выполнения, а именно: thrust::seq и thrust::device.Ниже приведены некоторые графики, показывающие разницу в производительности.

Эффективность была оценена на Kepler K20c и на Maxwell GeForce GTX 850M.