cuda数组任意长度归约求和

最新推荐文章于 2024-04-14 13:36:16 发布

变羊法师

最新推荐文章于 2024-04-14 13:36:16 发布

阅读量477

点赞数 13

分类专栏： CUDA随笔文章标签：算法笔记

本文链接：https://blog.csdn.net/zerusz/article/details/135407879

版权

CUDA随笔专栏收录该内容

5 篇文章 0 订阅

订阅专栏

cuda数组任意长度求和

既是记录自己书写的代码也是方便以后用到代码的时候可以直接复制粘贴，有什么问题或者好的建议可以私信我我会耐心回复

宏

本文出现的一些常用宏定义

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

#define MYID \
    CUDA_ID; \
    const int id = bid.x * bdim.x + tid.x

#define MYSRD(T) extern SRD int srd[];  T *tsrd = (T *)srd

#define CDTSYNC __syncthreads()

#define IFBIGEQ(eq, big) if(eq < big)  eq = big
#define IFSMLEQ(eq, small) if(eq > small) eq = small

#define _IN_OUT_
#define _IN_
#define _OUT_

#define _ARR_
#define _VAL_

工具方法

默认线程块和线程格的y和z都是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 / flt(maxThreadsPerBlock));
    }
}

调用的cuda的内核方法

方法命名还有待改善

template <typename T>
GLB void summation_reduction(_IN_ARR_ const T *data, _IN_VAL_ const int *datasize, _OUT_ARR_ T *dest)
{
    MYID;

    if (id < *datasize)
    {
        MYSRD(T);

        tsrd[tid.x] = data[id];
        CDTSYNC;

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

        if(tid.x == 0)
        
            dest[bid.x] = tsrd[0];
        
    }
}

主要方法

找max和找min同理，只需要把内部的大小写逻辑修改一下就可以了

代码还不算完善可以再增加一个变量来确定寻找的是最大值还是最小值

template <typename T>
void summation(_IN_ARR_ const T *data, _IN_VAL_ cint size, _OUT_VAL_ T *dest)
{
    CDACCPY(T, d_arr, size, data);

    CDCCPYV(int, d_arr_size, size);

    dim3 grid, block;
    calculateScalar_grid_block(size, grid, block);

    int srdsize = block.x * sizeof(T);
    
    CDANV(T, d_sum, grid.x);

    while (true)
    {
        summation_reduction<<<grid, block, srdsize>>>(d_arr, d_arr_size, d_sum);
        BASE_SYNC;

        if (grid.x == 1)
        {
            break;
        }
        else
        {
            CDACPYDTD(T, d_arr, d_sum, grid.x);
            CDCPYHTDV(int, d_arr_size, grid.x);
            calculateScalar_grid_block(grid.x, grid, block);
            srdsize = block.x * sizeof(T);
        }
    }

    CDCPYDTH(T, dest, d_sum);

    cudaFree(d_arr);
    cudaFree(d_arr_size);
    cudaFree(d_sum);
}