Быстрые параллельные сокращения на Kepler

Question

Я просто начинающий CUDA и пытаюсь использовать Faster Parallel Reductions on Kepler в своей программе, но я не получил результат, ниже функция того, что я делаю, вывод 0, I было бы полезно узнать, в чем моя ошибка?

#ifndef __CUDACC__ 
#define __CUDACC__ 
#endif 

#include <cuda.h> 
#include <cuda_runtime.h> 
#include "device_launch_parameters.h" 
#include <iostream> 
#include <cuda_runtime_api.h> 
#include <device_functions.h> 
#include <stdio.h> 
#include <math.h> 

__inline__ __device__ 
float warpReduceSum(float val) { 
    for (int offset = warpSize/2; offset > 0; offset /= 2) 
    val += __shfl_down(val, offset); 
    return val; 
} 

__inline__ __device__ 
float blockReduceSum(float val) { 

    static __shared__ int shared[32]; // Shared mem for 32 partial sums 
    int lane = threadIdx.x % warpSize; 
    int wid = threadIdx.x/warpSize; 

    val = warpReduceSum(val);  // Each warp performs partial reduction 

    if (lane==0) shared[wid]=val; // Write reduced value to shared memory 

    __syncthreads();    // Wait for all partial reductions 

    //read from shared memory only if that warp existed 
    val = (threadIdx.x < blockDim.x/warpSize) ? shared[lane] : 0; 

    if (wid==0) val = warpReduceSum(val); //Final reduce within first warp 

    return val; 
} 

__global__ void deviceReduceKernel(float *in, float* out, size_t N) 
{ 
    float sum = 0; 
    //reduce multiple elements per thread 
    for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < N; i += blockDim.x * gridDim.x) 
    { 
    sum += in[i]; 
    } 
    sum = blockReduceSum(sum); 
    if (threadIdx.x==0) 
    out[blockIdx.x]=sum; 
} 

int main() 
{ 
    int n = 1000000; 
    float *b = new float[1](); 
    float *d = new float[1](); 
    float *a ; 


    int blocks = (n/512)+1; 
    float *d_intermediate; 

    cudaMalloc((void**)&d_intermediate, n*sizeof(float)); 
    cudaMalloc((void**)&a, n*sizeof(float)); 

    cudaMemset(a, 1, n*sizeof(float)); 

    deviceReduceKernel<<<blocks, 512>>>(a, d_intermediate, n); 
    deviceReduceKernel<<<1, 1024>>>(d_intermediate, &b[0], blocks); 
    cudaMemcpy(d, b, sizeof(float), cudaMemcpyDeviceToHost); 
    cudaFree(d_intermediate); 
    std::cout << d[0]; 
    return 0; 

} 

Answer 1

Существуют различные проблемы с вашим кодом:

1. Каждый раз, когда у вас возникли проблемы с кодом CUDA, вы должны использовать proper cuda error checking и запустить свой код с cuda-memcheck, перед просят другие для Помогите. Даже если вы не понимаете выход ошибки, это будет полезно для других, пытающихся вам помочь. Если вы сделали это с помощью этого кода, вам будут сообщены различные ошибки/проблемы

2. Любой указатель, переданный ядру CUDA, должен быть допустимым указателем устройства CUDA. Ваш b указатель является указателем хост:

   float *b = new float[1](); 
   

   так что вы не можете использовать его здесь:

   deviceReduceKernel<<<1, 1024>>>(d_intermediate, &b[0], blocks); 
                  ^
   

   , так как вы, очевидно, хотите использовать его для хранения одного float количества на устройстве, мы можем легко повторите использование указателя a.

3. По той же причине, что это не имеет смысла:

   cudaMemcpy(d, b, sizeof(float), cudaMemcpyDeviceToHost); 
   

   в этом случае как b и d являются указателями хоста. Это не будет копировать данные с устройства на хост.

4. Это, вероятно, не делать то, что вы думаете:

   cudaMemset(a, 1, n*sizeof(float)); 
   

   Я полагаю, вы думаете, это будет заполнить float массив с количеством 1, но это не будет. cudaMemset, как memset, заполняет байт и принимает байт количество. Если вы используете его для заполнения массива float, вы фактически создаете массив, заполненный 0x01010101. Я не знаю, какое значение это преобразует, когда вы преобразуете бит-шаблон в число float, но оно не даст вам значения float 1. Мы исправим это, заполнив обычный массив хоста контуром, а затем перенося эти данные на устройство, которое должно быть уменьшено.

Вот модифицированный код, который имеет вышеуказанные проблемы решены, и работает правильно для меня:

$ cat t1290.cu 
#include <iostream> 
#include <stdio.h> 
#include <math.h> 

__inline__ __device__ 
float warpReduceSum(float val) { 
    for (int offset = warpSize/2; offset > 0; offset /= 2) 
    val += __shfl_down(val, offset); 
    return val; 
} 

__inline__ __device__ 
float blockReduceSum(float val) { 

    static __shared__ int shared[32]; // Shared mem for 32 partial sums 
    int lane = threadIdx.x % warpSize; 
    int wid = threadIdx.x/warpSize; 

    val = warpReduceSum(val);  // Each warp performs partial reduction 

    if (lane==0) shared[wid]=val; // Write reduced value to shared memory 

    __syncthreads();    // Wait for all partial reductions 

    //read from shared memory only if that warp existed 
    val = (threadIdx.x < blockDim.x/warpSize) ? shared[lane] : 0; 

    if (wid==0) val = warpReduceSum(val); //Final reduce within first warp 

    return val; 
} 

__global__ void deviceReduceKernel(float *in, float* out, size_t N) 
{ 
    float sum = 0; 
    //reduce multiple elements per thread 
    for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < N; i += blockDim.x * gridDim.x) 
    { 
    sum += in[i]; 
    } 
    sum = blockReduceSum(sum); 
    if (threadIdx.x==0) 
    out[blockIdx.x]=sum; 
} 

int main() 
{ 
     int n = 1000000; 
     float b; 
     float *a, *a_host; 
     a_host = new float[n]; 

     int blocks = (n/512)+1; 
     float *d_intermediate; 

     cudaMalloc((void**)&d_intermediate, blocks*sizeof(float)); 
     cudaMalloc((void**)&a, n*sizeof(float)); 
     for (int i = 0; i < n; i++) a_host[i] = 1; 
     cudaMemcpy(a, a_host, n*sizeof(float), cudaMemcpyHostToDevice); 

     deviceReduceKernel<<<blocks, 512>>>(a, d_intermediate, n); 
     deviceReduceKernel<<<1, 1024>>>(d_intermediate, a, blocks); 
     cudaMemcpy(&b, a, sizeof(float), cudaMemcpyDeviceToHost); 
     cudaFree(d_intermediate); 
     std::cout << b << std::endl; 
     return 0; 
} 
$ nvcc -arch=sm_35 -o t1290 t1290.cu 
$ cuda-memcheck ./t1290 
========= CUDA-MEMCHECK 
1e+06 
========= ERROR SUMMARY: 0 errors 
$