【参加CUDA线上训练营】多种存储单元与优化

达达今天学习了嘛

于 2023-02-15 03:43:39 发布

阅读量156

点赞数

分类专栏： CUDA学习文章标签： ai Powered by 金山文档

本文链接：https://blog.csdn.net/weixin_43250007/article/details/129035262

版权

CUDA学习专栏收录该内容

5 篇文章 0 订阅

订阅专栏

一、多种存储单元意义

由于GPU硬件构成拥有多种存储单元，含有多种类型的内存，每种内存分别具有不同的容量和延迟。一般来说延迟低的内存容量小，延迟高的内存容量大。相对于GPU来说，CUDA编程模型向程序员提供更多的控制权，程序员可以根据存储单元的性质选择更好的优化方式以更好的加速场景。因此，对CUDA编程来说，熟悉其内存的分级组织是非常重要的。对于多种存储单元的介绍如下：

a.全局内存（global memory）

全局内存（如在数组相加例子中的指针d_x即是指向全局内存）的含义是核函数中的所有线程都能够访问其中的数据，全局内存由于没有存放在GPU的芯片上，因此具有较高的延迟和较低的访问速度。然而，它的容量是所有设备内存中最大的。其容量基本上就是显存容量。

b.常量内存（Constant Memory）

常量内存仅可读，不可写。常量内存是有常量缓存的全局内存。固定内存空间驻留在设备内存中，并缓存在固定缓存中。数量有限，仅有64KB。它的可见范围和生命周期与全局内存相同。由于有缓存，常量内存的访问速度比全局内存高，但是得到高访问速度的前提是一个线程束中的线程要读取相同的常量内存数据。

c.纹理内存（Texture Memory）

纹理内存驻留在GPU片上内存中，并且只可读不可写。纹理内存专门为哪些在内存呢访问模式中存在大量空间局部性的图形应用程序而设计的。换而言之，当任务中计算的线程“非常接近”时，使用纹理内存进行优化会得到很不错的结果。

d.寄存器（Registers）

寄存器时GPU中最快的存储器，核函数中没有特殊声明的自动变量都是存放在寄存器当中的。当数组的索引是常量（constant）类型且在编译期能被确定的话，就是内置类型，数组也是放在寄存器中。其中，寄存器变量是每个线程私有的，一旦线程执行结束，寄存器变量就会失效。寄存器是稀有资源。

--maxrregcount

可以设置寄存器的大小，并且不同的设备架构的寄存器数量是不同的。

e.共享内存（shared memory）

共享内存和寄存器类似，存在于芯片上，速度与寄存器接近，数量也有限。共享内存对整个线程块可见，生命周期也与整个线程块一致。也就是说，每个线程块拥有一个共享内存变量的副本。共享内存变量的值在不同的线程块中可以不同。一个线程块中的所有线程都可以访问该线程块的共享内存变量副本，但是不能访问其他线程块的共享内存变量副本。共享内存的作用是减少对全局内存的访问，或者改善对全局内存的访问模式。使用__shared__修饰符修饰的变量存放在共享内存之中。使用共享内存进行优化对加速是极有用处的结构如下所示：

f.局部内存（local memory）

局部内存的用法和寄存器几乎一样，核函数中定义的不加任何限定符的变量有可能在寄存器中，也有可能在局部内存中。寄存器中放不下的变量，以及索引值不能再编译时确定的数组，都有可能放在局部内存中。局部内存只是全局内存的一部分。所以，局部内存的延迟也很高。每个线程最多能使用高达512KB的局部内存，但使用过多会降低程序的性能。

g.固定内存（constant Memory）

固定内存空间驻留在设备内存中，并缓存在固定缓存中。当一个wrap中所有线程都从一个Memory地址读取数据时，constant Memory表现会非常好，会触发广播机制。

二、共享存储单元优化矩阵乘法

共享内存的速度仅次于寄存器，数量也有限。共享内存对整个线程块可见，生命周期也与整个线程块一致。所以像使用共享存储单元进行矩阵乘法优化一定要注意使用好有限的资源，避免使得warp中的线程访问同一个bank中的不同地址时发生bank conflict。

在使用GPU调用时使用共享内存对所需要计算的矩阵进行分块操作，以防止共享内存资源占用溢出。

三、使用共享内存优化矩阵乘法代码

#include <stdio.h>
#include <math.h>
#include "error.cuh"

#define BLOCK_SIZE 16

__global__ void gpu_matrix_mult(int *a,int *b, int *c, int m, int n, int k)
{ 
    int row = blockIdx.y * blockDim.y + threadIdx.y; 
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int sum = 0;
    if( col < k && row < m) 
    {
        for(int i = 0; i < n; i++) 
        {
            sum += a[row * n + i] * b[i * k + col];
        }
        c[row * k + col] = sum;
    }
} 

__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) 
    {
        idx = row * n + sub * BLOCK_SIZE + threadIdx.x;
        tile_a[threadIdx.y][threadIdx.x] = row<n && (sub * BLOCK_SIZE + threadIdx.x)<n? d_a[idx]:0;
        idx = (sub * BLOCK_SIZE + threadIdx.y) * n + col;
        tile_b[threadIdx.y][threadIdx.x] = col<n && (sub * BLOCK_SIZE + threadIdx.y)<n? d_b[idx]:0;

        __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;
    }
}

void cpu_matrix_mult(int *h_a, int *h_b, int *h_result, int m, int n, int k) {
    for (int i = 0; i < m; ++i) 
    {
        for (int j = 0; j < k; ++j) 
        {
            int tmp = 0.0;
            for (int h = 0; h < n; ++h) 
            {
                tmp += h_a[i * n + h] * h_b[h * k + j];
            }
            h_result[i * k + j] = tmp;
        }
    }
}

int main(int argc, char const *argv[])
{
    int m=100;
    int n=100;
    int k=100;

    int *h_a, *h_b, *h_c, *h_cc, *h_cs;
    CHECK(cudaMallocHost((void **) &h_a, sizeof(int)*m*n));
    CHECK(cudaMallocHost((void **) &h_b, sizeof(int)*n*k));
    CHECK(cudaMallocHost((void **) &h_c, sizeof(int)*m*k));
    CHECK(cudaMallocHost((void **) &h_cc, sizeof(int)*m*k));
    CHECK(cudaMallocHost((void **) &h_cs, sizeof(int)*m*k));
    
    cudaEvent_t start, stop,stop_share;
    CHECK(cudaEventCreate(&start));
    CHECK(cudaEventCreate(&stop));
    CHECK(cudaEventCreate(&stop_share));


    for (int i = 0; i < m; ++i) {
        for (int j = 0; j < n; ++j) {
            h_a[i * n + j] = 1;
        }
    }

    for (int i = 0; i < n; ++i) {
        for (int j = 0; j < k; ++j) {
            h_b[i * k + j] = 0;
        }
    }

    int *d_a, *d_b, *d_c, *d_c_share;
    CHECK(cudaMalloc((void **) &d_a, sizeof(int)*m*n));
    CHECK(cudaMalloc((void **) &d_b, sizeof(int)*n*k));
    CHECK(cudaMalloc((void **) &d_c, sizeof(int)*m*k));
    CHECK(cudaMalloc((void **) &d_c_share, sizeof(int)*m*k));

    CHECK(cudaEventRecord(start));
    // copy matrix A and B from host to device memory
    CHECK(cudaMemcpy(d_a, h_a, sizeof(int)*m*n, cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(d_b, h_b, sizeof(int)*n*k, cudaMemcpyHostToDevice));

    unsigned int grid_rows = (m + BLOCK_SIZE - 1) / BLOCK_SIZE;
    unsigned int grid_cols = (k + BLOCK_SIZE - 1) / BLOCK_SIZE;
    dim3 dimGrid(grid_cols, grid_rows);
    dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);
    

    gpu_matrix_mult<<<dimGrid, dimBlock>>>(d_a, d_b, d_c, m,n,k);    

    CHECK(cudaMemcpy(h_c, d_c, (sizeof(int)*m*k), cudaMemcpyDeviceToHost));
    //cudaThreadSynchronize();
    CHECK(cudaEventRecord(stop));
    CHECK(cudaEventSynchronize(stop));
    
    gpu_matrix_mult_shared<<<dimGrid, dimBlock>>>(d_a, d_b, d_c_share, n);
    CHECK(cudaMemcpy(h_cs, d_c_share, (sizeof(int)*m*k), cudaMemcpyDeviceToHost));
    
    CHECK(cudaEventRecord(stop_share));
    CHECK(cudaEventSynchronize(stop_share));
    
    float elapsed_time, elapsed_time_share;
    CHECK(cudaEventElapsedTime(&elapsed_time, start, stop));
    CHECK(cudaEventElapsedTime(&elapsed_time_share, stop, stop_share));
    printf("Time_global = %g ms.\n", elapsed_time);
    printf("Time_share = %g ms.\n", elapsed_time_share);

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

    cpu_matrix_mult(h_a, h_b, h_c, m, n, k);

    int ok = 1;
    for (int i = 0; i < m; ++i)
    { 
        for (int j = 0; j < k; ++j)
        {
            if(fabs(h_cs[i*k + j] - h_c[i*k + j])>(1.0e-10))
            {
                printf("hcs: %d hc: %d  ",h_cs[i*k + j], h_c[i*k + j]);
                ok = 0;
            }
        }
    }

    if(ok)
    {
        printf("Pass!!!\n");
    }
    else
    {
        printf("Error!!!\n");
    }
    
    // free memory
    CHECK(cudaFree(d_a));
    CHECK(cudaFree(d_b));
    CHECK(cudaFree(d_c));
    CHECK(cudaFreeHost(h_a));
    CHECK(cudaFreeHost(h_b));
    CHECK(cudaFreeHost(h_c));
    return 0;
}