CUDA 学习笔记

最新推荐文章于 2025-03-16 15:25:20 发布

ipanda_zll

最新推荐文章于 2025-03-16 15:25:20 发布

阅读量568

点赞数

本文链接：https://blog.youkuaiyun.com/ipanda_huanhuan/article/details/107300804

版权

本文介绍了CUDA编程，包括在Windows10环境下使用VS2019和CUDA11.0进行编程的基础操作，如数据传输。重点讨论了统一内存的概念和使用，以及计算运行时间和错误检查。同时，深入探讨了CUDA的优化策略，如合并内存访问、减少线程束中的指令分歧、循环展开和使用cudaOccupancyMaxPotentialBlockSize计算最佳块大小。还提到了线程束、寄存器和共享内存在GPU架构中的重要性，强调了优化寄存器使用和利用共享内存提高性能的方法。

摘要生成于 C知道，由 DeepSeek-R1 满血版支持，前往体验 >

windows10 VS2019 CUDA11.0
1.主流的标准操作：创建一些额外的数组并显式地使用cudaMemcpy（）函数，在主机和设备间通过PCIe总线进行数据的输入和输出。

//程序改编自书籍：《CUDA高性能并行计算》
#ifndef distance_h
#define distance_h

void distanceArray(float* out, float* in, float ref, int n);
float scale(int i, int n);

#endif

#include <stdio.h>
//也可不加此头文件，因为CUDA内联文件已经包含了math.h，程序可以正常运行
//只是不加此头文件，VS会因为识别不了sqrt()函数而报错
#include <cmath>
//也可不加以下头文件，程序可以正常运行，只是VS会报错
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#define TPB 32

//从核函数中调用并在GPU上执行，因此需要添加__device__标识。
__device__ float distance(float x1, float x2)//左右各两个下划线
{
	return sqrt((x2 - x1) * (x2 - x1));
}

//如所有的核函数一样，是从主机上调用而在设备上执行，因此需要使用修饰符__global__，并且设置返回类型为void。
__global__ void distanceKernal(float* d_out, float* d_in, float ref)
{
	const int i = blockIdx.x * blockDim.x + threadIdx.x;//计算出线程索引号，然后分散到各线程里进行并行计算
	const float x = d_in[i];
	d_out[i] = distance(x, ref);
	printf("i=%2d: dist from %f to %f is %f.\n", i, ref, x, d_out[i]);
}

//以下为主流的标准操作：创建一些额外的数组并显式地使用cudaMemcpy（）函数，在主机和设备间通过PCIe总线进行数据的输入和输出。

//函数distanceArray()调用核函数distanceKernel()。这样的函数被称为封装（warpper）或者启动函数（launcher）。
__host__ void distanceArray(float* out, float* in, float ref, int n)
{
	//声明指向设备端数组的指针
	float* d_in = NULL;
	float* d_out = NULL;
	//在设备端为设备端数组分配内存
	cudaMalloc(&d_in, n * sizeof(float));
	cudaMalloc(&d_out, n * sizeof(float));
	//把主机端上的数组复制到设备端对应的数组中
	cudaMemcpy(d_in, in, n * sizeof(float), cudaMemcpyHostToDevice);
	//启动核函数进行相关的计算
	distanceKernal << <n / TPB, TPB >> > (d_out, d_in, ref);
	//把设备端的输出数组复制到主机端对应的数组中
	cudaMemcpy(out, d_out, n * sizeof(float), cudaMemcpyDeviceToHost);
	//释放分配给设备端数组的内存
	cudaFree(d_in);
	cudaFree(d_out);
}

//在CPU执行的for循环内调用，来生成输入数据。合适的修饰符是__host__,但是我们可以不添加它，因为它是默认的标识符。
__host__ float scale(int i, int n)
{
	return (float)i / (n - 1);
}

#include <iostream>
#include "distance.h"
#define N 64
int main(int argc, char** argv)
{
	float* in = (float*)calloc(N, sizeof(float));
	float* out = (float*)calloc(N, sizeof(float));
	const float ref = 0.5f;

	for (int i = 0; i < N; i++)
	{
		in[i] = scale(i, N);
	}

	distanceArray(out, in, ref, N);
	free(in);
	free(out);
	return 0;
}

2.统一内存：它可以使得一个数组在设备和主机上都能访问统一内存。最关键的部分是函数cudaMallocManaged（），使用这个函数来为托管内存分配空间（在定义了指针以后，与cudaMalloc（）函数使用方法类似）。托管内存数组会带来一个开发效率和执行效率的折衷。

#include <stdio.h>
#include <cmath>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <helper_cuda.h>//需要手动包含此文件，不然VS识别不了
#define TPB 32
#define N 1024

//从核函数中调用并在GPU上执行，因此需要添加__device__标识。
__device__ float distance(float x1, float x2)//左右各两个下划线
{
	return sqrt((x2 - x1) * (x2 - x1));
}

//如所有的核函数一样，是从主机上调用而在设备上执行，因此需要使用修饰符__global__，并且设置返回类型为void。
__global__ void distanceKernal(float* d_out, float* d_in, float ref)
{
	const int i = blockIdx.x * blockDim.x + threadIdx.x;//计算出线程索引号，然后分散到各线程里进行并行计算
	const float x = d_in[i];
	d_out[i] = distance(x, ref);
	printf("i=%2d: dist from %f to %f is %f.\n", i, ref, x, d_out[i]);
}

//在CPU执行的for循环内调用，来生成输入数据。合适的修饰符是__host__,但是我们可以不添加它，因为它是默认的标识符。
__host__ float scale(int i, int n)
{
	return((float)i) / (n - 1);
}


//统一内存：它可以使得一个数组在设备和主机上都能访问统一内存。
//最关键的部分是函数cudaMallocManaged（），使用这个函数来为托管内存分配空间（在定义了指针以后，与cudaMalloc（）函数使用方法类似）。
//托管内存数组会带来一个开发效率和执行效率的折衷。
int main(int argc, char** argv)
{

	//声明指向在设备和主机上访问统一内存的输入、输出数组的指针
	float* in = NULL;
	float* out = NULL;

	//错误检查有一定代价，那就是降低了程序的性能
	//为输入、输出数组分配统一内存
	checkCudaErrors(cudaMallocManaged(&in, N * sizeof(float)));
	checkCudaErrors(cudaMallocManaged(&out, N * sizeof(float)));
	
	const float ref = 0.5f;
	for (int i = 0; i < N; i++)
	{
		in[i] = scale(i, N);
	}

	//启动核函数进行相关的计算
	distanceKernal << <N / TPB, TPB >> > (out, in, ref);
	checkCudaErrors(cudaGetLastError());
	//checkCudaErrors(cudaPeekAtLastError());
	//调用cudaDeviceSynchronize（）保证在返回之前，GPU上的计算已经完成
	checkCudaErrors(cudaDeviceSynchronize());

	//释放分配给设备端数组的内存
	checkCudaErrors(cudaFree(in));
	checkCudaErrors(cudaFree(out));

	return 0;
}

3.计算运行时间与检查错误

#include <stdio.h>
#include <cmath>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <helper_cuda.h>
#define TPB 16

__host__ float scale(int i, int n)
{
	return (float)i / (n - 1);
}
__device__ float distance(float x1, float x2)//左右各两个下划线
{
	return sqrt((x2 - x1) * (x2 - x1));
}

__global__ void distanceKernal(float* d_out, float* d_in, float ref)
{
	const int i = blockIdx.x * blockDim.x + threadIdx.x;//计算出线程索引号，然后分散到各线程里进行并行计算
	const float x = d_in[i];
	d_out[i] = distance(x, ref);
	printf("i=%2d: dist from %f to %f is %f.\n", i, ref, x, d_out[i]);
}


__host__ void distanceArray(float* out, float* in, float ref, int n)
{
	//创建用于计时的事件变量
	cudaEvent_t StartMemcpy, StopMemcpy;
	cudaEvent_t StartKernal, StopKernal;
	checkCudaErrors(cudaEventCreate(&StartMemcpy));
	checkCudaErrors(cudaEventCreate(&StopMemcpy));
	checkCudaErrors(cudaEventCreate(&StartKernal));
	checkCudaErrors(cudaEventCreate(&StopKernal));


	float* d_in = NULL;
	float* d_out = NULL;

	checkCudaErrors(cudaMalloc(&d_in, n * sizeof(float)));
	checkCudaErrors(cudaMalloc(&d_out, n * sizeof(float)));

	//记录事件
	checkCudaErrors(cudaEventRecord(StartMemcpy));

	checkCudaErrors(cudaMemcpy(d_in, in, n * sizeof(float), cudaMemcpyHostToDevice));
	checkCudaErrors(cudaEventRecord(StopMemcpy));

	checkCudaErrors(cudaEventRecord(StartKernal));

	distanceKernal << <n / TPB, TPB >> > (d_out, d_in, ref);
	checkCudaErrors(cudaEventRecord(StopKernal));


	checkCudaErrors(cudaMemcpy(out, d_out, n * sizeof(float), cudaMemcpyDeviceToHost));

	//确保函数全部执行完毕
	checkCudaErrors(cudaEventSynchronize(StopMemcpy));
	checkCudaErrors(cudaEventSynchronize(StopKernal));

	//计算并输出运行时间
	float MemcpyTime = 0;
	checkCudaErrors(cudaEventElapsedTime(&MemcpyTime, StartMemcpy, StopMemcpy));
	float KernalTime = 0;
	checkCudaErrors(cudaEventElapsedTime(&KernalTime, StartKernal,StopKernal));
	printf("Kernal time is %f ms.\n", KernalTime);
	printf("Data transfer time is %f ms.\n", MemcpyTime);

	checkCudaErrors(cudaFree(d_in));
	checkCudaErrors(cudaFree(d_out));
}

4.摘录自 “基于结构光投影的运场物体高速实时三维测量方法研究”。
在这里插入图片描述
一个SM可能包含多个线程块。
SM是GPU架构的核心，而寄存器和共享内存是SM中的稀缺资源。CUDA将这些资源分配到SM中的所有常驻线程里。因此，这些有限的资源限制了在SM上活跃的线程束数量，而活跃的线程束数量对应于SM上的并行量(link)。

5.优化方向
1）合并内存访问：当一个线程束中全部的32个线程访问一个连续的内存块时，满足合并内存访问，可以最大化全局内存吞吐量。
2）尽可能使得同一个 warp (线程束)中的32个线程的结束时间相同，否则，因为处理器是以 warp 为单位来调度线程的，提前完成任务的线程会闲置，直到整个 warp 的线程都完成了任务。典型例子是分支语句。
3）展开循环。
4）优化核函数。
5）cudaOccupancyMaxPotentialBlockSize，用于计算当SM占有率最高时，blocksize的大小。占用率是指SM的实际的活动线程束数量与每个多处理器的总的线程束数量的比值。

6.摘自“摄影测量影像快速立体匹配关键技术研究”。
1)线程束(warp)是 GPU 的基本执行单元，包含32个线程，GPU每次调用线程都是以线程束为单位的.
2)在线程核函数中使用较少的寄存器可使SM中有更多的常驻线程块，增加使用率和性能。而如果一个核函数使用了超过硬件限制数量的寄存器，系统将会用本地内存代替多出的寄存器，这会降低算法性能。
3)共享内存：共享内存是片上内存，与本地内存和全局内存相比具有更高的带宽和更低的延迟。共享内存由线程块分配，线程块中的每个线程都可以共享其存储空间。常用的合作方式是在计算前，将全局内存读进共享内存中，而读取的方式是每个线程负责读取某一个位置的数据，读完之后线程块内的所有线程都能够使用整个共享内存中的数据。因为共享内存的延迟低带宽高，所以这种方式比直接读取全局内存要高效。
4)全局内存：全局内存常驻于设备内存中，是 GPU 中空间最大、延迟最高、最常使用的内存，