CUDA中的条件约简

Question

CUDA中的条件约简

performancecudasynchronizationsum

5

我需要对存储在数组中的约 100000 个值进行求和，但需要满足一些条件。

是否有一种方法可以在CUDA中实现快速结果？

有人能够发布一小段代码来完成此操作吗？

- Roshan

1个回答

网页内容由stack overflow 提供, 点击上面的

可以查看英文原文，
原文链接

- Vitality · Accepted Answer

我认为，要执行条件约简，可以将条件直接引入为加数的 0（假）或 1（真）相乘。换句话说，假设您希望满足的条件是加数小于 10.f。在这种情况下，借用 M. Harris 在《CUDA中优化并行约简》的第一种代码，则上述意味着：

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

如果您希望使用CUDA Thrust执行条件约简，可以使用thrust::transform_reduce完成相同的操作。或者，您可以创建一个新向量d_b，通过thrust::copy_if将所有满足谓词的d_a元素复制到其中，然后在d_b上应用thrust::reduce。我还没有检查哪种解决方案表现最佳。也许，在稀疏数组上第二种解决方案会表现更好。下面是两种方法的实现示例。

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

}