cuda数组找出最值

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_
#define _IN_VAL_
#define _OUT_VAL_
#define _IN_ARR_
#define _OUT_ARR_
#define _IO_ARR_
#define _IO_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 T3>
GLB void findARR3maxmin(_IN_ _ARR_ const T3 *data, _IN_ _VAL_ cintp datasize, _OUT_ _ARR_ T3 *dest, bool is_max = true)
{
    MYID;

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

        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)
                {
                    if (is_max)
                    {
                        IFBIGEQ(tsrd[tid.x].x, tsrd[tid.x + a].x);
                        IFBIGEQ(tsrd[tid.x].y, tsrd[tid.x + a].y);
                        IFBIGEQ(tsrd[tid.x].z, tsrd[tid.x + a].z);
                    }
                    else
                    {
                        IFSMLEQ(tsrd[tid.x].x, tsrd[tid.x + a].x);
                        IFSMLEQ(tsrd[tid.x].y, tsrd[tid.x + a].y);
                        IFSMLEQ(tsrd[tid.x].z, tsrd[tid.x + a].z);
                    }
                }
                CDTSYNC;
            }
        }

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

主要方法

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

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

template <typename T3>
void getArray_maxmin(_IN_ _ARR_ const T3 *data, _IN_ _VAL_ cintr size, _OUT_ _ARR_ T3 *dest, bool is_max = true)
{
    CDACCPY(T3, d_arr, size, data);

    CDCCPYV(int, d_arr_size, size);

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

    int srdsize = block.x * sizeof(T3);

    CDANV(T3, d_max, grid.x);

    while (true)
    {
        findARR3maxmin<<<grid, block, srdsize>>>(d_arr, d_arr_size, d_max, is_max);
        BASE_SYNC;

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

    CDCPYDTH(T3, dest, d_max);

    cudaFree(d_arr);
    cudaFree(d_arr_size);
    cudaFree(d_max);
}

核心思想就是使用cuda多线程的优势，将数组分成多段 每一段同时独立的进行排序 将排序的结果组成一个新的数组 再次将新数组 分段独立排序 直到只剩下一个
因为可以每一段分组可以同时进行排序 在速度上会比cpu更有优势一些