CUDA koşullu azaltma

Question

Bir dizide saklanan 100000 değerlerini, ancak koşullarla özetlemem gerekiyor.

Hızlı sonuçlar elde etmek için CUDA'da bunu yapmanın bir yolu var mı?

Bunu yapmak için küçük bir kod gönderen var mı?

Answer 1

Koşullu azaltmayı gerçekleştirmek için, koşulu doğrudan eklere 0 (false) veya 1 (true) ile çarpma olarak tanıtabilirsiniz. Diğer bir deyişle, karşılamak istediğiniz koşulun eklerin 10.f'dan daha küçük olması gerektiğini varsayalım. Eğer CUDA koşullu azalma gerçekleştirmek için Thrust kullanmak isterseniz bu durumda, Optimizing Parallel Reduction in CUDA by M. Harris ilk kod ödünç sonra yukarıdaki

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

anlamına geleceğini, sen thrust::transform_reduce kullanarak aynısını yapabilirsiniz. Alternatif olarak, 'un thrust::copy_if numaralı yüklemeyi karşılayan tüm öğelerin d_b kopyalanmasını venumaralı telefona thrust::reduce uygulanmasını sağlayabilirsiniz. Hangi çözümün en iyi sonucu verdiğini kontrol etmedim. Belki ikinci çözüm seyrek dizilerde daha iyi performans gösterecektir. Aşağıda, her iki yaklaşımın da uygulanmasına dair bir örnek verilmiştir.

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

}