Riduzione condizionale in CUDA

Question

Devo sommare circa i valori 100000 memorizzati in un array, ma con condizioni.

C'è un modo per farlo in CUDA per produrre risultati rapidi?

Qualcuno può pubblicare un piccolo codice per farlo?

Answer 1

Penso che, per eseguire la riduzione condizionale, è possibile introdurre direttamente la condizione come moltiplicazione per 0 (falso) o 1 (true) per gli addendi. In altre parole, supponiamo che la condizione che vorresti incontrare sia che gli addendi siano più piccoli di 10.f. In questo caso, prendendo in prestito il primo codice a Optimizing Parallel Reduction in CUDA by M. Harris, poi il sopra significherebbe

__global__ void reduce0(int *g_idata, int *g_odata) { 

    extern __shared__ int sdata[]; 

    // each thread loads one element from global to shared mem 
    unsigned int tid = threadIdx.x; 
    unsigned int i = blockIdx.x*blockDim.x + threadIdx.x; 
    sdata[tid] = g_idata[i]*(g_data[i]<10.f); 
    __syncthreads(); 

    // do reduction in shared mem 
    for(unsigned int s=1; s < blockDim.x; s *= 2) { 
     if (tid % (2*s) == 0) { 
      sdata[tid] += sdata[tid + s]; 
     } 
     __syncthreads(); 
    } 

    // write result for this block to global mem 
    if (tid == 0) g_odata[blockIdx.x] = sdata[0]; 
} 

Se si desidera utilizzare CUDA di spinta per eseguire la riduzione condizionale, è possibile fare lo stesso utilizzando thrust::transform_reduce. In alternativa, è possibile creare una nuova copia di vettore d_b in cui tutti gli elementi di d_a soddisfano il predicato tramite thrust::copy_if e quindi l'applicazione di thrust::reduce su d_b. Non ho controllato quale soluzione offra il meglio. Forse, la seconda soluzione funzionerà meglio su array sparsi. Di seguito è riportato un esempio con un'implementazione di entrambi gli approcci.

#include <thrust/host_vector.h> 
#include <thrust/device_vector.h> 
#include <thrust/reduce.h> 
#include <thrust/count.h> 
#include <thrust/copy.h> 

// --- Operator for the first approach 
struct conditional_operator { 
    __host__ __device__ float operator()(const float a) const { 
    return a*(a<10.f); 
    } 
}; 

// --- Operator for the second approach 
struct is_smaller_than_10 { 
    __host__ __device__ bool operator()(const float a) const { 
     return (a<10.f); 
    } 
}; 

void main(void) 
{ 
    int N = 20; 

    // --- Host side allocation and vector initialization 
    thrust::host_vector<float> h_a(N,1.f); 
    h_a[0] = 20.f; 
    h_a[1] = 20.f; 

    // --- Device side allocation and vector initialization 
    thrust::device_vector<float> d_a(h_a); 

    // --- First approach 
    float sum = thrust::transform_reduce(d_a.begin(), d_a.end(), conditional_operator(), 0.f, thrust::plus<float>()); 
    printf("Result = %f\n",sum); 

    // --- Second approach 
    int N_prime = thrust::count_if(d_a.begin(), d_a.end(), is_smaller_than_10()); 
    thrust::device_vector<float> d_b(N_prime); 
    thrust::copy_if(d_a.begin(), d_a.begin() + N, d_b.begin(), is_smaller_than_10()); 
    sum = thrust::reduce(d_b.begin(), d_b.begin() + N_prime, 0.f); 
    printf("Result = %f\n",sum); 

    getchar(); 

}