cuda 的并行规约计算平均数

宏

本文会把一些宏罗列出来 关于创建变量和复制变量的宏请移步这篇博客查看

#define CUDA_ID            \
    uint3 tid = threadIdx; \
    uint3 bid = blockIdx;  \
    uint3 bdim = blockDim; \
    uint3 gdim = gridDim

#define SRD __shared__
#define GLB __global__

#define CDTSYNC __syncthreads()
#define CDDSYNC cudaDeviceSynchronize()

工具方法

计算cuda内核函数需要传入的前两个参数

• 线程格（有多少个线程块）（所有线程块的规模（x,y,z））
• 线程块中的线程数量（每个线程块中的线程规模（x,y,z））

// 并没有使用到yz 即 所有的线程块和线程格为(n,1,1)
void calculateScalar_grid_block(const int size, dim3 &grid, dim3 &block, int device_num)
{
    cudaSetDevice(device_num);
    cudaDeviceProp property;
    cudaGetDeviceProperties(&property, device_num);

    grid.x = grid.y = grid.z = block.x = block.y = block.z = 1;

    int maxThreadsPerBlock = property.maxThreadsPerBlock;
    if (size <= maxThreadsPerBlock)
    {
        block.x = size;
    }
    else
    {
        block.x = maxThreadsPerBlock;
        grid.x = std::ceil(size / maxThreadsPerBlock);
    }
}

调用的cuda内核方法

内核负责把前两个参数所代表的数组利用共享内存规约计算存储到第三个数组对应线程块的下标中。

GLB void kernel_mean(const float3 *data, const int *datasize, float3 *dest)
{
    CUDA_ID;
    const int id = bid.x * bdim.x + tid.x;
    extern SRD int srd[];  // 一个页面的共享内存只能有一个名字
    					 // 因为原页里有一些其他的方法使用了int类型
    					 // 所以此处先声明了int,再进行了类型强转

    float3 *d_srd = (float3 *)srd;

    d_srd[tid.x] = data[id];
    CDTSYNC;  // 所有线程同步 本体看宏

    for (int a = 1; a < bdim.x; a *= 2)
    {
        if (tid.x % (a * 2) == 0)
        {
            d_srd[tid.x].x += d_srd[tid.x + a].x;
            d_srd[tid.x].y += d_srd[tid.x + a].y;
            d_srd[tid.x].z += d_srd[tid.x + a].z;
        }
        CDTSYNC;
    }

    dest[bid.x] = d_srd[0];
}

主要方法

将所有数组放在新建最大线程数的线程块中，然后新建一个线程块数量的数组A，每一次将每一个线程块内全部线程的求和结果放在新建的数组A中，while循环进行迭代，即数组需要进行 ceil(size / maxThreadsPerBlock)次计算

void calc_mean_point2(const float3 *data, const int *datasize, float3 &dest)
{
    const int size = *datasize;

    CDACCPY(float3, d_arr, size, data);

    CDNV(int, d_size);

    dim3 grid;
    dim3 block;
    int tmpsize = size;

    CDACCPY(float3, d_mid_arr, tmpsize, data);

    CDNV(int, d_mid_size);

    bool flag = true;
    while (flag)
    {
        CDCPYHTDV(int, d_mid_size, tmpsize);

        calculateScalar_grid_block(tmpsize, grid, block);

        CDANV(float3, d_tmp_arr, grid.x);
        kernel_mean<<<grid, block>>>(d_mid_arr, d_mid_size, d_tmp_arr);
        CDDSYNC;

        CDACPYDTD(float3, d_mid_arr, d_tmp_arr, grid.x);
        cudaFree(d_tmp_arr);

        if (grid.x == 1)
            break;

        tmpsize = grid.x;
    }
    cudaFree(d_mid_size);

    float3 mean;
    CDCPYDTHV(float3, mean, d_mid_arr);

    dest = {mean.x / size, mean.y / size, mean.z / size};

    cudaFree(d_arr);
    cudaFree(d_size);
    cudaFree(d_mid_arr);
}

文章中出现的代码都是经过调试并且成功的，有问题可以评论区咨询或者给一些有利于代码的建议，本人会虚心学习。