Nvidia Summer Camp Day3&4 个人心得

最新推荐文章于 2024-09-14 17:02:52 发布

George WJC

最新推荐文章于 2024-09-14 17:02:52 发布

阅读量480

点赞数

文章标签： c# c++

本文链接：https://blog.csdn.net/wjc2021/article/details/125737310

版权

本文主要探讨CUDA的存储单元，包括register、shared memory、local memory、constant memory和texture memory。重点讲述了shared memory在矩阵乘法优化中的应用，通过代码示例展示其优势。由于内存访问速度差异，使用shared memory位于GPU芯片上，能提升计算效率。遗憾的是，由于服务器关闭，无法提供实验数据。最后，作者提到颈部疲劳，后续将分享原子操作和CUDA加速库的内容。

摘要由CSDN通过智能技术生成

第三篇总结先从CUDA的储存单元说起

shared memory：用__shared__修饰的变量储存在shared memory中，同一个block中共享一块shared memory，shared memory比较小需要节省着使用。

local memory：无法确定索引是否为常量的数组、消耗太多寄存器资源的结构或数组、如果内核使用了多于可用寄存器的任何变量（寄存器溢出），这几种情况，都会使用local memory。register不够时会使用local memory替代。

constant memory：固定内存空间驻留在设备内存中，并缓存在固定缓存中。constant范围是全局的，针对所有kernel。

texture memory：实际上与global memory在一块，但他有自己专有的只读cache。

memory按照access speed排序如上，只有register与shared memory是放在在GPU chip上。

关于shared memory，今天会着重讲。

按行优先读取一般不会出现conflict，但是按照列有限的话我们就要注意了！既然目的是加速，那就要做到极致。Don't make bank conflict！用代码来展示就是这样：

int x_id = blockdim.x * blockIdx.x + threadIdx.x ;

int y_id = blockdim.y * blockIdx.y + threadIdx.y ;

int index = y_id * col + x_id ;//column

//do like this you can avoid bank conflict : "+ 1"

__shared__ float data[blockSize][blockSize + 1] ;

可以看一下图片帮助你理解。

下面是今天重点，如何使用shared memory优化我们的矩阵乘法。

不多说，直接上代码：

#define BLOCK_SIZE 16

__global__ void gpu_matrix_mult_shared(int *d_a, int *d_b, int *d_result, int n) 
{
    __shared__ int tile_a[BLOCK_SIZE][BLOCK_SIZE];
    __shared__ int tile_b[BLOCK_SIZE][BLOCK_SIZE];

    int row = blockIdx.y * BLOCK_SIZE + threadIdx.y;
    int col = blockIdx.x * BLOCK_SIZE + threadIdx.x;
    int tmp = 0;
    int idx;

    for (int sub = 0; sub < gridDim.x; ++sub) sub < gridDim.x 结合图更好理解
    {
        idx = row * n + sub * BLOCK_SIZE + threadIdx.x;//a的索引
        if(idx >= n*n)
        {
            tile_a[threadIdx.y][threadIdx.x] = 0;
        }
        else
        {
            tile_a[threadIdx.y][threadIdx.x] = d_a[idx];
        }

        idx = (sub * BLOCK_SIZE + threadIdx.y) * n + col;//b的索引
        if(idx >= n*n)
        {
            tile_b[threadIdx.y][threadIdx.x] = 0;
        }  
        else
        {
            tile_b[threadIdx.y][threadIdx.x] = d_b[idx];
        }
        __syncthreads();//同步，不懂的可以看之前的文章
       
            for (int k = 0; k < BLOCK_SIZE; ++k) 
        {
            tmp += tile_a[threadIdx.y][k] * tile_b[k][threadIdx.x];
        }
        __syncthreads();
    }
    if(row < n && col < n)
    {
        d_result[row * n + col] = tmp;
    }
}

好像还没有解释为什么shared memory优化的矩阵乘法速度会更快，其实说白了，就是shared memory是在GPU chip上的，而global memory不是，memory之间的区别导致了速度的差异，可惜实验用的服务器节点已经关闭，没有实验的数据了。

本来准备继续写原子操作与CUDA的加速库，但是作者颈肌劳损了，等我休息休息继续更新。本期内容请各位看代码时多看图，结合着图片更好理解！如有问题欢迎各位指正，谢谢！