CUDA On Arm Platfrom ---Day02简要摘录

最新推荐文章于 2024-08-27 13:00:46 发布

NVDEV-Ruan

最新推荐文章于 2024-08-27 13:00:46 发布

阅读量858

点赞数 9

分类专栏： CUDA_On_Arm_2022夏令营文章标签： arm 嵌入式硬件人工智能

本文链接：https://blog.csdn.net/m0_46197553/article/details/125624761

版权

CUDA_On_Arm_2022夏令营专栏收录该内容

6 篇文章 8 订阅

订阅专栏

感谢伟大的NV开发者社区，感谢参加本次活动并细心指导我们的各位大犇

课程内容

一、GPU存储单元

关于配有CUDA的GPU详细信息可以通过编译并运行CUDA Samples中的deviceQuery.exe程序查看

在这里插入图片描述

存储单元	位置	cache	访问速度（时钟周期）	权限	作用域
Register	on chip	N/A	0.19	R/W	thread
Local Memory	off chip	无	203	R/W	thread
Shared memory	on chip	N/A	47	R/W	block
Constant memory	off chip	有	110	R	grid
Global memory	off chip	有	218	R/W	grid
Texture Memory	off chip	有	115	R	grid

1、Register（寄存器）

在这里插入图片描述
寄存器是速度最快的存储单元，位于GPU芯片的SM上，用于存储局部变量。每个SM上有成千上万（65536）的32位寄存器，当kernel函数启动后，这些寄存器被分配给指定的线程来使用。

2、Local Memory

在这里插入图片描述

Local Memory本身在硬件中没有特定的存储单元，而是从Global Memory虚拟出来的地址空间。Local Memory是为寄存器无法满足存储需求的情况而设计的，主要是用于存放单线程的大型数组和变量。Local Memory是线程私有的，线程之间是不可见的。由于GPU硬件单位没有Local Memory的存储单元，所以，针对它的访问是比较慢的。从上面的表格中，也可以看到跟Global Memory的访问速度是接近的。

3、Shared Memory

在这里插入图片描述
Shared Memory位于GPU芯片上，访问延迟仅次于寄存器。Shared Memory是可以被一个Block中的所有Thread来进行访问的，可以实现Block内的线程间的低开销通信。

4、Constant Memory

在这里插入图片描述

Constant Memory类似于Local Memory，也是没有特定的存储单元的，只是Global Memory的虚拟地址。因为它是只读的，所以简化了缓存管理，硬件无需管理复杂的回写策略。Constant Memory启动的条件是同一个warp所有的线程同时访问同样的常量数据。

5、Global Memory

在这里插入图片描述

在这里插入图片描述
Global Memory是GPU中最大的存储单元，Host memory与GPU之间的数据交互均会通过Global Memory进行保存，它也是读取速度最慢的组件。GPU中所有计算单元均可以访问该存储单元。

6、Texture Memory

在这里插入图片描述

Texture Memory是GPU的重要特性之一，也是GPU编程优化的关键。Texture Memory实际上也是Global Memory的一部分，但是它有自己专用的只读cache。这个cache在浮点运算很有用，Texture Memory是针对2D空间局部性的优化策略，所以thread要获取2D数据就可以使用texture Memory来达到很高的性能。从读取性能的角度跟Constant Memory类似。

Texture Memory 、Constant Memory 、Local Memory 都是 Global Memory的一部分

二、GPU内存的分配与释放

1、cudaMalloc(void** devPtr , size_t s)

在这里插入图片描述
作用是在GPU上分配内存，分配线性大小的内存，devPtr是返回一个指向已经分配内存的指针，也就是CPU的devPtr所在的内存单元存的是GPU分配的显存首地址。

	cudaMalloc((void**)&d_a, sizeof(int) * m * n);//在设备端为d_a分配大小为m*n个int的大小
    cudaMalloc((void**)&d_b, sizeof(int) * n * k);
    cudaMalloc((void**)&d_c, sizeof(int) * m * k);

2、cudaMemcpy(void dst , const void src , size_t count , cudaMemcpyKind kind)

在这里插入图片描述

用于在主机（Host）和设备（Device）之间往返的传递数据

该函数是同步执行函数，在未完成数据的转移操作之前会锁死并一直占有CPU进程的控制权，所以不用再添加cudaDeviceSynchronize()函数

①Host To Device

	cudaMemcpy(d_a, h_a, sizeof(int) * m * n, cudaMemcpyHostToDevice);//将主机端的h_a拷贝到设备端的d_a
    cudaMemcpy(d_b, h_b, sizeof(int) * n * k, cudaMemcpyHostToDevice);

②Device To Host

	cudaMemcpy(h_c, d_c, sizeof(int) * m * k, cudaMemcpyDeviceToHost);//将设备端的d_c拷贝到主机端的h_c

3、cudaFree(void* devPtr)

在这里插入图片描述
接收void* devptr指针。释放指针指向的显存。如果输入参数时0，不会进行释放操作。

注意F的大写

三、实战环节（GPU实现矩阵乘）

在这里插入图片描述

1、CPU端的实现

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;
            for (int h = 0; h < n; ++h)
            {
                tmp += h_a[i * n + h] * h_b[h * k + j];
            }
            h_result[i * k + j] = tmp;
        }
    }
}

2、GPU端的实现

在这里插入图片描述

有了CPU的代码，思考如何修改成为GPU代码：替换CPU代码中的外两层循环，通过线程的index(索引)计算出i和j,即代码中的row和col，有了这两个值就可以确定当前线程正在计算结果矩阵中的哪个数值，根据矩阵乘法即可算出当前值

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

掌握了通过GPU进行矩阵乘计算后，就初步了解了2D计算框架，通过修改即可完成各种类似任务，如通过CUDA进行二维图像的灰度处理

__global__ void im2gray(uchar3 *in, unsigned char *out, int height, int width)
{
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;
 
    if (x < width && y < height)
    {
        uchar3 rgb = in[y * width + x];
 
        out[y * width + x] = 0.30f * rgb.x + 0.59f * rgb.y + 0.11f * rgb.z;
    }
}

3、GPU函数的调用

m * n的矩阵与n * k的矩阵相乘之后得到的是一个k * m的矩阵，所以执行设置要确保包含k与m，采用向上取整的方法

	dim3 GridDim((k - 1 + BLOCK_SIZE) / BLOCK_SIZE, (m - 1 + BLOCK_SIZE) / BLOCK_SIZE, 1);
    gpu_matrix_mult << <GridDim, { BLOCK_SIZE, BLOCK_SIZE, 1 } >> > (d_a, d_b, d_c, m, n, k);

Q&A 与自我反思

1、block_size 大小的的设定

在这里插入图片描述
从上图可知一个block中的最大线程数为1024，设定时要保证三个方向上之积不超过1024。

注意三个方向上的分配数值的正确性，防止网格大小与预期不相同，导致计算的缺失

2、blockdim 与 warp 的关系

在这里插入图片描述

一个Warp中的线程必然在同一个block中，如果block所含线程数目不是Warp大小的整数倍，那么多出的那些thread所在的Warp中，会剩余一些inactive的thread，也就是说，即使凑不够Warp整数倍的thread，硬件也会为Warp凑足，只不过那些thread是inactive状态，需要注意的是，即使这部分thread是inactive的，也会消耗SM资源。由于warp的大小一般为32，所以block所含的thread的大小一般要设置为32的倍数。

3、野指针的预防（越界）

由于实际中，启动的总线程数总是大于所需线程数，如果不加判断的操作指针所指向的数值，可能会出现越界问题,在运算之前要做好判断

if (row < m && col < k) {  //外层的if判断 有效组织了d_c[]指向域外
        for (int i = 0; i < n; i++) {
            d_c[row * k + col] += d_a[row * n + i] * d_b[col + i * k];
        }
}