zoukankan      html  css  js  c++  java
  • CUDA计算实列:浅析n(n等于1或大于1)维数组计算方法与grid、block、thread之间关系(上篇)

      本文主要针对与我一样的小白想使用CUDA加速方法,而经过Visual Studio编译器对cuda代码反复试验与调整得到以下成果。

    目前本文主要使用一维数组对CUDA的grid与block的组合尝试,已实现CUDA一维数组的计算,将其记录如下。

    而我也将会在下篇继续使用多维数组使用CUDA计算。

    本文可以帮助大家尽快上手CUDA实现简单代码编写,了解如何利用CUDA 线程实现每个元素处理(若为图片则是像素处理)。

    本文结构:1.基本原理;2.一维数组代码;3.实现结果展示。

    1.基本原理

    原理引用博客为https://blog.csdn.net/tiao_god/article/details/107181883

    CUDA会把线程分为Block和Grid:

     

    这里方便介绍我们用的都是二维的结构,其实Block和Grid也可以设置成三维的。

    blockDim.x:Block的x方向的维度,这里是5,即每行5个线程。

    blockDim.y:Block的y方向的维度,这里是3,即每列3个线程。

    blockIdx.x:Block在x方向的位置,图中放大的Block是2,即为第2个。

    blockIdx.y:Block在y向的位置,图中放大的Block是2,即为第2个。

    注意blockIdx中的Idx是表示index的缩写,而不是表示x方向的ID。

    在CUDA里计算线程索引一般都是:

    const int X = blockIdx.x * blockDim.x + threadIdx.x;
    const int Y = blockIdx.y * blockDim.y + threadIdx.y;
    对应图中放大的区域的Thread(3,1):

    计算式:X = 2*5+3 Y = 2*3+1

    2.一维数组代码

    // 一个维度的计算方法
    __global__ void one_vector_add(float* a_device, float* b_device, float* c_device) {
        
        int tid = threadIdx.y+blockDim.x*threadIdx.x;
        
        c_device[tid] = a_device[tid] + b_device[tid];
    
    }
    
    
    // 一维数组相加
    
    void OneDim_add() {
    
        const int length = 4;                                      // 数组长度为16
        float a[length], b[length], c[length];                      // host中的数组
    
        // 变量初始化赋值
        for (int i = 0; i < length; i++) {
            a[i] = i + 1;
            b[i] = a[i] * 10;
    
        }
    
        // 变量初始化打印
        cout << "a=" << '	';
        for (int i = 0; i < length; i++) { cout << a[i] << '	'; }
        cout << '
    ' << "b=" << '	';
        for (int i = 0; i < length; i++) { cout << b[i] << '	'; }
    
    
        //构造cuda指针变量并分配内存
        float *a_device, *b_device, *c_device;                      // device中的数组
        cudaSetDevice(0);  // 设置cuda设备
        cudaMalloc((void**)&a_device, length * sizeof(float));
        cudaMalloc((void**)&b_device, length * sizeof(float));
        cudaMalloc((void**)&c_device, length * sizeof(float));
        // 将host数组的值拷贝给device数组
        cudaMemcpy(a_device, a, length * sizeof(float), cudaMemcpyHostToDevice);    
        cudaMemcpy(b_device, b, length * sizeof(float), cudaMemcpyHostToDevice);
        cout << '
    ' << "length:" << length << "	分配大小:" << length * sizeof(float) << endl; 
        // 配置grid,block参数
        dim3 grid(1, 1, 1), block(1,length , 1);                    // 设置grid block 参数
        // 重点:调用kernel函数
        
        one_vector_add <<<grid, block >>> (a_device, b_device, c_device);           // 启动kernel
        
        // 将gpu上得到结果值返回host变量                                                                                  
        cudaMemcpy(c, c_device, length * sizeof(float), cudaMemcpyDeviceToHost);    // 将结果拷贝到host
        cout << '
    ' << "c=" << '	';
        for (int i = 0; i < length; i++) { cout << c[i] << '	'; }  //打印cuda计算的值
    
    
    
    
    }
    
    
    
    
    
    //block-thread 1D-3D
    __global__ void TwoBlock1Thread2Way1(float *c, const float *a, const float *b)
    {
        int i = threadIdx.x+threadIdx.y*blockDim.x + threadIdx.z*blockDim.x*blockDim.y; //3d 也适用2D
        c[i] = b[i] + a[i];
    }
    
    
    
    //block-thread 2D-3D
    __global__ void TwoBlock2Thread3Way1(float *c, const float *a, const float *b)
    {
        int thread2d = threadIdx.x + threadIdx.y*blockDim.x;
        int block2d = blockIdx.x+blockIdx.y*gridDim.x; // block是2D形式,但同样适用block为1D形式(blockIdx.x)
        int i = thread2d + (blockDim.x*blockDim.y)*block2d;
        
        c[i] = b[i] + a[i];
    }
    
    
    
    
    
    void TwoDimway1_add() {
        //将二维矩阵化为一维数组,使用一维数组进行block分配
        const int row = 4;
        const int col = 6;
        const int length = row*col;                                      // 数组长度为16
        float a[length], b[length], c[length];                      // host中的数组
    
        // 变量初始化赋值
        for (int i = 0; i < length; i++) {
            a[i] = i + 1;
            b[i] = a[i] * 100;
        }
    
        // 变量初始化打印
        cout << "变量a=" <<endl;
        for (int i = 0; i < row; i++) { 
            for (int i = 0; i < col; i++) {
            cout<< '	' << a[i] ; 
            }
            cout << endl;
        }
        
        cout << "变量b=" << endl;
        for (int i = 0; i < row; i++) {
            for (int i = 0; i < col; i++) {
                cout << '	' << b[i];
            }
            cout << endl;
        }
    
    
        //构造cuda指针变量并分配内存
        float *a_device, *b_device, *c_device;                      // device中的数组
        cudaSetDevice(0);  // 设置cuda设备
        cudaMalloc((void**)&a_device, length * sizeof(float));
        cudaMalloc((void**)&b_device, length * sizeof(float));
        cudaMalloc((void**)&c_device, length * sizeof(float));
        // 将host数组的值拷贝给device数组
        cudaMemcpy(a_device, a, length * sizeof(float), cudaMemcpyHostToDevice);
        cudaMemcpy(b_device, b, length * sizeof(float), cudaMemcpyHostToDevice);
        
        // 配置grid,block参数
        //dim3 grid(1, 1, 1), block(row, col/2, col / 2);                    // 设置block为3D参数
        //TwoBlock1Thread2Way1 <<<grid, block >>> (c_device,a_device, b_device );  // 重点:调用kernel函数
        //dim3 grid(1, 1, 1), block(row, col , 1);                    // 设置block为2D参数                                           
        //TwoBlock1Thread2Way1 << <grid, block >> > (c_device, a_device, b_device);// 重点:调用kernel函数
        
        dim3 grid(2, 1, 1), block(row/2, col, 1);                    // 设置block为2D参数                                           
        TwoBlock2Thread3Way1 << <grid, block >> > (c_device, a_device, b_device);// 重点:调用kernel函数
    
        // 将gpu上得到结果值返回host变量                                                                                  
        cudaMemcpy(c, c_device, length * sizeof(float), cudaMemcpyDeviceToHost);    // 将结果拷贝到host
        
        
        cout << "计算结果 c=" << endl;
        for (int i = 0; i < row; i++) {
            for (int i = 0; i < col; i++) {
                cout << '	' << c[i];
            }
            cout << endl;
        }
        
    
    
    }
    
    
    
    
    int main() {
        
        cout << "
    一维数组cuda计算结果:
    " << endl;
        OneDim_add();
        //TwoDim_add();
        
        cout << "
    
    二维数组一维方式cuda计算结果:
    
    " << endl;
        TwoDimway1_add();
    
    
    
        return 0;
    }

    3.实现结果展示

    采用加法运算,将其变量a与变量b位置元素一一对应相加结果:

    处理算法通用的辅助的code,如读取txt文件,读取xml文件,将xml文件转换成txt文件,读取json文件等
  • 相关阅读:
    2-SAT 问题及图论解法
    割点、强连通分量
    欧拉路、欧拉回路
    DFS 树
    【题解】[AGC 034 F] RNG and XOR【异或卷积 FWT】
    【题解】[UOJ 62] UR #5 怎样跑得更快【莫比乌斯反演】
    【题解】LOJ #2085 / 洛谷 P1587「NOI2016」循环之美【莫比乌斯反演】
    【题解】LOJ #6052 「雅礼集训 2017 Day11」DIV【莫比乌斯反演】
    【题解】[HDU 5382] GCD?LCM【莫比乌斯反演 差分 线性筛】
    【题解】[51Nod 1847] 奇怪的数学题【min_25筛 杜教筛 莫比乌斯反演】
  • 原文地址:https://www.cnblogs.com/tangjunjun/p/14969161.html
Copyright © 2011-2022 走看看