Bedingte Reduktion in CUDA

Question

Ich muss über 100000 Werte in einem Array gespeichert, aber mit Bedingungen.

Gibt es eine Möglichkeit, in CUDA schnelle Ergebnisse zu erzielen?

Kann jemand einen kleinen Code schreiben, das zu tun?

Answer 1

Ich denke, dass, um Conditional Reduction durchzuführen, Sie direkt die Bedingung als eine Multiplikation von 0 (false) oder 1 (True) zu den Summanden einführen können. Mit anderen Worten, angenommen, die Bedingung, die Sie erfüllen möchten, ist, dass die Summanden kleiner als sind. In diesem Fall wird der erste Code in Optimizing Parallel Reduction in CUDA by M. Harris borgen, dann über die

__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]; 
} 

bedeuten würde, wenn Sie es wünschen CUDA verwenden Thrust bedingte Reduktion durchzuführen, können Sie das gleiche tun, indem thrust::transform_reduce verwenden. Alternativ können Sie einen neuen Vektor d_b Kopieren, dass alle Elemente der d_a erfüllt das Prädikat erstellen, indem thrust::copy_if und dann thrust::reduce auf d_b Anwendung. Ich habe nicht überprüft, welche Lösung am besten funktioniert. Vielleicht wird die zweite Lösung auf Sparse-Arrays besser funktionieren. Im Folgenden finden Sie ein Beispiel mit einer Implementierung beider Ansätze.

#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(); 

}