CUDA入门之利用GPU寻找一组数据中最大的k个元素（二）

还能再吃一碗w

已于 2022-07-09 21:21:33 修改

阅读量708

点赞数

分类专栏： CUDA入门

于 2022-07-09 19:56:02 首次发布

本文链接：https://blog.csdn.net/Not_answer/article/details/125697601

版权

算法矩阵开发语言

CUDA入门专栏收录该内容

6 篇文章 0 订阅

订阅专栏

参加CUDA On Arm Platform 线上夏令营学习笔记（五）

代码部分
讲解部分
实验结果比较
总结反思

CUDA入门之利用GPU寻找一组数据中最大的k个元素（一）
的中介绍了如何利用GPU寻找一组数据中最大的k个元素的基本方法，并在文章的最后，作者留下了一个思考：上文中采用two-pass的方法进行求解，但是第二次grid等于1时，申请的线程数其实是远小于输入数据的数量，这样在ken老师提供的insert_value函数会在最初赋值的时候调用很多次，如果多调用几次核函数，使最后一次就final的时候申请的线程数大于等于输入的数据，那么执行速度会更快吗？
本文中采取了3-pass的方式重新执行了代码，对比2-pass结果发现有一定加速效果。

代码部分

#include <stdio.h>
#include <stdlib.h>
#include <time.h>   
#include "error.cuh"

#define BLOCK_SIZE 256
#define N 2000000
#define GRID_SIZE  ((N + BLOCK_SIZE - 1) / BLOCK_SIZE) 
#define topk 20


__managed__ int source_array[N];
//申请的中间变量
__managed__ int _1pass_results[topk * GRID_SIZE];
__managed__ int _2pass_results[ topk * ((topk * GRID_SIZE + BLOCK_SIZE - 1) / BLOCK_SIZE)];

__managed__ int final_results[topk];

__device__ __host__ void insert_value(int* array, int k, int data)
{
    for (int i = 0; i < k; i++)
    {
        if (array[i] == data)
        {
            return;
        }
    }
    if (data < array[k - 1])
        return;
    for (int i = k - 2; i >= 0; i--)
    {
        if (data > array[i])
            array[i + 1] = array[i];
        else {
            array[i + 1] = data;
            return;
        }
    }
    array[0] = data;
}

__global__ void top_k(int* input, int count, int* output, int k)
{
    __shared__ int topk_per_thread[BLOCK_SIZE * topk];

    int top[topk] = {INT_MIN};
    
    for(int idx = threadIdx.x + blockDim.x * blockIdx.x; idx < count; idx += gridDim.x * blockDim.x){
        insert_value(top, k, input[idx]);
    }
    
    for(int i = 0; i < topk; ++i){
        topk_per_thread[k * threadIdx.x + i] = top[i];
    }
    __syncthreads();
    
    for(int length = BLOCK_SIZE / 2; length >= 1; length /= 2){
        
        if(threadIdx.x < length){
            for(int i = 0; i < topk; ++i){
                insert_value(top, k, topk_per_thread[k * (threadIdx.x + length) + i]);
            }
        }
        __syncthreads();
        
        if(threadIdx.x < length){
            for(int i = 0; i < topk; ++i){
                topk_per_thread[topk * threadIdx.x + i] = top[i];
            }        
        }
        __syncthreads();
    }
   
    if(blockIdx.x * blockDim.x < count){
        for(int i = 0; i < topk; ++i){
            output[topk * blockIdx.x + i] = topk_per_thread[i];
        }
    }
}

void cpu_result_topk(int* input, int count, int* output)
{
    /*for (int i = 0; i < topk; i++)
    {
        output[i] = INT_MIN;
    }*/
    for (int i = 0; i < count; i++)
    {
        insert_value(output, topk, input[i]);

    }
}

void _init(int* ptr, int count)
{
    srand((unsigned)time(NULL));
    for (int i = 0; i < count; i++) ptr[i] = rand();
}

int main(int argc, char const* argv[])
{
    int cpu_result[topk] = { 0 };
    cudaEvent_t start, stop,stop_test;
    CHECK(cudaEventCreate(&start));
    CHECK(cudaEventCreate(&stop));
    CHECK(cudaEventCreate(&stop_test));
    //Fill input data buffer
    _init(source_array, N);


    printf("\n***********GPU RUN**************\n");
    CHECK(cudaEventRecord(start));
    top_k << <GRID_SIZE, BLOCK_SIZE >> > (source_array, N, _1pass_results, topk);
    CHECK(cudaGetLastError());
    top_k << <1, BLOCK_SIZE >> > (_1pass_results, topk * GRID_SIZE, final_results, topk);
    CHECK(cudaGetLastError());
    CHECK(cudaDeviceSynchronize());

    CHECK(cudaEventRecord(stop));
    CHECK(cudaEventSynchronize(stop));
    // 3-pass运行结果
    top_k << <GRID_SIZE, BLOCK_SIZE >> > (source_array, N, _1pass_results, topk);
    CHECK(cudaGetLastError());
    top_k << <((topk * GRID_SIZE + BLOCK_SIZE - 1) / BLOCK_SIZE), BLOCK_SIZE >> > (_1pass_results, topk * GRID_SIZE, _2pass_results, topk);
    CHECK(cudaGetLastError());
    top_k << <1, BLOCK_SIZE >> > (_2pass_results, topk * ((topk * GRID_SIZE + BLOCK_SIZE - 1) / BLOCK_SIZE), final_results, topk);
    CHECK(cudaDeviceSynchronize()); 

    CHECK(cudaEventRecord(stop_test));
    CHECK(cudaEventSynchronize(stop_test));

    float elapsed_time,elapsed_time_mul;
    CHECK(cudaEventElapsedTime(&elapsed_time, start, stop));
    CHECK(cudaEventElapsedTime(&elapsed_time_mul, stop, stop_test));
    printf("Time = %g ms.\n", elapsed_time);
    printf("Time_mul = %g ms.\n", elapsed_time_mul);

    CHECK(cudaEventDestroy(start));
    CHECK(cudaEventDestroy(stop));

    cpu_result_topk(source_array, N, cpu_result);

    int ok = 1;
    for (int i = 0; i < topk; ++i)
    {
        printf("cpu top%d: %d; gpu top%d: %d \n", i + 1, cpu_result[i], i + 1, final_results[i]);
        if (fabs(cpu_result[i] - final_results[i]) > (1.0e-10))
        {

            ok = 0;
        }
    }

    if (ok)
    {
        printf("Pass!!!\n");
    }
    else
    {
        printf("Error!!!\n");
    }
    return 0;
}

讲解部分

本次的程序与《CUDA入门之利用GPU寻找一组数据中最大的k个元素（一）》中的程序在核函数的调用上并无差别，差别在于主函数中kernel调用了3次而非两次，虽然此时依旧无法保证在最后一次核函数调用时申请的线程数大于等于输入数据的元素个数，但是在中间一次的核函数调用中，将输入数据的个数减少了一部分，这样可以使得最后一次调用核函数时，输入数据的元素个数小于2-pass时输入数据的元素个数。

//2-pass运行结果
top_k << <GRID_SIZE, BLOCK_SIZE >> > (source_array, N, _1pass_results, topk);
CHECK(cudaGetLastError());
top_k << <1, BLOCK_SIZE >> > (_1pass_results, topk * GRID_SIZE, final_results, topk);
CHECK(cudaGetLastError());
CHECK(cudaDeviceSynchronize());

// 3-pass运行结果
top_k << <GRID_SIZE, BLOCK_SIZE >> > (source_array, N, _1pass_results, topk);
CHECK(cudaGetLastError());
top_k << <((topk * GRID_SIZE + BLOCK_SIZE - 1) / BLOCK_SIZE), BLOCK_SIZE >> > (_1pass_results, topk * GRID_SIZE, _2pass_results, topk);
CHECK(cudaGetLastError());
top_k << <1, BLOCK_SIZE >> > (_2pass_results, topk * ((topk * GRID_SIZE + BLOCK_SIZE - 1) / BLOCK_SIZE), final_results, topk);
CHECK(cudaDeviceSynchronize());

实验结果比较

在这里插入图片描述
通过实验对比发现，3-pass的运行时间仅比2-pass稍短，在咨询cuda夏令营的老师和气氛组得到如下解答：
“如果你原始问题规模非常大。例如有16M组中间结果，那么直接将16M组中间结果只用1个block处理就比较吃亏了。此时再上一组很多个blocks的中间kernel处理一次就比较好了。当然，考虑到NX只有6个SM，最坏的情况下只能中间的额外kernel，同时驻留6个blocks（和blocks的资源使用，线程规模有关，不一定）。在NX上的最坏情况中间kernel只能提供6X的并行度。这里可能会影响和直接上最终1个block的时间对比。在台式卡上 + 中间结果很多组的时候，应当3次效果比较好。具体哪种在目标问题上更好，不能直接得到答案。但是你总是可以去实验。如果是在Nano上，情况会更加糟糕。根据block使用的资源不同，和线程规模不同。你中间执行的那个kernel，最坏情况只能有1个block，上在只有1个SM的nano身上。此时最坏的中间的真正的并行度只有1X。和直接上最终的kernel没太大区别了。”
因此可能是几个因素之间的综合叠加使得速度并未得到显著的提升，相反如果输入数据的元素个数少一些，可能会使得调用核函数的时间大于损失的时间，因此需要综合考量一下。

总结反思

开始尝试优化的时候，尝试是在核函数外套一个循环，直至上一个输出向量的大小等于或小于batchsize，但是因为作者使用的是nano，中间向量的设定早就设定好了，如果在主函数中重新设定中间向量的话，还需要进行 cudaMemcpy();，在运行时间上耽误了更久，所以初步实验仅仅是测试了3-pass与2-pass之间的运行时间差异，所有的中间向量都是作者事先设定好的。

__managed__ int source_array[N];
//申请的中间变量
__managed__ int _1pass_results[topk * GRID_SIZE];
__managed__ int _2pass_results[ topk * ((topk * GRID_SIZE + BLOCK_SIZE - 1) / BLOCK_SIZE)];

__managed__ int final_results[topk];