CUDA 编程：相关概念，程序结构

最新推荐文章于 2024-03-10 13:47:26 发布

AI路漫漫

最新推荐文章于 2024-03-10 13:47:26 发布

阅读量393

点赞数

分类专栏： cuda编程文章标签： c++ cuda

本文链接：https://blog.csdn.net/weixin_46192930/article/details/112033439

版权

cuda编程专栏收录该内容

2 篇文章 0 订阅

订阅专栏

CUDA 介绍

计算统一设备架构（Compute Unified Device Architecture CUDA），是一套并行计算平台和编程模型。支持英伟达的GPU 卡。CUDA 可以使用简单的编程API 在图形处理单元（GPU）上创建大规模并行应用程序。程序猿可以通过使用 CUDA C 和 C 艹利用GPU 的性能加速应用程序。CUDA 编程就类似与C ，就是添加需要利用GPU 并行性的关键字。允许程序猿指定CUDA 代码那个部分在 CPU 上执行还是在GPU 上执行。

并行处理

在时钟速度趋于饱和的现在，需要新的计算模式。比如选择更快的时钟，使用更多的晶体管，但是功耗限制了这两个方法。第三种就是使用拥有许多可以并行执行任务的更小更简单的处理器。这就是GPU 的计算模式，不是可以执行复杂任务的强大处理器，而是有许多小而简单的可以并行工作的处理器。

GPU 架构和 CUDA 介绍

现在的GPU 除渲染图形图像还可以进行数学计算，被称为通用GPU（GPGPU），CPU 是具有复制控制硬件和较少数据计算硬件。GPU 则是具有简单的控制硬件和更多的数据计算硬件，使之具有并行计算能力。串行的CPU 是优化延迟的（完成给定任务所花费的时间），GPU 是优化吞吐量的（给定时间完成任务的数量）。好像看起来一样啊。。。改进一个另一个方面也会提高，但看待这个改进的方式不同。。。GPU 的这种优化就很适合图像处理和计算机视觉应用，在给定的时间内处理更多的像素。现在为了与GPU 交互，并利用并行计算能力，需要一个由 CUDA 提供的简单的并行编程架构。

CUDA 架构

CUDA 架构包含几个专门为GPU 通用计算而设计的特性。一个 unified shedder 管道 ，允许GPU 芯片上的所有算术逻辑单元（ALU） 被一个 CUDA 程序编组。指令集也适合与一般用途的计算，而不是特定于像素的计算。它还允许对内存的任意读写访问，这些特性使CUDA GPU 架构在通用应用程序中非常有用。。。。为什么我大学没选计算机。。。

GPU 具有很多并行处理单元（核心，Core），硬件上这些 core 分为流处理器和流多处理器。 GPU 有这些流多处理器的网格，软件方面，CUDA 程序是作为一系列并行运行的多线程（Thread）来执行的。每个线程都在不同核心上执行，可以看作 GPU 是由多个块 block 的组合，每个块可以执行多个线程。每个块绑定到GPU 上的不同流多处理器。调度器完成块和流多处理器之间的映射。来自同一块的线程可以相互通信。GPU 有一个分层的内存结构，处理一个块和多个块内线程之间的通信。

CUDA 包含主机（CPU 和其内存）和设备（GPU和其内存）的代码。主机代码由C 或 C艹的编译器在CPU 上编译。设备代码由 GPU 编译器在 GPU 上编译。主机代码通过内核调用设备代码，在设备上并行启动多个线程，设备上启动多少线程由程序猿决定。

和普通的C 代码相比，它类似于正常的串行C 代码，但是实在大量内核上并行执行，而且要使这段代码工作，需要设备显存上的数据。所有在启动线程之前，主机将数据从主机内存复制到设备显存，线程处理来自显存的数据，并将结果保存在显存中，最后这些数据复制回主机的内存进行进一步处理。

所以 CUDA C 程序的开发步骤：

为主机和设备显存中的数据分配内存
将数据从主角内存中复制到显存中
通过指定并行度启动内核
所有线程完成，将数据从设备显存复制回主机内存
释放主机和设备上使用的所有内存。

CUDA 在计算机视觉的应用

计算机视觉和图像处理算法使计算密集型的，需要实时处理这些大图像。这些算法通过CUDA 加速，图像分割，目标检测，分类等应用的实时帧率性能可以大大提升。而且对深度神经网络和深度学习算法进行更快的训练。英伟达的 Jetson TK1 , Jetson TK2 硬件平台可以加速计算机视觉应用。

配置加检测

配置：脑补链接，，，win10 的
完成后再 kernel.cn 文件中替换代码：（虽然看不懂）

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "iostream"

//打印GPU设备信息
void print_GPU_device_info()
{
	int deviceCount;
	cudaGetDeviceCount(&deviceCount);
	for (int i = 0; i < deviceCount; i++)
	{
		cudaDeviceProp devProp;
		cudaGetDeviceProperties(&devProp, i);
		std::cout << "使用GPU device " << i << ": " << devProp.name << std::endl;
		std::cout << "设备全局内存总量： " << devProp.totalGlobalMem / 1024 / 1024 << "MB" << std::endl;
		std::cout << "SM的数量：" << devProp.multiProcessorCount << std::endl;
		std::cout << "每个SM的最大线程数：" << devProp.maxThreadsPerMultiProcessor << std::endl;
		std::cout << "每个SM的最大线程束数：" << devProp.maxThreadsPerMultiProcessor / 32 << std::endl;
		std::cout << "每个Block的共享内存大小：" << devProp.sharedMemPerBlock / 1024.0 << " KB" << std::endl;
		std::cout << "每个Block的最大线程数：" << devProp.maxThreadsPerBlock << std::endl;
		std::cout << "每个Block中可用的32位寄存器数量： " << devProp.regsPerBlock << std::endl;
		std::cout << "======================================================" << std::endl;

	}
}


int main()
{
	print_GPU_device_info();

	while (1);
	return 0;
}

一个基本的CUDA C 程序

要记得，主机代码由标准C 编译器编译，设备代码由 NVIDIA GPU 编译器编译执行，NVIDIA 工具将主机代码提供给标准C 编译器，例如我的 VS2019。另外 GPU 编译器可以再没有任何设备代码的情况下运行CUDA 代码，所有CUDA 代码必须保存为 *.cu 的扩展名。

#include "iostream"
__global__ void myfirstkernel(void){}
int main(void)
{
	myfirstkernel <<<1, 1>>> ();
	printf("hello ,cuda!\n") ; 
	return 0;
}
// 基本没用过c艹啥的，，哎还是python香啊。。

与普通C代码相比，使用了前缀 __global__ ，使用 <<1 , 1>> 调用函数。

__global__ 使CUDA C 再标准 C 中添加的限定符，告诉编译器这个限定符后面的函数定义应该在设备上而不是主机上运行。所以这个 myfirstkernel() 就会运行再设备上。

内核调用：一个CUDA C 的技巧，从主机代码调用设备代码。尖括号内的值表示我们希望再运行时从主机传递给设备的参数。基本上，表示块的数量和将再设备上并行运行的线程数。<< <1,1> >> 就表示 myfirstkernel 这个函数将运行再设备上的一个块和一个线程（每个块几个线程）上。

CUDA 程序结构

CUDA 程序是在主机或 GPU 设备上执行的函数的组合。显示数据并行性的函数就在GPU上执行。GPU 编译器在编译期间隔离这些函数。CUDA 代码基本上与基础C 代码相同，就是添加了一些开发数据并行性所需要的关键字。

CUDA C 中的双变量加法程序

#include <iostream>
#include <cuda.h>
#include <cuda_runtime.h>
// 定义增加两个变量的函数
__global__ void gpuAdd(int d_a, int d_b, int* d_c)
{
	*d_c = d_a + d_b;
}
int main()
{
	int h_c, * d_c;   // 一个主机变量来保存结构，定义一个 设备指针
	cudaMalloc((void**)&d_c, sizeof(int));    // 给设备指针分配内存
	// 一个块，每个块一个线程来执行这个函数
	gpuAdd <<<1, 1>>> (1, 4, d_c);      // 注意这里三个阔折号间没有空格，，
	// 从设备内存中复制结果到主机内存中
	cudaMemcpy(&h_c, d_c, sizeof(int), cudaMemcpyDeviceToHost);
	printf("1 + 4 = %d\n", h_c);
	cudaFree(d_c);  // 是否内存
	return 0;
}

约定 h_ 前缀为主机变量，d_ 前缀为设备变量。

内核调用

使用基础的C 关键字和CUDA 扩展关键字编写的设备代码称为内核。它是主机代码通过内核调用的方式来启动的。简单来说内核调用就是从主机代码启动设备代码。内核调用通常会生成大量的块 block 和线程 thread 来再GPU 上并行地处理数据。与普通的C 函数不同的就是这段代码由多个线程并行执行。

内核启动的语法：kernel <<< number of blocks , number of threads per block, size of shared memory >>> (parameters for kernel)

这个 kernel 需要确保是使用 global 关键字定义的，具有 <<< >>> 内核启动配置，第三个参数是可选的，指定内核使用的共享内存的大小。上例中gpuAdd <<<1, 1>>> (1, 4, d_c) ，其中需要注意 d_c 是指向设备内存的指针，而且参数传递给内核的指针应该仅指向设备显存，如果指向了主机内存，程序就会崩溃。内核执行完成后，设备指针指向的结果可以复制回主机内存，

内核参数的配置

就是启动多线程，可以指定块的数量和每个块中线程的数量，每个块再流多处理器上运行，一个块的线程可以通过共享内存彼此通信。计算机视觉应用程序中，需要处理二维和三维图像，GPU 支持三维网格块和三维线程块：mykernel <<< dim3(Nbx,Nby,Nbz), dim3(Ntx, Nty, Ntz)>>>()，这里 Nbx, Nby, Nbz，表示网格中沿 x,y ,z 轴方向的块数。Ntx 。。就是一个块中沿 x,y,z 方向的线程数。每个没有指定 y,z 的维数默认为1。例如处理一个图像，你可以启动一个 16 * 16 的块网络，每个块 16 * 16 个线程mykernel <<< dim3(16,16), dim3(16,16) >>>()

CUDA API函数

_global_ : 与 _device_ 和 _host_ 一起是三个限定符关键字，这个关键字表示这就是一个设备函数，从主机上调用时将在设备上执行。注意，这个函数只能从主机调用，要在设备上执行函数并从设备函数中调用函数，必须使用 _device_ 关键字。_host_ 关键字用于定义只能从其他主机函数调用的主机函数。类似于普通的C 函数。默认所有函数都是主机函数，_host_ and _device_ 都可以同时用于定义任何类型函数，生成同一个函数的两个副本，一个在主机上执行，另一个在设备上执行。
cudaMalloc : 动态内存分配，用于在设备上分配特定大小的内存块。cudaMalloc(void ** d_pointer,size_t size) example : cudaMalloc((void**) s&d_c, sizeof(int)); 分配一个整形变量大小的内存块。并返回指向内存位置的指针。
cudaMemcpy ：将一个内存区域复制到主机或设备上的其他区域 cudaMemcpy( void * dst_ptr, const void * src_ptr, size_t size, enum cudaMemcpyKind kind ) example : cudaMemcpy( &h_c, d_c, sizeof(int), cudaMemcpyDeaviceToHost) ; 第一个参数是目标指针，第二个参数是源指针，分布指向主机内存或设备显存位置，第三个参数表示数据复制的大小，最后一个参数表示数据复制的方向，（主机到主机，设备到主机之类的。。）需要注意前两个参数指针的方向要一致。
cudaFree : 释放内存空间

将参数传递给 CUDA 函数

按值或引用传递参数给内核函数，

按值传递参数
gupAdd<<<1,1>>>(1,4,d_c)
通过引用传递参数

#include <iostream>
#include <cuda.h>
#include <cuda_runtime.h>
// 定义增加两个变量的函数
__global__ void gpuAdd(int *d_a, int *d_b, int* d_c)
{
	*d_c = *d_a + *d_b;
}
int main()
{
	int h_a, h_b, h_c;
	int* d_a, * d_b, * d_c;
	h_a = 1;
	h_b = 4;   // 初始化主机变量
	cudaMalloc((void**)&d_a, sizeof(int));   // 给设备指针分配内存
	cudaMalloc((void**)&d_b, sizeof(int));
	cudaMalloc((void**)&d_c, sizeof(int));   
	cudaMemcpy(d_a, &h_a, sizeof(int), cudaMemcpyHostToDevice);   // 将主机变量复制到设备显存中
	cudaMemcpy(d_b, &h_b, sizeof(int), cudaMemcpyHostToDevice);
	gpuAdd <<<1, 1>>>  (d_a, d_b, d_c);
	cudaMemcpy(&h_c, d_c, sizeof(int), cudaMemcpyDeviceToHost);   // 将结果从设备显存中复制会主机
	printf(" %d + %d = %d\n",h_a,h_b,h_c);
	cudaFree(d_a);  
	cudaFree(d_b);
	cudaFree(d_c);
	return 0;
}

在使用设备指针并将其传递给内核时，需要遵守一些限制，使用cudaMalloc 分配内存的设备指针只能用于从设备显存中读写。可以作为参数传递给设备函数，而不应该用于从主机函数读写内存。应该使用设备指针从设备函数中读取和写入设备显存，并且应该使用主机指针从主机函数中读取和写入主机内存。

在设备上执行线程

#include <iostream>
#include <stdio.h>
__global__ void kernel(void)
{
	printf("hello, I' m thread in block : %d\n ", blockIdx.x);   // 块ID 
}
/*
可以认为并行启动了16个 kernel 代码的线程副本，每个线程副本有一个属于自己的块ID 和 线程ID 。
threadIdx.x blockIdx.x ，另外多运行几次可以发现，线程块是随机执行的。
*/

int main()
{ 
	kernel <<<16, 1 >>> ();   // 16个并行块
	cudaDeviceSynchronize(); // 这个函数等待所有线程都完成后才结束。
	/*
	启动内核是一个异步操作，发布了内核启动命令，不等内核执行完成控制权就会交回给调用内核的CPU 线程
	这个程序中，CPU线程返回然后执行 printf 不等内核部分执行完成这个程序就完成了。如果不加这个同步函数
	就看不见输出结果了。通过这个同步函数就使程序在内核执行完成之后才退出。
	*/
	printf("All threads are finished ! \n ");  
	return 0;
}

>>>
hello, I' m thread in block : 7
 hello, I' m thread in block : 4
 hello, I' m thread in block : 13
 hello, I' m thread in block : 1
 hello, I' m thread in block : 10
 hello, I' m thread in block : 8
 hello, I' m thread in block : 6
 hello, I' m thread in block : 5
 hello, I' m thread in block : 3
 hello, I' m thread in block : 14
 hello, I' m thread in block : 12
 hello, I' m thread in block : 2
 hello, I' m thread in block : 0
 hello, I' m thread in block : 11
 hello, I' m thread in block : 9
 hello, I' m thread in block : 15
 All threads are finished !