2016-06-28 3 views
-2

я использую пример кода, поставляемый Роберт Crovella:GPU работает только один раз

thrust::max_element slow in comparison cublasIsamax - More efficient implementation?

который очень быстро кода сокращения. Я изменил его, чтобы также вернуть индекс max во входном массиве поплавков. Когда я использую его в своем коде, он будет выполняться только один раз. Если я попробую снова вызвать процедуру, она не найдет новое максимальное значение, оно просто вернет предыдущий макс. Есть ли что-то в энергозависимой глобальной памяти, которую использует программа, которую нужно сбросить, прежде чем ее можно будет снова вызвать?

#include <cuda.h> 
#include <cublas_v2.h> 
#include <thrust/extrema.h> 
#include <thrust/device_ptr.h> 
#include <thrust/device_vector.h> 
#include <stdio.h> 
#include <stdlib.h> 

#define DSIZE 4096*4 // nTPB should be a power-of-2 
#define nTPB 512 
#define MAX_KERNEL_BLOCKS 30 
#define MAX_BLOCKS ((DSIZE/nTPB)+1) 
#define MIN(a,b) ((a>b)?b:a) 
#define FLOAT_MIN -1.0f 

#include <helper_functions.h> 
#include <helper_cuda.h> 

// this code has been modified to return the index of the max instead of the actual max value - for my application 
__device__ volatile float blk_vals[MAX_BLOCKS]; 
__device__ volatile int blk_idxs[MAX_BLOCKS]; 
__device__ int blk_num = 0; 

//template <typename T> 
__global__ void max_idx_kernel(const float *data, const int dsize, int *result){ 

    __shared__ volatile float vals[nTPB]; 
    __shared__ volatile int idxs[nTPB]; 
    __shared__ volatile int last_block; 
    int idx = threadIdx.x+blockDim.x*blockIdx.x; 
    last_block = 0; 
    float my_val = FLOAT_MIN; 
    int my_idx = -1; 
    // sweep from global memory 
    while (idx < dsize){ 
    if (data[idx] > my_val) {my_val = data[idx]; my_idx = idx;} 
    idx += blockDim.x*gridDim.x;} 
    // populate shared memory 
    vals[threadIdx.x] = my_val; 
    idxs[threadIdx.x] = my_idx; 
    __syncthreads(); 
    // sweep in shared memory 
    for (int i = (nTPB>>1); i > 0; i>>=1){ 
    if (threadIdx.x < i) 
     if (vals[threadIdx.x] < vals[threadIdx.x + i]) {vals[threadIdx.x] = vals[threadIdx.x+i]; idxs[threadIdx.x] = idxs[threadIdx.x+i]; } 
    __syncthreads();} 
    // perform block-level reduction 
    if (!threadIdx.x){ 
    blk_vals[blockIdx.x] = vals[0]; 
    blk_idxs[blockIdx.x] = idxs[0]; 
    if (atomicAdd(&blk_num, 1) == gridDim.x - 1) // then I am the last block 
     last_block = 1;} 
    __syncthreads(); 
    if (last_block){ 
    idx = threadIdx.x; 
    my_val = FLOAT_MIN; 
    my_idx = -1; 
    while (idx < gridDim.x){ 
     if (blk_vals[idx] > my_val) {my_val = blk_vals[idx]; my_idx = blk_idxs[idx]; } 
     idx += blockDim.x;} 
    // populate shared memory 
    vals[threadIdx.x] = my_val; 
    idxs[threadIdx.x] = my_idx; 
    __syncthreads(); 
    // sweep in shared memory 
    for (int i = (nTPB>>1); i > 0; i>>=1){ 
     if (threadIdx.x < i) 
     if (vals[threadIdx.x] < vals[threadIdx.x + i]) {vals[threadIdx.x] = vals[threadIdx.x+i]; idxs[threadIdx.x] = idxs[threadIdx.x+i]; } 
     __syncthreads();} 
    if (!threadIdx.x) 
     *result = idxs[0]; 
    } 
} 



int main(){ 

    int nrElements = DSIZE; 
    float *d_vector, *h_vector; 

    StopWatchInterface *hTimer = NULL; 
    sdkCreateTimer(&hTimer); 
    double gpuTime; 
    int k; 
    int max_index; 
    int *d_max_index; 
    cudaMalloc(&d_max_index, sizeof(int)); 


    h_vector = new float[DSIZE]; 
    for(k=0; k < 5; k++){ 
    for (int i = 0; i < DSIZE; i++) h_vector[i] = rand()/(float)RAND_MAX; 
     h_vector[10+k] = 10; // create definite max element that changes with each loop iteration 
    cublasHandle_t my_handle; 
    cublasStatus_t my_status = cublasCreate(&my_handle); 
    cudaMalloc(&d_vector, DSIZE*sizeof(float)); 
    cudaMemcpy(d_vector, h_vector, DSIZE*sizeof(float), cudaMemcpyHostToDevice); 

     max_index = 0; 
     sdkResetTimer(&hTimer); 
     sdkStartTimer(&hTimer); 
     //d_vector is a pointer on the device pointing to the beginning of the vector, containing nrElements floats. 
     thrust::device_ptr<float> d_ptr = thrust::device_pointer_cast(d_vector); 
     thrust::device_vector<float>::iterator d_it = thrust::max_element(d_ptr, d_ptr + nrElements); 
     max_index = d_it - (thrust::device_vector<float>::iterator)d_ptr; 
     cudaDeviceSynchronize(); 
     gpuTime = sdkGetTimerValue(&hTimer); 
     std::cout << "loop: " << k << " thrust time: " << gpuTime << " max index: " << max_index << std::endl; 

     max_index = 0; 
     sdkResetTimer(&hTimer); 
     sdkStartTimer(&hTimer); 
     my_status = cublasIsamax(my_handle, DSIZE, d_vector, 1, &max_index); 
     cudaDeviceSynchronize(); 
     gpuTime = sdkGetTimerValue(&hTimer); 
     std::cout << "loop: " << k << " cublas time: " << gpuTime << " max index: " << max_index-1 << std::endl; 

     max_index = 0; 
     sdkResetTimer(&hTimer); 
     sdkStartTimer(&hTimer); 
     max_idx_kernel<<<MIN(MAX_KERNEL_BLOCKS, ((DSIZE+nTPB-1)/nTPB)), nTPB>>>(d_vector, DSIZE, d_max_index); 
     cudaMemcpy(&max_index, d_max_index, sizeof(int), cudaMemcpyDeviceToHost); 
     gpuTime = sdkGetTimerValue(&hTimer); 
     std::cout << "loop: " << k << " idx kern time: " << gpuTime << " max index: " << max_index << std::endl; 
     std::cout << std::endl; 

    } // end for loop on k 

    cudaFree(d_max_index); 
    cudaFree(d_vector); 

    return 0; 
} 
+3

Можете ли вы предоставить полный пример кода, который может воспроизвести проблему? – kangshiyin

ответ

1

Основная проблема в повторном использовании этого кода для нескольких циклов, как есть в этой статической инициализации устройства (глобальной) переменной:

__device__ int blk_num = 0; 

Это нормально, если вы только собираетесь для запуска процедуры один раз. Но если вы намерены повторно использовать его, вам нужно будет повторно инициализировать эту переменную до нуля перед каждым вызовом ядра.

Мы могли бы это исправить, поставив явную инициализацию этого переменный на ноль перед каждым вызовом к ядру редукции:

cudaMemcpyToSymbol(blk_num, &max_index, sizeof(int)); 

(я использую max_index здесь просто потому, что это удобный хост int переменного, только что был установлен на ноль.)

Это единственное изменение, необходимое для получения кода «работа».

Однако введение цикла создало некоторые другие «проблемы», которые я бы указал. Эти 3 строки кода:

cublasHandle_t my_handle; 
cublasStatus_t my_status = cublasCreate(&my_handle); 
cudaMalloc(&d_vector, DSIZE*sizeof(float)); 

не принадлежат внутри для цикла на k. Это эффективно создает утечку памяти и излишне повторно инициализирует библиотеку cublas.

Следующий код имеет те изменения, и, кажется, работает для меня:

$ cat t1183.cu 
#include <cuda.h> 
#include <cublas_v2.h> 
#include <thrust/extrema.h> 
#include <thrust/device_ptr.h> 
#include <thrust/device_vector.h> 
#include <stdio.h> 
#include <stdlib.h> 

#define DSIZE 4096*4 // nTPB should be a power-of-2 
#define nTPB 512 
#define MAX_KERNEL_BLOCKS 30 
#define MAX_BLOCKS ((DSIZE/nTPB)+1) 
#define MIN(a,b) ((a>b)?b:a) 
#define FLOAT_MIN -1.0f 

#include <helper_functions.h> 
#include <helper_cuda.h> 

// this code has been modified to return the index of the max instead of the actual max value - for my application 
__device__ volatile float blk_vals[MAX_BLOCKS]; 
__device__ volatile int blk_idxs[MAX_BLOCKS]; 
__device__ int blk_num; 

//template <typename T> 
__global__ void max_idx_kernel(const float *data, const int dsize, int *result){ 

    __shared__ volatile float vals[nTPB]; 
    __shared__ volatile int idxs[nTPB]; 
    __shared__ volatile int last_block; 
    int idx = threadIdx.x+blockDim.x*blockIdx.x; 
    last_block = 0; 
    float my_val = FLOAT_MIN; 
    int my_idx = -1; 
    // sweep from global memory 
    while (idx < dsize){ 
    if (data[idx] > my_val) {my_val = data[idx]; my_idx = idx;} 
    idx += blockDim.x*gridDim.x;} 
    // populate shared memory 
    vals[threadIdx.x] = my_val; 
    idxs[threadIdx.x] = my_idx; 
    __syncthreads(); 
    // sweep in shared memory 
    for (int i = (nTPB>>1); i > 0; i>>=1){ 
    if (threadIdx.x < i) 
     if (vals[threadIdx.x] < vals[threadIdx.x + i]) {vals[threadIdx.x] = vals[threadIdx.x+i]; idxs[threadIdx.x] = idxs[threadIdx.x+i]; } 
    __syncthreads();} 
    // perform block-level reduction 
    if (!threadIdx.x){ 
    blk_vals[blockIdx.x] = vals[0]; 
    blk_idxs[blockIdx.x] = idxs[0]; 
    if (atomicAdd(&blk_num, 1) == gridDim.x - 1) // then I am the last block 
     last_block = 1;} 
    __syncthreads(); 
    if (last_block){ 
    idx = threadIdx.x; 
    my_val = FLOAT_MIN; 
    my_idx = -1; 
    while (idx < gridDim.x){ 
     if (blk_vals[idx] > my_val) {my_val = blk_vals[idx]; my_idx = blk_idxs[idx]; } 
     idx += blockDim.x;} 
    // populate shared memory 
    vals[threadIdx.x] = my_val; 
    idxs[threadIdx.x] = my_idx; 
    __syncthreads(); 
    // sweep in shared memory 
    for (int i = (nTPB>>1); i > 0; i>>=1){ 
     if (threadIdx.x < i) 
     if (vals[threadIdx.x] < vals[threadIdx.x + i]) {vals[threadIdx.x] = vals[threadIdx.x+i]; idxs[threadIdx.x] = idxs[threadIdx.x+i]; } 
     __syncthreads();} 
    if (!threadIdx.x) 
     *result = idxs[0]; 
    } 
} 



int main(){ 

    int nrElements = DSIZE; 
    float *d_vector, *h_vector; 

    StopWatchInterface *hTimer = NULL; 
    sdkCreateTimer(&hTimer); 
    double gpuTime; 
    int k; 
    int max_index; 
    int *d_max_index; 
    cudaMalloc(&d_max_index, sizeof(int)); 


    h_vector = new float[DSIZE]; 
    cublasHandle_t my_handle; 
    cublasStatus_t my_status = cublasCreate(&my_handle); 
    cudaMalloc(&d_vector, DSIZE*sizeof(float)); 
    for(k=0; k < 5; k++){ 
    for (int i = 0; i < DSIZE; i++) h_vector[i] = rand()/(float)RAND_MAX; 
     h_vector[10+k] = 10; // create definite max element that changes with each loop iteration 
    cudaMemcpy(d_vector, h_vector, DSIZE*sizeof(float), cudaMemcpyHostToDevice); 

    max_index = 0; 
    sdkResetTimer(&hTimer); 
    sdkStartTimer(&hTimer); 
     //d_vector is a pointer on the device pointing to the beginning of the vector, containing nrElements floats. 
    thrust::device_ptr<float> d_ptr = thrust::device_pointer_cast(d_vector); 
    thrust::device_vector<float>::iterator d_it = thrust::max_element(d_ptr, d_ptr + nrElements); 
    max_index = d_it - (thrust::device_vector<float>::iterator)d_ptr; 
    cudaDeviceSynchronize(); 
    gpuTime = sdkGetTimerValue(&hTimer); 
    std::cout << "loop: " << k << " thrust time: " << gpuTime << " max index: " << max_index << std::endl; 

    max_index = 0; 
    sdkResetTimer(&hTimer); 
    sdkStartTimer(&hTimer); 
    my_status = cublasIsamax(my_handle, DSIZE, d_vector, 1, &max_index); 
    cudaDeviceSynchronize(); 
    gpuTime = sdkGetTimerValue(&hTimer); 
    std::cout << "loop: " << k << " cublas time: " << gpuTime << " max index: " << max_index-1 << std::endl; 

    max_index = 0; 
    sdkResetTimer(&hTimer); 
    sdkStartTimer(&hTimer); 
    cudaMemcpyToSymbol(blk_num, &max_index, sizeof(int)); 
    max_idx_kernel<<<MIN(MAX_KERNEL_BLOCKS, ((DSIZE+nTPB-1)/nTPB)), nTPB>>>(d_vector, DSIZE, d_max_index); 
    cudaMemcpy(&max_index, d_max_index, sizeof(int), cudaMemcpyDeviceToHost); 
    gpuTime = sdkGetTimerValue(&hTimer); 
    std::cout << "loop: " << k << " idx kern time: " << gpuTime << " max index: " << max_index << std::endl; 
    std::cout << std::endl; 

    } // end for loop on k 

    cudaFree(d_max_index); 
    cudaFree(d_vector); 

    return 0; 
} 
$ nvcc -I/usr/local/cuda/samples/common/inc t1183.cu -o t1183 -lcublas 
$ cuda-memcheck ./t1183 
========= CUDA-MEMCHECK 
loop: 0 thrust time: 2.806 max index: 10 
loop: 0 cublas time: 0.441 max index: 10 
loop: 0 idx kern time: 0.395 max index: 10 

loop: 1 thrust time: 1.298 max index: 11 
loop: 1 cublas time: 0.419 max index: 11 
loop: 1 idx kern time: 0.424 max index: 11 

loop: 2 thrust time: 1.303 max index: 12 
loop: 2 cublas time: 0.43 max index: 12 
loop: 2 idx kern time: 0.419 max index: 12 

loop: 3 thrust time: 1.291 max index: 13 
loop: 3 cublas time: 0.423 max index: 13 
loop: 3 idx kern time: 0.415 max index: 13 

loop: 4 thrust time: 1.299 max index: 14 
loop: 4 cublas time: 0.423 max index: 14 
loop: 4 idx kern time: 0.417 max index: 14 

========= ERROR SUMMARY: 0 errors 
$ 
+0

Благодарим вас за исправление. – brumby

+0

Не накладные расходы при вызове cudaMemcpyToSymbol(), так как это копия с Host на Device? Было бы лучше просто сбросить blk_num до нуля, когда мы достигнем секции if (last_blk) {} кода ядра? – brumby

+0

Да, вы также сможете сделать эту работу. В этом случае вам нужно добавить обратно в статическую инициализацию для 'blk_num' (' = 0; '), которую я удалил в приведенном выше коде. –

Смежные вопросы