Réductions parallèles plus rapides sur Kepler

Question

Je suis juste un débutant CUDA et essayant d'utiliser Faster Parallel Reductions on Kepler sur mon programme, mais je n'ai pas obtenu le résultat, ci-dessous est une fonction de ce que je fais, la sortie est 0, je serait apprécié de savoir quelle est mon erreur?

#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

Il y a plusieurs problèmes avec votre code:

1. Chaque fois que vous rencontrez des problèmes avec un code CUDA, vous devez utiliser proper cuda error checking et exécuter votre code avec cuda-memcheck, avant demander aux autres Aidez-moi. Même si vous ne comprenez pas le résultat de l'erreur, il sera utile aux autres qui essaient de vous aider. Si vous l'aviez fait avec ce code, vous seriez informé de diverses erreurs/problèmes

2. Tout pointeur passé à un noyau CUDA devrait être un pointeur de périphérique CUDA valide. Votre b pointeur est un pointeur hôte:

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

   vous ne pouvez pas l'utiliser ici:

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

   puisque vous voulez évidemment l'utiliser pour le stockage d'une seule quantité float sur l'appareil, nous pouvons facilement réutilisez le pointeur a pour cela.

3. Pour une raison similaire, ce n'est pas raisonnable:

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

   dans ce cas, les deux b et d sont des pointeurs hôtes. Cela ne copiera pas les données de l'appareil vers l'hôte.

4. Cela ne fait probablement pas ce que vous pensez:

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

   Je suppose que vous pensez que cela va remplir un tableau float avec la quantité 1, mais il ne sera pas. cudaMemset, comme memset, remplit octets et prend un octet quantité. Si vous l'utilisez pour remplir un tableau float, vous créez effectivement un tableau rempli de 0x01010101. Je ne sais pas quelle valeur se traduit lorsque vous convertissez le modèle de bits en une quantité float, mais cela ne vous donnera pas une valeur de float de 1. Nous allons résoudre ce problème en remplissant un tableau hôte ordinaire avec une boucle, puis transférer ces données à l'appareil à réduire.

est ici un code modifié qui a les questions ci-dessus adressées, et fonctionne correctement pour moi:

$ 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 
$