NVIDIA CUDA 高性能计算笔记（一）cuda编程简介及矩阵赋值案例

NVIDIA CUDA 高性能计算笔记（一）

​ CUDA （Compute Unified Device Architecture）是NIVIDIA 推出的通用并行计算平台，支持C，C++，Python等语言，实现CPU和GPU协同计算。其架构采用Grid-Blocks-Threads线程层次结构和SIMT并行模式，在给出CUDA的编程实例之前，需要给出模型的基础知识做个简单的介绍。

1.1CUDA编程模型简介

​ CUDA编程模型是一个异构模型，需要GPU和CPU协同工作。在CUDA架构中，我们用host端指代CPU及其内存的，用device指代GPU及其内存。CUDA程序中即包含Host程序，又包含device程序，它们分别在CPU与GPU上运行。同时，host与device之间进行通信，这样它们之间可以进行数据拷贝。典型的CUDA程序的执行的程序的流程为：

  1. 分配host内存，并进行数据初始化；
  2. 分配device内存（显存、共享内存），并从host端将数据拷贝到device端；
  3. 调用CUDA的核函数在device函数上完成指定的运算；
  4. 将device上的运算结果拷贝到host上；
  5. 释放device和host上分配的内存。

​ 由于CUDA编程模型实际上是异构编程模型，所以需要区分host和device上的代码，在CUDA中是通过函数类型限定词区别开host和device上的函数，主要的三个函数类型限定词如下：

  • __global__: 在device端上执行，从host中调用（一些特定的GPU也可以从device上调用），返回类型必须为 void , 不支持可变参数，不能成为类成员函数。注意__global__ 定义的kernel是异步的，这意味着host端不会等待kernel执行完就执行下一步；
  • __device__: 在device端上执行，但仅可以从device中调用，不可以和 __global__ 同时用；
  • __host__: 在host上执行，仅可以从host中调用，一般省略不写，不可以和 __global__同时用，但可以和 __device__，此时函数会在device和host都编译。

​ 上面的流程中最重要的一个过程是调用CUDA的核函数来执行并行计算，kernel是CUDA中的一个重要的概念，kernel是在device上线程中并行执行的函数，在调用时需要用 <<<grid,block>>> 来指定kernel要执行的线程数量，在CUDA中，每个线程都要执行核函数，并且每个线程会分配一个唯一的\(thread\space ID\) ,这个\(ID\) 值可以通过核函数的内置变量 thread Idx 来获得。

​ 要深刻理解\(kernel\)，必须要对\(kernel\) 的线程层次结构有一个清晰的认识。首先，\(GPU\)上很多并形化的轻量级线程。\(kernel\) 在device上执行时实际上是启动很多线程，一个\(kernel\) 所启动的所有线程称为网格\(grid\) ，同一个网格的线程共享相同的全局内存空间，grid是线程结构的第一个层次，而网格又可以分为很多线程块(block)，一个线程块里面包含很多线程，这是第二个层次。 为了编程方便，\(grid\) 和\(block\) 都是定义为 dim3 类型的变量，dim3 可以看成是包含三个无符号整数\((x,y,z)\) 成员的结构体变量，在定义时，缺失值初始化为1。因此，grid和block可以灵活地定义为1-dim，2-dim以及3-dim的结构，对于，\(knernel\)在定义调用时也必须通过执行配置 <<<grid,block>>>来指定kernel所使用的线程数及结构。

fHaSP2zNs

​ 所以，为了方便编程，CUDA中使用了 dim3 类型（dim3 是基于unit3定义的矢量类型，相当于由3个 unsigned int类型组成的结构体）的内建变量 threadIdx 和 blockIdx。这样，就可以使用一维、二维或三维的索引来标识线程，构成 一维、二维或三维线程块。使得线程组织形式对各种域（向量、矩阵，或者高维张量）中数据的划分变得直观、自然。

  • 对于一维的block，线程的\(threadID\)就是\(threadId.x\);
  • 对于大小为\((Dx,Dy)\)的二维线程块block，线程的\(threadID\) 是 \((threadIdx.x+threadIdx.x\times{Dx})\);
  • 对于大小为\((Dx,Dy,Dz)\)的三维线程块block, 线程的\(threadID\)是（\(threadIdx.x+threadIdx.y\times{Dx}+threadIdx.z\times{Dx}\times{Dy}\)）;

另外，线程还有内置变量gridDim，用于获取网格块各个维度的大小。

​ 此外，这里简单介绍一下CUDA的内存模块，如图所示。可以看到，每个线程有自己的私有本地内存（\(Local Memory\)）, 而每个线程块有包含共享内存（\(Shared \space Memory\)）。还可以访问一些只读内存块：常用内存（\(Constant \space Memory\)）和纹理内存 （\(Texture \space Memory\)）。内存结构涉及到程序优化，这里就过多讨论。

fHaRlklWa

​ 还有重要一点，你需要对\(GPU\)的硬件实现有一个基本的认识。上面说到了\(kernel\)的线程组织层次，那么一个\(kernel\) 实际上会启动很多线程，这些线程是逻辑上是并行的，但是在物理层也是无法却并不一定。这其实和CPU的多线程有类似之处，多线程如果没有多核支持，在物理层也无法实现并行的。但是好在\(GPU\) 存在很多CUDA核心，充分利用CUDA核心可以充分发挥GPU的并行计算能力。GPU硬件的一个核心组件是SM，前面已经说过，SM是Streaming Multiprocessor，SM的核心组件包括的CUDA核心、共享内存、寄存器等，SM可以并发的执行上，一个线程块只能在一个SM上被调度。SM一般可以调度多个线程块，这要看SM本身的能力。那么有可能一个kernel的各个线程块被分配多个SM，所以grid只是逻辑层，而SM才是执行的物理层。SM采用的是SIMT(Single-Instruction, Multiple-Thread，单指令多线程)架构，基本的执行单元是线程束（warps)，线程束包含32个线程，这些线程同时执行相同的指令，但是每个线程都包含自己的指令地址计数器和寄存器状态，也有自己独立的执行路径。所以尽管线程束中的线程同时从同一程序地址执行，但是可能具有不同的行为，比如遇到了分支结构，一些线程可能进入这个分支，但是另外一些有可能不执行，它们只能死等，因为GPU规定线程束中所有线程在同一周期执行相同的指令，线程束分化会导致性能下降。当线程块被划分到某个SM上时，它将进一步划分为多个线程束，因为这才是SM的基本执行单元，但是一个SM同时并发的线程束数是有限的。这是因为资源限制，SM要为每个线程块分配共享内存，而也要为每个线程束中的线程分配独立的寄存器。所以SM的配置会影响其所支持的线程块和线程束并发数量。总之，就是网格和线程块只是逻辑划分，一个kernel的所有线程其实在物理层是不一定同时并发的。所以kernel的grid和block的配置不同，性能会出现差异，这点是要特别注意的。还有，由于SM的基本执行单元是包含32个线程的线程束，所以block大小一般要设置为32的倍数。

内存类型： 内存作用：
全局内存（Global Memory） 容量最大（通常数GB），所有线程可访问，但延迟高（400-800周期）
共享内存（shared Memory） 片上内存，速度比全局内存快100倍，但容量有限（每SM通常16-64KB）
寄存器（Registers） 最快的存储，每个线程私有
常量内存（Constant Memory） 只读缓存，适合广播数据
纹理内存（Texture Memory） 专为图形处理优化的特殊缓存

​ 内存访问特性比较：

内存类型 物理位置 作用域 带宽、速度 使用场景 显式控制关键字
寄存器 GPU芯片寄存器 线程私有 最高（1周期） 高频访问的私有变量（如循环计数器） 自动分配（局部变量）
共享内存 GPU芯片上的SM处理器 线程块共享 高（1-32周期） 线程协作（如规约运算、矩阵分块） __share__
本地内存 实际在全局内存中内存 线程私有 中低（\(\approx\)全局内存） 大数组或寄存器不足时的溢出变量 编译器自动分配
全局内存 GPU设备显存 所有线程+主机 中（400~800周期） 大规模数据存储，需要频繁访问时需合并访问优化 cudaMalloc分配
常量内存 GPU芯片上的缓存 所有线程只读 中（缓存加速） 需要广播给所有线程的至多 __constant__
纹理内存 GPU专用缓存 所有线程 中 （优化访存） 图形处理、具有空间局部性的非对齐访问 纹理API绑定
主机内存 CPU内存 主机+设备（需要拷贝） 最低（PCLe瓶颈） CPU-GPU数据传输的中间存储 malloc、cudaHostAlloc

下面我将详细地介绍CUDA中各种内存管理函数的功能、参数和使用方法。

CUDA是一种用于异构并行计算的编程模型，经常需要在主机端（host）和设备端（Device）之间进行数据传输。这是因为CUDA核函数传入的必须是指向其中处理GPU显存的三个关键的API：cudaMalloc,cudaMemcpy和 cudaFree。

  • cudaMalloc：
其接口API形式： cudaError_t cudaMalloc(void ** devPtr,size_t size )
函数功能： 在设备上分配线性内存size字节，并通过指针返回分配的内存devPtr。分配的内存对应任何类型的变量。记忆没有被清除。失败时返回 cudaErrorMemoryAllocation。
参数： devPtr 设备内存分配指针；size ：分配的字节数
返回值： cudaSuccess , cudaErrorMemoryAllocation

注意事项：

分配的内存

  • cudaMemcpy ：

    其接口形式： cudaError_t cudaMemcpy(void * dist, const void * src,size_t count,CudaMemcpyKind kind)
    函数功能： 将指向的内存区域的字节复制到指向的存储区域
    参数： dist-目的存储地址；src -源内存地址；count-复制内存的字节数； kind-传输类型
    返回值： cudaSuccess，cudaErrorInvalidValue，cudaErrorInvalidDevicePointer，cudaErrorInvalidMemcpyDirection

  • cudaFree：

    其接口形式： cudaError_t cudaFree(void * devPtr)
    函数功能： 释放由 指向的内存空间，该空间必须是之前调用cudaMalloc()或cudaMallocPitch()时返回过的。否则，或者如果cudaFree()之前已被调用过，则返回错误。如果 为 0，则不执行作。cudaFree() 在失败时返回cudaErrorInvalidDevicePointer。
    参数: devPtr -设备指针指向内存释放
    返回值： cudaSuccess，cudaErrorInvalidDevicePointer, cudaErrorInitialization

    ​

1.2 CUDA的第一个程序—矩阵赋值(Matrix Assign)

​ 在本节通过一个矩阵赋值（matrix Assign）例子开始真正的CUDA程序实现，本例是在SDK中template程序的基础上修改得到的。\(template\) 是 \(NVIDIA\) 公司提供的CUDA程序模板，也就是CUDA程序最基本的框架。要创建一个CUDA程序，可以把整个template文件复制一份。在一个CUDA程序中，基本的主机端代码主要完成以下的功能：

  • 启动CUDA，使用多卡时应该时应该加上设备号，或使用\(cudaSetDevice()\)设备GPU设备；
  • 为输入数据分配内存空间；
  • 初始化输入数据；
  • 为GPU分配内存，用于存放输入数据；
  • 将内存中的输入数据拷贝到显存；
  • 为GPU分配显存，用于存放输出数据；
  • 调用device端的kernel进行计算，将结果写到显存中的对应区域；
  • 为CPU分配内存，用于存放GPU传回来的输出数据；
  • 将显存中的结果读取到内存；
  • 释放内存和显存空间；
  • 退出CUDA；

最简单的设备端代码主要完成以下功能：

  • 从显存读取数据到GPU片内；

  • 对数据进行处理；

  • 将处理后的数据写回显存；

    其整个工程包含了三：

    （1）主程序文件CPU-Host端程序（example1main.cu）；

​ （2）GPU设备端函数的处理函数头文件（example_matrixassign_kernel.cuh）；

​ （3）GPU设备端函数的处理函数文件（example_matrixassign_kernel.cu）；

File1：主程序文件CPU-Host端程序（example1main.cu）；

#include<stdio.h> //系统头文件
#include<stdlib.h>
#include<string.h>
#include<math.h>

#include"cuda_runtime.h" //cuda项目头文件
#include"device_launch_parameters.h"
#include"example_matrixassign_kernel.cuh"  //核函数的数据的头文件


void runTest(int argc, char** argv);

int main(int argc,char** argv){

	runTest(argc,argv);

}

void runTest(int argc, char** argv){

	unsigned int num_blocks = 4;  //定义网格中的线程块数量
	unsigned int num_threads= 4;  //定义每个线程块中的线程数量

	unsigned int mem_size = sizeof(float) * num_blocks * num_threads; //为了数据分配的存储器大小，这里每一个人线程计算一个flaot

	//在host端分配内存，h_表示host端，i表示input，o表示output
	float* h_idata = nullptr;
	float* h_odata = nullptr;

	h_idata =(float *)malloc(mem_size);
	h_odata = (float*)malloc(mem_size);

	if(h_idata != nullptr) {
	   memset(h_idata, 0, mem_size);
	}else{
		return;
	}
	if(h_odata!=nullptr){
		memset(h_odata, 0, mem_size);
	}else{
		return;
	}
	
	//在device端分配显存，d_表示device端，i表示input，o表示output
	float* d_idata = nullptr;
	float* d_odata = nullptr;

	cudaError_t cudaStatus;  //cuda状态判断

	cudaStatus=cudaMalloc((void**)&d_idata, mem_size);
	if(cudaStatus != cudaSuccess){
		printf("d_idata is cudaMalloc failed!\n");
		return;
	}
	cudaStatus=cudaMalloc((void**)&d_odata, mem_size);
	if(cudaStatus!=cudaSuccess){
		printf("d_odata is cudaMalloc failed!\n");
		return;
	}
	
	//初始化内存中的值
	for(unsigned int i = 0; i < num_threads * num_blocks;i++){
		h_idata[i] =1.0f;
	}//end for(unsigned int i = 0; i < num_threads * num_blocks;i++)

	//将内存中的输入数据读入设备端显存，这样就完成了主机对设备的数据写入
	cudaStatus=cudaMemcpy(d_idata,h_idata,mem_size,cudaMemcpyHostToDevice);

	//设置运行参数，即网格的形状和线程块的形状
	dim3 grid(num_blocks,1,1);
	dim3 block(num_threads,1,1);

	// 运行核函数，调用GPU进行运算
	testMatrixAssignKernel <<<grid, block>>> (d_idata,d_odata);

	//将结果从显存写入内存
	cudaStatus = cudaMemcpy(h_odata,d_odata,mem_size,cudaMemcpyDeviceToHost);

	//打印结果
	printf("赋值前的矩阵：\n");
	for (unsigned int iblock = 0; iblock < num_blocks; iblock++) {
		for (unsigned int ithread = 0; ithread < num_threads; ithread++) {
			printf("%5.0f", h_idata[iblock * num_threads + ithread]);
		}//end for(unsigned int ithread = 0; ithread < num_threads; ithread++)
		printf("\n");
	}//end for(unsigned int iblock = 0; iblock < num_blocks; iblock++)

	printf("赋值后的矩阵：\n");
	for(unsigned int iblock = 0; iblock < num_blocks; iblock++){
		for(unsigned int ithread = 0; ithread < num_threads; ithread++){
			printf("%5.0f",h_odata[iblock*num_threads+ithread]);
		}//end for(unsigned int ithread = 0; ithread < num_threads; ithread++)
		printf("\n");
	}//end for(unsigned int iblock = 0; iblock < num_blocks; iblock++)

	//输出存储器指针
	free(h_idata);
	free(h_odata);
	cudaFree(d_idata);
	cudaFree(d_odata);
}

从代码中看出，CUDA的主机端代码与C语言非常相似。但也有一部分C语言中没有的语句，下面逐一进行分析。

​ （1）cudaMalloc(size)在显存global memory上分配大小为size字节的线性空间。需要注意的是，与malloc和free一样，cudaMalloc() 也必须与cudaFree()成对使用，否则无法释放显存空间，运行几次程序以后显卡上就没有显存可供分配，程序也就无法正常运行了。另外，为了杜绝指针指费的情况现象，最好在程序结束前将指针赋空并摧毁。

​ （2） cudaMemcpy()用于拷贝存储器中的数据，其中第二参数是指向目标的指针，第二个参数是指向源的指针，第三个参数是需要拷贝的字节数，第四个参数是拷贝操作的类型。拷贝操作类型共有三种：

  • cudaMemcpyDeviceToHost 将显存中的数据拷贝内存中；
  • cudaMemcpyHostToDevice 将内存中的数据拷贝到显存中；
  • cudaMemcpyDeviceToDevice将global memory中的数据拷贝到同一个CUDA上下文的global的另一个区域中；

​ （3）<<<>>>运算符对kernel函数完整的执行参数配置形式是<<<Dg,Db,Ns,S>>>，其中各个参数的含义是：

  • 参数Dg用于定义整个grid的维度和尺寸，为dim3类型，但实际上只有前两维可以不为1。Dim3 Dg(Dg.x,Dg.y,1)中每行有Dg.x个block，每列有Dg.y个block的维度，第三维恒为1。
  • 参数Db为dim3类型，用于定义每个block的维度与尺寸。Dim3 Db(Db.x,Db.y,Db.z) 中每行有Db.x个thread，每列Db.y个thread，高为Db.z，可以定义三维尺寸。整个block中共有Db.x*Db.y*Db.z 个线程；
  • 参数Ns是一个可选参数，用于设置每个block的共享内存shared memory以外，最多能够动态分配的shared memory大小，单位为Byte。
  • 参数\(s\)是一个cudaStream_t类型的可选参数，初始值为0。在本案例中没有用到Stream的相关内容因此这个参数不填，默认为0号流。

File2：主程序文件CPU-Host端程序（example1main.cu）；

​


#pragma once
#ifndef EXAMPLE_MATRIXASSIGN_KERNEL_H
#define EXAMPLE_MATRIXASSIGN_KERNEL_H

#include<stdio.h>
#include"cuda_runtime.h"

__global__ void testMatrixAssignKernel(float* data_input, float* data_output);


#endif // !_EXAMPLE_MATRIXASSIGN_KERNEL_H_

File2：主程序文件CPU-Host端程序（example1main.cu）；


__global__ void testMatrixAssignKernel(float *data_input,float *data_output){

	//shared memory,extern表示大小由host端的Ns参数确定
	extern __shared__ float sdata[];

	const unsigned int bid = blockIdx.x; //线程所在的block的索引号
	const unsigned int tid_in_block = threadIdx.x; //线程在block中的位置
	const unsigned int tid_in_grid = blockDim.x * blockIdx.x + threadIdx.x;

	//按行划分任务时，线程在整个grid中的位置

  // 将数据从global memory读入shared memory
	sdata[tid_in_block] = data_input[tid_in_grid];
	//读入数据后进行一次同步，保证计算时所有数据均已到位
	__syncthreads();

	// 计算
	sdata[tid_in_block] = (float)tid_in_grid;
	//  sdata[tid_in_block] *= (float)tid_in_block;
	//  sdata[tid_in_block] *= (float)tid_in_grid;

	  //进行同步，确保要写入的数据已经被更新
	__syncthreads();

	// 将shared memory中的数据写到global memory
	data_output[tid_in_grid] = sdata[tid_in_block];


}

由上可知，最简单的__gloabal__程序由以下的过程组成：

  1. 分配\(shared \space memory\)；
  2. 将\(global\space memory\) 中的数据读入\(shared \space memory\);
  3. 将进行计算，将结果写到\(shared \space memory\);
  4. 将\(shared\)中的结果写到\(global \space memory\) ;

​ 进行一次GPU计算，要在多种存储器进行几次数据传输，要消耗相当多的时间。这导致了较大的延迟，这导致使\(GPU\) 不适合处理一些实时性要求很高的应用。不同存储器间的数据传输速率和使用方法有很大差异，开发人员需要根据硬件的特点来设计算法，以优化存储器访问。在理想情况下，在所有的存储器传输进行的同时，GPU的各个核心也始终在进行计算，这样就能够很好的隐藏各种访问延迟。CUDA 并不是一种完全硬件透明的语言，程序员需要根据硬件特征将任务进行合理的分解，在编程时对数据传输和寄存器访问进行优化。

​ __global__前缀表示这一段代码是cuda GPU端内核函数。内核函数运行在设备上，其返回类型必须为void。__global__函数中是每一个线程要执行的语句，但由于\(shared\space memory\)和同步的存在，在最好将__global__函数理解为对每一个block的行为的描述。

​ 在这一端内核函数中，首先定义了\(shared \space memory\) 中的变量；然后根据内建变量定义每一个block和thread的索引，对任务进行划分；最后，每一个线程执行了相同的求和运算，但处理数据不同，由线程的索引决定的。程序员在编写__global__函数之前，要先对任务进行划分，设计各个block的工作流程后，做到成竹在胸。

​ 由于CUDA采用了两层并行，因此本例在划分任务时，每个thread在grid中的索引\(tid\_in\_grid\) 是由thread所在block内编号tid计算得来的。计算出每个线程的索引后，就可以根据索引处理线程中不同的数据，请读者好好体会这一点。

​ extern __shared__ float sdata[] 在shared memory中为数组data动态分配了空间。extern 在设备端和主机端有不同的含义：__device__和__global__ 函数中表示动态分配，而在主机端函数中表示外部变量。如果要静态分配一块 shared memory，那么在__shared__之前就不加extern，还必须在[]中写上要分配的字节数。动态分配的shared memory大小，是<<<>>>的执行参数中第三个参数规定的大小。关于shared memory大小。

​ CUDA定义了一些内建变量如下：

  1. gridDim： 网格的维度的变量，dim3类型
  2. blockIdx: 块的索引变量，unit3类型
  3. blockDim：块的维度变量，dim3类型
  4. threadIdx：块内的线程索引变量，unit3类型
  5. warpSize：线程中的warp大小，int类型

其输出结果：

输出结果

​

posted @ 2026-01-05 18:33  GeoFXR  阅读(144)  评论(0)    收藏  举报