CUDA编程（五）获得GPU加速的关键

grace 1314

已于 2023-10-23 22:01:47 修改

阅读量284

点赞数

文章标签：人工智能

于 2023-10-23 20:07:26 首次发布

本文链接：https://blog.csdn.net/qq_43616651/article/details/133991998

版权

上一节中检查程序的正确性是一个问题，那么程序的性能/执行速度怎么去提高呢。

研究CUDA程序的性能则需要对程序进行比较精确的计时。

1.用CUDA事件计时

在 C++ 中，有多种可以对一段代码进行计时的方法，包括使用 GCC 和 MSVC 都有的 clock 函数和与头文件<chrono> 对应的时间库、GCC 中的 gettimeofday 函数及 MSVC 中QueryPerformanceCounter 和 QueryPerformanceFrequency 函数等。

CUDA提供了基于CUDA事件（CUDA event）的计时方式，可用来给一段CUDA代码（可能包含了主机代码和设备代码）计时。

CUDA计时方式：


cudaEvent_t start,stop;//定义cuda事件类型变量


CHECk(cudaEventCreate(&start));
CHECK(cudaEventCreate(&stop)); 
//使用cudaEventCreate来初始化，


CHECK(cudaEventRecord(start));  //将start传入cudaEventRecord函数，
//在需要计时的代码块之前记录一个代表开始的事件。

//在TCC驱动模式可以省略，但在WDDM驱动模式下的GPU必须保留。
cudaEventQuery(start); //此处不能用CHECK宏函数。

//需要计时的代码块

//可以是主机与设备代码混合


CHECK(cudaEventRecord(stop));  //将stop传入cudaEventRecord函数，
//在需要计时的代码块之后记录一个代表结束的事件。

ChECK(cudaEventSynchronize(stop));
//cudaEventSynchronize 函数让主机等待事件stop被记录完毕。
float elapsed_time;
CHECK(cudaEventElapsedTime(&elapsed_time,start,stop));

printf("Time = %g ms.\n",elapsed_time);
//elapsed_time "已消耗时间",程序执行时间
//调用cudaEventElapsedTime 函数计算 start 和 stop 这两个事件之间的时
//间差（单位是ms）并输出到屏幕

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

//调用cudaEventDestroy 函数销毁start和stop这两个CUDA事件。

上述代码中一些问题：

GPU驱动模式TCC与WDDM：（有请GPT）

VIDIA的GPU驱动支持两种不同的驱动模式：TCC（Tesla Compute Cluster）和WDDM（Windows Display Driver Model）。这两种模式具有不同的特性和适用于不同的使用情况。

TCC 模式（Tesla Compute Cluster）：
- 用途：TCC 模式是专门为GPU计算而设计的模式，它不支持图形输出。这意味着在TCC模式下，GPU不能用于显示图形，而是专注于计算任务。
- 特性：TCC模式下的GPU通常能够提供更高的计算性能，因为它不需要处理与图形相关的任务。此模式适用于需要大规模计算的应用程序，如科学计算、深度学习训练等。
- 要求：要在TCC模式下使用GPU，通常需要专门的Tesla系列GPU（如NVIDIA的Tesla、Quadro或部分GeForce卡），以及在NVIDIA驱动中选择TCC模式。
WDDM 模式（Windows Display Driver Model）：
- 用途：WDDM模式是通用模式，支持图形输出。在WDDM模式下，GPU不仅用于计算，还用于图形渲染，适用于通用桌面应用程序、游戏和多媒体。
- 特性：WDDM模式允许GPU在计算和图形渲染之间切换，使其成为一种通用的图形和计算解决方案。然而，与TCC模式相比，WDDM模式下的计算性能可能略有降低。
- 要求：大多数NVIDIA GPU支持WDDM模式，并且它是默认模式。通常，没有额外的设置就可以在WDDM模式下使用GPU。

要选择GPU的驱动模式，通常需要在NVIDIA控制面板中进行设置，尤其是对于特定的计算任务，可以选择TCC模式以获得更高的计算性能。需要注意的是，选择模式通常需要重新启动系统以使更改生效。

总结，TCC专注计算，WDDM既能做计算，显示图形，做渲染还能用在应用程序，游戏里。

什么是CUDA事件（CUDA event）：

CUDA事件是一种用于测量GPU操作执行时间的机制。它们是NVIDIA CUDA平台的一部分，允许开发者精确地测量CUDA核函数和其他GPU操作的执行时间。CUDA事件通常用于性能分析、优化和调试，以确保GPU计算任务能够在预期的时间内完成。

以下是一些关于CUDA事件的重要信息：

创建和记录事件：你可以使用CUDA API函数来创建和记录事件。通常，你会在代码中插入事件记录的代码，以便在核函数执行之前和之后记录事件。
计算执行时间：记录了事件后，你可以使用CUDA API函数来计算事件之间的时间差，以测量核函数的执行时间。这使你可以了解核函数的性能表现。
异步操作：CUDA事件是异步操作，它们不会阻塞主机线程。这意味着你可以在GPU执行核函数的同时进行其他任务，然后在需要时查询事件以获取执行时间信息。
多GPU支持：如果你使用多个GPU，可以创建和记录事件以分别测量不同GPU上的操作的执行时间。

c语言输出%g是什么格式：

其中%g 是一种格式控制符，通常用于格式化输出浮点数。它会自动选择 %e（科学计数法）或 %f（定点表示法），以显示浮点数的精度，从而避免不必要的零位数。

在.cu文件中可以对C++程序和CUDA程序进行计时。

C++可以自己尝试下，我这里对CUDA程序测试一下。

#include "error.cuh"

#include <stdio.h>

#include "cuda_runtime.h"

#include "device_launch_parameters.h"

#include<math.h>

#include<stdlib.h>


const double EPSILON = 1.0e-15;
const double a = 1.23;
const double b = 2.34;
const double c = 3.75;
void __global__ add(const double* x, const double* y, double* z, int N);
void check(const double* z, const int N);


int main(void) {

	//Host 变量，数组并进行初始化
	const int N = 1000000;
	const int M = sizeof(double) * N;
	double* h_x = (double*)malloc(M);
	double* h_y = (double*)malloc(M);
	double* h_z = (double*)malloc(M);

	for (int n = 0; n < N; ++n) {
		h_x[n] = a;
		h_y[n] = b;
	}

	double* d_x, * d_y, * d_z;

	//计时准备
	cudaEvent_t start, stop;
	float elapsedTime;

	cudaEventCreate(&start);
	cudaEventCreate(&stop);

	cudaEventRecord(start);

	//使用宏CHECK(cuda函数)
	CHECK(cudaMalloc((void**)&d_x, M));
	CHECK(cudaMalloc((void**)&d_y, M));
	CHECK(cudaMalloc((void**)&d_z, M));


	CHECK(cudaMemcpy(d_x, h_x, M, cudaMemcpyHostToDevice));
//	CHECK(cudaMemcpy(d_x, h_x, M, cudaMemcpyDeviceToHost));  //此处为修改错误，

	CHECK(cudaMemcpy(d_y, h_y, M, cudaMemcpyHostToDevice));

	//核函数执行配置参数
	const int block_size = 128; //128个一维线程块
	const int grid_size = (N + block_size - 1) / block_size; //10^8/128个线程块
	


	//调用核函数计算
	add << <grid_size, block_size >> > (d_x, d_y, d_z, N);
	//核函数检查
	CHECK(cudaGetLastError());
	CHECK(cudaDeviceSynchronize());

	//在传输数据时，cudaMemcpy()函数起到了隐式的同步主机与设备的作用。
	CHECK(cudaMemcpy(h_z, d_z, M, cudaMemcpyDeviceToHost));

	check(h_z, N);

	free(h_x);
	free(h_y);
	free(h_z);


	CHECK(cudaFree(d_x));
	CHECK(cudaFree(d_y));
	CHECK(cudaFree(d_z));

	//记录计时
	cudaEventRecord(stop);
	cudaEventSynchronize(stop);

	cudaEventElapsedTime(&elapsedTime, start, stop);
	printf("执行时间: %f ms\n", elapsedTime);

	cudaEventDestroy(start);
	cudaEventDestroy(stop);

	return 0;

}

void __global__ add(const double* x, const double* y, double* z, int N) {
	//核函数中数据与线程的对应
	//单指令-多线程
	//n 表示数组元素指标
	const int n = blockDim.x * blockIdx.x + threadIdx.x;
	if (n < N) {
		z[n] = x[n] + y[n];
	}

}

void check(const double* z, const int N) {
	bool has_error = false;
	for (int n = 0; n < N; ++n) {
		if (fabs(z[n] - c) > EPSILON) {
			has_error = true;
		}
	}
	printf("%s\n", has_error ? "Has error" : "No errors");

}

代码+结果如上。

这里面数据是双精度的。

我们再改一下，改成单精度会如何。即把代码中的float全都改成double

第一测的时间会多点，所以多调试几次后再截的图，可以看出时间还是可以的。

双精度19ms，单精度13ms，我的电脑是垃圾显卡GT1030。

书上说的一些：

对于数组相加的问题，其执行速度是由显存带宽决定的，而不是由浮点数运算峰值决定的。

有效显存带宽定义为GPU在单位时间内访问设备内存的字节数。

有效显存带宽略小于理论显存带宽，进一步说明该问题是访存主导的，即该问题中的浮点数运算所占比例可以忽略不计。

上面代码是cuda内存申请，cuda数据复制，核函数运行，cuda内存释放的过程时间。

下面改一下，试试只测核函数运行时间。

//计时准备
cudaEvent_t start, stop;
float elapsedTime;

cudaEventCreate(&start);
cudaEventCreate(&stop);

cudaEventRecord(start);

//调用核函数计算
add << <grid_size, block_size >> > (d_x, d_y, d_z, N);


//记录计时
cudaEventRecord(stop);
cudaEventSynchronize(stop);

cudaEventElapsedTime(&elapsedTime, start, stop);
printf("执行时间: %f ms\n", elapsedTime);

cudaEventDestroy(start);
cudaEventDestroy(stop);

单精度；

双精度

双精度大概是2倍。

所以可以看到核函数运行时间占的很小，大部分都是数据复制存储所需要的时间。

也可以用CUDA工具箱，nvprof的可执行文件，，可用于对CUDA程序进行更多的性能剖析。

下面有请GPT:

在 Visual Studio 2022 中使用 NVIDIA 的 CUDA 工具集（nvcc 和 nvprof）需要进行一些配置和设置。以下是步骤：

安装 CUDA 工具包：首先，确保你已经安装了NVIDIA CUDA工具包，其中包括CUDA编译器nvcc和性能分析工具nvprof。你可以从NVIDIA官方网站下载并安装CUDA工具包。
配置 Visual Studio 2022：
- 打开 Visual Studio 2022。
- 打开或创建一个CUDA项目，或者在现有项目中添加CUDA支持。
- 在项目属性中配置CUDA：
  - 在 "属性管理器" 中，右键单击你的项目，选择 "属性"。
  - 转到 "配置属性" > "常规"，确保 "CUDA C/C++" 已经启用。
  - 在 "配置属性" > "CUDA C/C++" > "常规" 中，设置 "CUDA Toolkit Custom Dir" 为 CUDA 工具包的安装目录，例如：C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v11.0。你需要根据你安装的CUDA版本进行相应设置。
  - 在 "配置属性" > "CUDA C/C++" > "命令行" 中，确保 "生成器" 设置为 Host (x64)。
使用 nvcc 编译器：
- 在 Visual Studio 2022 中，你可以编写CUDA代码并将其编译成可执行文件。当你构建项目时，Visual Studio 会自动调用 nvcc 编译器来编译 CUDA 文件。
使用 nvprof 进行性能分析：
- 在 Visual Studio 2022 中，你可以使用 nvprof 来进行性能分析。在项目属性中配置启动参数，以便在启动时调用 nvprof。
- 在 "配置属性" > "调试" > "命令参数" 中，添加 --nvprof 或其他 nvprof 选项。
运行项目：构建项目并运行它。你可以在Visual Studio中查看CUDA代码的输出和性能数据。

请注意，上述步骤的具体设置可能因你的项目和CUDA版本而异。确保按照你的需求进行相应的配置。此外，你可以在Visual Studio中使用调试工具进行CUDA程序的调试和性能分析，以便更好地了解代码的性能和行为。

按照上面试了下，看不到。

但是再cuda toolkit文件里确实有这些，nvcc，和nvprof。

有会在win上用的可以评论分享。

2.几个影响GPU加速的关键因素

数据传输的比例：

简单的问题在GPU上运行可能比在CPU中运行还慢，这是因为花在数据传输（从cpu将数据传输到GPU中）上的时间比计算本身还要很多很多。

GPU计算核心和设备内存之间数据传输的峰值理论带宽要远高于GPU和CPU之间数据传输的带宽。

GPU显存带宽一般几百gb每秒，而常用的GPU和CPU内从的PCIe总线只有16GB带宽

它们相差几十倍。要获得可观的GPU加速，就必须尽量缩减数据传输所花时间的比例。有时候，即使有些计算在GPU中的速度并不高，也要尽量在GPU中实现，避免过多的数据经由PCIe传递。这是CUDA编程中较重要的原则之一。

只要数据传输的时间比例小，尽管在GPU中运行速度不高，也要尽量在GPU中实现。

1Gb = 1吉比特 = 10^9 比特

1GB = 1吉字节 = 10^9 字节

1字节 Byte = 8 比特 bit

假设计算任务不是做一次数组相加的计算，而是做次数组相加的计算，而且只需要在程序的开始和结束部分进行数据传输，那么数据传输所占的比例将可以忽略不计。此时，整个CUDA程序的性能就大为提高。

例如，在大量的数据计算时，仅在程序的开始部分将一些数据从主机复制到设备，然后在程序的中间部分偶尔将一些在GPU中计算的数据复制到主机。对这样的计算，用CUDA就有可能获得可观的加速。

算术强度：

问题的算术强度越高，得到的加速比就越高（指单纯cpu计算与GPU计算）。

一个计算问题的算术强度指的是其中算术操作的工作量与必要的内存操作的工作量之比。

算术操作的工作量与必要内存操作的工作量。

例如，在数组相加的问题中，在对每一对数据进行求和时需要先将一对数据从设备内存中取出来，然后对它们实施求和计算，最后再将计算的结果存放到设备内存。这个问题的算术强度其实是不高的，因为在取两次数据、存一次数据的情况下只做了一次求和计算。在CUDA中，设备内存的读、写都是代价高昂（比较耗时）的。

Fused Multiply-Add（FMA，融合乘加）是一种计算机指令，通常由硬件执行单元支持，用于在单个操作中执行浮点数的乘法和加法操作。这个指令对于高性能计算和科学计算等领域非常有用，因为它可以提高计算的速度和精度。

FMA 指令执行以下操作：result = (A * B) + C，其中 A、B 和 C 是浮点数。这是一种将乘法和加法操作合并在一起执行的方式，有助于减少舍入误差，提高计算的精度，并且通常比分开执行乘法和加法操作更快。

FMA 指令还支持三个操作数的形式，其中 A、B 和 C 都可以是浮点数，使其能够执行更复杂的数学运算。例如，你可以执行 result = (A * B) + (C * D) 这样的操作。

例子：复杂运算

下面两个代码都是.cu文件，一个是GPU版，一个是CPU版，都用cudaEvent计时。

1.gpu版本

# ifdef USE_DP
	typedef double real;         //双精度
	const real EPSILON = 1.0e-15;  
#else
	typedef float real;          //单精度
	const real EPSILON = 1.0e-6f;
#endif // USE_DP

//自定义，如果USE——DP有定义就双精度，没定义就单精度。

#include "error.cuh"
#include <stdio.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include<math.h>
#include<stdlib.h>

const real x0 = 100.0;

void __global__ arithmetic(real* d_x, const real x0, const int N);

int main(void) {

	printf("Gpu计算：");

	//Host 变量，数组并进行初始化
	const int N = 1000000;
	const int M = sizeof(real) * N;
	real* h_x = (real *)malloc(M);

	for (int n = 0; n < N; ++n) {
		h_x[n] = n;

	}

	real* d_x;

	//使用宏CHECK(cuda函数)
	CHECK(cudaMalloc((void**)&d_x, M));

	CHECK(cudaMemcpy(d_x, h_x, M, cudaMemcpyHostToDevice));
	//	CHECK(cudaMemcpy(d_x, h_x, M, cudaMemcpyDeviceToHost));  //此处为修改错误，

	//核函数执行配置参数
	const int block_size = 128; //128个一维线程块
	const int grid_size = (N + block_size - 1) / block_size; //10^8/128个线程块

	//计时准备
	cudaEvent_t start, stop;
	float elapsedTime;

	cudaEventCreate(&start);
	cudaEventCreate(&stop);

	cudaEventRecord(start);

	//调用核函数计算
	arithmetic << <grid_size, block_size >> > (d_x, x0, N);

	//记录计时
	cudaEventRecord(stop);
	cudaEventSynchronize(stop);

	cudaEventElapsedTime(&elapsedTime, start, stop);
	printf("执行时间: %f ms\n", elapsedTime);

	cudaEventDestroy(start);
	cudaEventDestroy(stop);

	//核函数检查
	CHECK(cudaGetLastError());
	CHECK(cudaDeviceSynchronize());

	free(h_x);

	CHECK(cudaFree(d_x));

	return 0;
}

void __global__ arithmetic(real* d_x, const real x0, const int N) {
	const int n = blockDim.x * blockIdx.x + threadIdx.x;
	if (n < N) {
		real x_tmp = d_x[n];
		while (sqrt(x_tmp) < x0) {
			++x_tmp;
		}
		d_x[n] = x_tmp;
	}
}

结果：

上面代码还是需要检查错误的error.cuh文件。

2.cpu版本

# ifdef USE_DP
typedef double real;         //双精度
const real EPSILON = 1.0e-15;
#else
typedef float real;          //单精度
const real EPSILON = 1.0e-6f;
#endif // USE_DP

//自定义，如果USE——DP有定义就双精度，没定义就单精度。

#include "error.cuh"
#include <stdio.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include<math.h>
#include<stdlib.h>

const real x0 = 100.0;

void arithmetic(real* x, const real x0, const int N);

int main(void) {
	printf("cpu计算：");

	//Host 变量，数组并进行初始化
	const int N = 1000000;
	const int M = sizeof(real) * N;
	real* x = (real*)malloc(M);

	for (int n = 0; n < N; ++n) {
		x[n] = n;
	}

	//计时准备
	cudaEvent_t start, stop;
	float elapsedTime;

	cudaEventCreate(&start);
	cudaEventCreate(&stop);

	cudaEventRecord(start);

	arithmetic(x, x0, N);

	//记录计时
	cudaEventRecord(stop);
	cudaEventSynchronize(stop);

	cudaEventElapsedTime(&elapsedTime, start, stop);
	printf("执行时间: %f ms\n", elapsedTime);

	cudaEventDestroy(start);
	cudaEventDestroy(stop);

	free(x);
	return 0;

}

void arithmetic(real* x, const real x0, const int N) {
	for (int n = 0; n < N; ++n) {
		real x_tmp = x[n];
		while (sqrt(x_tmp) < x0) {
			++x_tmp;
		}
		x[n] = x_tmp;
	}

}

cpu版本结果：

数组长度为10^6，单精度情况下，gpu 20ms，cpu，220ms，差的挺多的。

单精度GPU相对于CPU的加速比分别为 220/20 = 11;

这里宏定义的USE_DP,在.cu文件，貌似没法定义，所以只能用单精度，目前还不会如何在vs中设置成双精度。

GPT中是这么说的：

在 Visual Studio 中，右键单击项目名称或 CUDA 文件，然后选择 "属性"。
转到 "C/C++" > "预处理器"。
在 "预处理器定义" 中，添加 USE_DP。

进去后，

没有，g。

如果有兄弟会，可以评论区交流。

此外也可以自己改改数组长度N,然后比较cpu和gpu运行效果。

对于算术强度很高的问题，在使用双精度浮点数时Tesla系列的GPU相对于GeForce系列的GPU有很大的优势，而在使用单精度浮点数时前者没有显著的优势。对于算术强度不高的问题（如前面的数组相加问题），Tesla系列的GPU在使用单精度浮点数或双精度浮点数时都没有显著的优势。在使用单精度浮点数时，GeForce系列的GPU具有更高的性价比。

并行规模：

并行规模可用GPU中总的线程数目来衡量。

从硬件的角度来看，一个GPU由多个流多处理器（streamingmultiprocessor，SM）构成，而每个SM中有若干CUDA核心。每个SM是相对独立的。

从开普勒架构到伏特架构，一个SM中最多能驻留（reside）的线程个数是2048。对于图灵架构，该数目是1024。一块GPU中一般有几个到几十个SM（取决于具体的型号）。

所以，一块GPU一共可以驻留几万到几十万个线程。如果一个核函数中定义的线程数目远小于这个数的话，就很难得到很高的加速比。

在数组元素个数N很大时，核函数的计算时间正比于N；在N很小时，核函数的计算时间不依赖于N的值，保持为常数。这两个极限情况都是容易理解的。当N足够大时，GPU是满负荷工作的，增加一倍的工作量就会增加一倍的计算时间。

反之，当N不够大时，GPU中是有空闲的计算资源的，增加N的值并不会增加计算时间。若要让GPU满负荷工作，则核函数中定义的线程总数要不少于某个值，该值在一般情况下和GPU中能够驻留的线程总数相当，但也有可能更小。只有在GPU满负荷工作的情况下， GPU中的计算资源才能充分地发挥作用，从而获得较高的加速比。

因为我们的CPU程序中的计算是串行的，其性能基本上与数组长度无关，所以GPU程序相对于CPU程序的加速比在小N的极限下几乎是正比于N的。在大N的极限下，GPU程序相对于CPU程序的加速比接近饱和。总之，对于数据规模很小的问题，用GPU很难得到可观的加速。

总结：

CUDA程序获得高性能要求：

1.数据传输比例所占时间小。

2.核函数的算术强度较高。

3.核函数中定义的线程数目较多。

所以，在编写与优化CUDA程序时，一定要想方设法（主要是指仔细设计算法）做到以下几点：

• 减少主机与设备之间的数据传输。

• 提高核函数的算术强度。

• 增大核函数的并行规模。

3.CUDA中的数学函数库：

先看看一个概念：（gpt）

"内建函数"（Intrinsic Function）和 "内联函数"（Inline Function）都是编程中用于优化代码执行的概念，但它们有不同的特点和用途。

内建函数（Intrinsic Function）：

内建函数是特定编程语言或编译器提供的函数，通常用于执行底层的硬件操作或高度优化的计算。
这些函数通常与底层硬件操作系统相关，通常由编译器将其翻译为特定的机器指令，从而实现高效的操作。
内建函数通常用于访问特定硬件功能，如处理器指令、SIMD（单指令多数据）操作、特定寄存器的访问等。
内建函数通常对程序员来说是不可见的，编译器负责将其转换为适当的机器代码。

示例：在C/C++中，__builtin_popcount 是一个内建函数，用于计算整数中置位（二进制位为1）的数量。

内联函数（Inline Function）：

内联函数是在编译时展开的函数，通常用于优化函数调用的开销。
当你声明一个函数为内联函数时，编译器会尝试将函数的代码插入到调用函数的地方，而不是实际执行函数调用。
内联函数适用于短小的函数，因为展开函数代码可能会增加程序的体积。
内联函数通常用于提高程序性能，减少函数调用的开销，但不适用于所有函数。编译器通常会根据函数的大小和复杂度来决定是否进行内联。

示例：在C/C++中，你可以使用 inline 关键字将函数声明为内联函数，以建议编译器将其展开。

总结：

内建函数通常由编程语言或编译器提供，用于底层的硬件操作和优化，程序员通常不直接操作它们。
内联函数是由程序员指定的，用于优化函数调用的性能，编译器会根据情况将其展开。
内建函数和内联函数都用于提高代码的性能和效率，但它们的应用领域和实现方式有所不同。

在CUDA数学库中，还有很多类似的数学函数，如幂函数、三角函数、指数函数、对数函数等。

CUDA数学库中的函数可以归纳如下：

1. 单精度浮点数内建函数和数学函数（singleprecisionintrinsics and math functions）。使用该类函数时不需要包含任何额外的头文件。

2. 双精度浮点数内建函数和数学函数（doubleprecisionintrinsicsandmathfunctions）。使用该类函数时不需要包含任何额外的头文件。

3. 半精度浮点数内建函数和数学函数（halfprecisionintrinsicsandmathfunctions）。使用该类函数时需要包含头文件。

4. 整数类型的内建函数（integerintrinsics）。使用该类函数时不需要包含任何额外的头文件。

5. 类型转换内建函数（typecasting intrinsics）。使用该类函数时不需要包含任何额外的头文件。

6. 单指令-多数据内建函数（SIMDintrinsics）。使用该类函数时不需要包含任何额外的头文件。

如，当x是双精度浮点数时，我们只可以用sqrt(x)；当x是单精度浮点数时，我们可以用sqrt(x)，也可以用sqrtf(x)。

在开发CUDA程序时，浮点数精度的选择及数学函数和内建函数之间的选择都要视应用程序的要求而定。

来看看GPT怎么说：

CUDA（Compute Unified Device Architecture）数学库是NVIDIA为其GPU计算平台提供的一组数学函数和库，用于高性能并行计算任务。这些库允许开发者在GPU上执行各种数学运算，包括向量和矩阵操作、快速傅立叶变换、随机数生成、常见数学函数等。以下是一些常见的CUDA数学库：

cuBLAS：CUDA基本线性代数子程序（cuBLAS）是用于线性代数计算的CUDA库。它提供了一系列高性能的矩阵运算，如矩阵乘法、矩阵-向量乘法、矩阵求逆和分解等。
cuSPARSE：cuSPARSE库用于稀疏矩阵操作。它支持各种稀疏矩阵格式，并提供了执行稀疏矩阵向量乘法等操作的函数。
cuFFT：CUDA快速傅立叶变换库（cuFFT）用于高效执行傅立叶变换操作，包括一维、二维和三维傅立叶变换。这对于信号处理、图像处理和其他频域分析任务非常有用。
cuRAND：cuRAND库用于生成随机数。它提供了各种随机数生成器，包括伪随机数生成器和真随机数生成器，以及各种分布函数。
Thrust：Thrust是一个模板库，用于高性能并行计算，包括各种常见的STL（标准模板库）风格的算法，如排序、归约、扫描等。Thrust库可以与CUDA一起使用，以执行高效的并行数据处理操作。
NPP（NVIDIA Performance Primitives）：NPP是一组高性能的图像和信号处理函数，用于图像处理、视频编解码、特征检测等应用。
CUTLASS：CUTLASS是一个通用矩阵库，可用于定义和执行各种矩阵操作。它旨在简化CUDA程序员的矩阵计算任务。

这些CUDA数学库提供了高性能的数学和数值计算功能，允许开发者充分利用GPU的并行计算能力。它们对于科学计算、深度学习、计算机视觉和其他需要大规模数据处理的领域非常有用。开发者可以通过包含适当的头文件和链接相应的库来使用这些库。

grace 1314

关注

0
点赞
踩
1

收藏

觉得还不错? 一键收藏
0
评论
CUDA编程（五）获得GPU加速的关键

CUDA程序获得高性能要求：1.数据传输比例所占时间小。2.核函数的算术强度较高。3.核函数中定义的线程数目较多。所以，在编写与优化CUDA程序时，一定要想方设法（主要是指仔细设计算法）做到以下几点：• 减少主机与设备之间的数据传输。• 提高核函数的算术强度。• 增大核函数的并行规模。
复制链接

扫一扫