CUDA编程（十一）CUDA流

grace 1314

已于 2023-11-02 10:08:26 修改

阅读量890

点赞数 2

文章标签： c++

于 2023-10-28 18:44:24 首次发布

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

版权

CUDA程序的并行层次主要有两个，一个是核函数内部的并行，一个是核函数外部的并行。我们之前讨论的都是核函数内部的并行。核函数外部的并行主要指：

1. 核函数计算与数据传输之间的并行。

2. 主机计算与数据传输之间的并行。

3. 不同的数据传输（回顾一下cudaMemcpy函数中的第4个参数）之间的并行。

4. 核函数计算与主机计算之间的并行。

5. 不同核函数之间的并行。

核函数计算，主机计算，数据传输，这里的核函数外部并行就是这三者的相互并行，不同的是核函数有好多可以并行，数据传输也是可以是并行。

这里再来看一下cudaMemcpy函数：

cudaError_t cudaMemcpy(void* dst, const void* src, size_t count, cudaMemcpyKind kind);

dst：指向目标内存的指针，可以是主机或设备内存。
src：指向源内存的指针，可以是主机或设备内存。
count：要传输的数据字节数。
kind：传输的类型，可以是以下值之一：
- cudaMemcpyHostToHost：主机到主机的传输。
- cudaMemcpyHostToDevice：主机到设备的传输。
- cudaMemcpyDeviceToHost：设备到主机的传输。
- cudaMemcpyDeviceToDevice：设备到设备的传输

第四个参数具有不同的数据传输类型，这种也可以做到并行。

一般来说，核函数外部的并行不是开发CUDA程序时考虑的重点。我们前面强调过，要获得较高的加速比，需要尽量减少主机与设备之间的数据传输及主机中的计算，尽量在设备中完成所有计算。如果做到了这一点，上述前4种核函数外部的并行就显得不那么重要了。另外，如果单个核函数的并行规模已经足够大，在同一个设备中同时运行多个核函数也不会带来太多性能提升，上述第五种核函数外部的并行也将不重要。不过，对有些应用，核函数外部的并行还是比较重要的。为了实现这种并行，需要合理地使用CUDA流（CUDA stream）。

1.页锁定主机内存

主机内存分配函数：

malloc() 标准c库函数

cudaHostAlloc() CUDA运行时提供自己独有的机制来分配主机内存

两个函数分配的内存之间的差异：

malloc()将分配标准的，可分页的（Pagable）主机内存，
cudaHostAlloc()将分配页锁定的主机内存（固定内存）

页锁定内存（Pinned Memory）是一种在操作系统中锁定的内存区域，通常用于与GPU设备之间的高性能数据传输，特别是在使用CUDA编程时。

1.不受虚拟内存管理的影响，通常，操作系统会根据需求将内存页面从主存（RAM）交换到磁盘以进行虚拟内存管理。这可能导致数据传输时的延迟和性能下降。页锁定内存不会被交换到磁盘，而是保存在物理内存，因此对内存访问速度更快，对于GPU设备之间的数据传输尤其有用。由于知道物理内存可通过DMA（直接访问）技术实现GPU和主机之间复制数据。

2.高带宽内存，页锁定内存通常与高带宽内存总线一起使用，以加速主机与GPU之间的数据传输。

3.减少数据传输延迟，页锁定内存允许异步数据传输，这意味着你可以在数据传输的同时进行计算，从而减少了数据传输的延迟。

4.适用于大规模数据，在处理大规模数据时，使用页锁定内存可以提供更高的性能，因为它减少了内存分页操作的开销。

在CUDA编程中，通常使用cudaMallocHost函数来分配页锁定内存。使用页锁定内存时，你可以通过cudaMemcpy等CUDA内存传输函数将数据直接从页锁定内存传输到GPU设备内存，以获得更高的性能和效率。

总结：将数据放在为页锁定内存上，对数据的访问更高效，数据不可分，就是一个整体，不在具有虚拟内存的功能。

2.CUDA流

一个CUDA流指的是由主机发出的在一个设备中执行的CUDA操作序列（即和CUDA有关的操作，如主机－设备数据传输和核函数执行）。除了主机端发出的流，还有设备端发出的流。

一个CUDA流中各个操作的次序是由主机控制的，按照主机发布的次序执行。然而，来自于两个不同CUDA流中的操作不一定按照某个次序执行，而有可能并发或交错地执行。

任何CUDA操作都存在于某个CUDA流中，要么是默认流（defaultstream），也称为空流（null stream），要么是明确指定的非空流。在之前的章节中，我们没有明确地指定CUDA流，那里所有的CUDA操作都是在默认的空流中执行的。

为了实现不同CUDA流之间的并发，主机在向某个CUDA流中发布一系列命令之后必须马上获得程序的控制权，不用等待该CUDA流中的命令在设备中执行完毕。这样，就可以通过主机产生多个相互独立的CUDA流。

支持设别重叠功能的GPU能在执行一个CUDA C核函数的同时，还能在设备与主机之间执行复制操作。可以使用多个流来实现这种计算与数据传输的重叠。

如何使用流：

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

#define N (1024*1024)
#define FULL_DATA_SIZE (N*20)


__global__ void kernel(int* a, int* b, int* c) {
	int idx = threadIdx.x + blockIdx.x * blockDim.x;
	if (idx < N) {
		int idx1 = (idx + 1) % 256;
		//因为一个块的大小为256，所以这里除以256，方便实现并行
		//因为并行是所有块同时运行线程也是
		//而idx是全局索引
		int idx2 = (idx + 2) % 256;
		float as = (a[idx] + a[idx1] + a[idx2]) / 3.0f;
		float bs = (b[idx] + b[idx1] + b[idx2]) / 3.0f;
		//这里as是将当前线程索引值以及后两个索引值做平均

		c[idx] = (as + bs) / 2;

		//因为这里只是从a，b中读出数据，运算完后统一放进c里面，所以没有读写冲突
	}
}

int main(void) {
	cudaDeviceProp prop;
	int whichDevice;
	CHECK(cudaGetDevice(&whichDevice));
	CHECK(cudaGetDeviceProperties(&prop, whichDevice));
	//判断，支持设备重叠功能
	if (!prop.deviceOverlap) {
		printf("设备不支持重叠，不能为cuda流加速\n");
		return 0;
	}
	//支持设别重叠功能的GPU能在执行一个CUDA C核函数的同时，还能在设备与主机之间执行复制操作。

	//创建事件和流，以及主机变量，设备变量
	cudaStream_t stream;
	int* host_a, * host_b, * host_c;
	int* dev_a, * dev_b, * dev_c;

	//计时模块初始化
	cudaEvent_t start, stop;
	float elapsedtime;
	CHECK(cudaEventCreate(&start));
	CHECK(cudaEventCreate(&stop));
	
	//初始化流
	CHECK(cudaStreamCreate(&stream));

	//分配设备内存，只申请了1/20数据大小的内存
	CHECK(cudaMalloc((void**)&dev_a, N * sizeof(int)));
	CHECK(cudaMalloc((void**)&dev_b, N * sizeof(int)));
	CHECK(cudaMalloc((void**)&dev_c, N * sizeof(int)));

	// 在这里申请固定内存不仅仅是为了让复制操作执行得更快
	// 要以异步的方式在主机和设备之间复制数据必须是固定内存
	// 申请内存大小为数据大小
	CHECK(cudaHostAlloc((void**)&host_a, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault));
	CHECK(cudaHostAlloc((void**)&host_b, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault));
	CHECK(cudaHostAlloc((void**)&host_c, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault));

	//填充申请的缓冲区host_a,host_b
	for (int i = 0; i < FULL_DATA_SIZE; i++) {
		host_a[i] = rand();
		host_b[i] = rand();
	}
	//开始计时
	CHECK(cudaEventRecord(start, 0));

	//循环整个数据，分成更小的块
	//我们不将输入缓冲区整体复制到GPU，而是将输入缓冲区划分成更小的块（分成20块），并在每个块上执行一个包含三个步骤的过程：
	//1.将一部分输入缓冲区复制到GPU ；
	// 2.在这部分缓冲区上运行核函数；
	// 3.然后将一部分输入缓冲区复制到GPU

	for (int i = 0; i < FULL_DATA_SIZE; i += N) {
		//复制页锁定内存数据以异步方式复制到设备  async 异步操作
		CHECK(cudaMemcpyAsync(dev_a, host_a + i, N * sizeof(int), cudaMemcpyHostToDevice, stream));
		//host_a + i隔20个数据复制一个，将host的数据复制到device

		CHECK(cudaMemcpyAsync(dev_b, host_b + i, N * sizeof(int), cudaMemcpyHostToDevice, stream));

		//核函数带有流参数
		//刚好N个线程有N个数据，线程不需要多次工作
		//gridsize N/256
		//blocksize 256，
		//shared_size,0,共享内存大小，这块为0说明没有使用共享内存
		kernel << <N/256,256,0,stream >> >(dev_a,dev_b,dev_c) ;

		//将数据从设备复制到锁定内存
		CHECK(cudaMemcpyAsync(host_c + i, dev_c, N * sizeof(int), cudaMemcpyDeviceToDevice, stream));
	}

	//复制从页锁定内存结果块到缓冲
	CHECK(cudaStreamSynchronize(stream));
	
	//计时模块
	CHECK(cudaEventRecord(stop, 0));
	CHECK(cudaEventSynchronize(stop));
	CHECK(cudaEventElapsedTime(&elapsedtime,start,stop));
	printf("Time taken: %3.1f ms\n",elapsedtime); 
	CHECK(cudaEventDestroy(start));
	CHECK(cudaEventDestroy(stop));

	//释放内存
	CHECK(cudaFreeHost(host_a));
	CHECK(cudaFreeHost(host_b));
	CHECK(cudaFreeHost(host_c));
	CHECK(cudaFree(dev_a));
	CHECK(cudaFree(dev_b));
	CHECK(cudaFree(dev_c));
	//清理流
	CHECK(cudaStreamDestroy(stream));

	return 0;
}

上面代码已经附上注释，可以自行理解。

这里说一下，异步与并行。

异步：异步（Asynchronous）：异步是指任务的执行是不按照固定的顺序，而是根据事件或任务的准备情况进行的。异步任务可以在后台执行，而不会阻塞主线程。异步编程旨在提高程序的响应性，允许同时执行多个任务。异步任务可以是并行执行的，但不一定是。例如，在单线程的环境中，异步任务可能是顺序执行的，但不会阻塞主线程。

并行：并行是指多个任务在同一时刻同时执行，它们可以在多个处理器、多个线程或多个计算单元上并发运行。并行任务旨在提高计算性能和资源利用率。并行任务通常要求具有多个执行单元，例如多核CPU或GPU。并行任务是同步的，因为它们需要协同工作以完成某个任务，但不一定是异步的。

总结，异步可以导致并行，但异步更侧重于任务的执行方式和顺序，而并行则侧重于同时执行多个任务。

当for循环结束时，队列中应该包含了很多等待GPU执行的工作。如果想要确保GPU只能执行完了计算和内存复制等操作。那么就需要将GPU与主机同步。也就是说主机在继续执行之前要先等待GPU完成。调用cudaStreamSynchronize()并指定想要等待的流。

主机与设备之间复制数据

cudaMemcpy()同步方式执行：意味着，当函数返回时，复制操作已经完成，并且在输出缓冲区包含了复制进去的内容。
新函数cudaMemcpyAsync()异步方式执行：与同步方式相反，在调用该函数时，只是放置一个请求，表示在流中执行一次内存复制操作，这个流是通过函数stream来指定的。当函数返回时，我们无法确保复制操作是否已经启动，更无法保证它是否已经结束。我们能够保证的是，复制操作肯定会当下一个被放入流中的操作之前执行。
任何一个传递给cudaMemcpyAsync()的主机内存指针都必须已经通过cudaHostAlloc()分配好内存。你只能已异步方式对固定内存进行复制操作。

异步数据传输：

异步数据传输指的是在主机（通常是CPU）和设备（通常是GPU）之间进行数据传输，而不会阻塞主机或设备上的其他操作。这意味着数据传输操作可以与其他计算操作同时进行，而不需要等待数据传输完成。
在异步数据传输中，数据可以被拷贝到设备，也可以从设备返回到主机，而不会阻止主机或设备上的其他操作执行。
异步数据传输通常使用CUDA流（cudaStream）来实现，允许将数据传输操作添加到流中，以便在流的上下文中异步执行

cuda异步传输函数

cudaMemcpyAsync：用于在主机和设备之间异步传输数据。您可以使用这个函数来复制数据从主机到设备或从设备到主机。它允许您指定目标流，以确定传输的异步性。

cudaMemcpyAsync(dest_device, src_host, size, cudaMemcpyHostToDevice, stream);

cudaMemsetAsync：用于在设备上异步设置内存。它允许您在目标流上异步执行内存设置操作。

cudaMemsetAsync(device_ptr, value, size, stream);

cudaMemcpyPeerAsync：用于在两个不同的GPU设备之间异步传输数据。您可以在目标流上执行跨设备的数据传输。

cudaMemcpyPeerAsync(dest_device, src_device, src_ptr, size, stream);

cudaHostRegister 和 cudaHostUnregister：这些函数允许您在主机端注册内存以在设备和主机之间进行异步数据传输。注册后，您可以使用异步传输函数在主机和设备之间传输数据。

cudaStreamSynchronize：虽然不是传输函数，但这个函数允许您等待一个特定流中的操作完成。您可以使用它来确保在继续执行主机端代码之前等待流中的异步传输操作完成。

带有流参数的核函数，此时核函数的调用是异步的。

上述代码中，将host数据，host_a,host_b,从lock内存复制,执行核函数，将结果搬回host_c,这3件事是异步进行的，意思就是给你一个指令你做你的，我继续向下执行代码，再把分派任务给其他人，你们接到任务做就行了。同一个流中，操作通常按照它们被提交到流中的顺序执行。也就是说，一个操作必须在前一个操作完成之后才会开始执行。但是这里使用的是cudaMemcpyAsync异步数据传输，核函数本身执行就是异步的，所以这里3件事都会异步做。

3.使用多个流

如何使用多个流对上面的操作进行改进呢：

这里是一些思路：

分块计算
内存复制和核函数执行的重叠

上图中，第0个流执行：核函数时，在第1个流中执行：输入缓冲区复制到GPU......
在任何支持内存复制和核函数的执行相互重叠的设备上，当使用多个流是，应用程序的整体性能都会提升。

当涉及到多个CUDA流时，执行顺序可以是不固定的。不同的CUDA流之间可以并行执行。开发者可以使用多个流来管理不同的任务，并充分利用GPU的并行性能。虽然流内的执行顺序相对固定，但多个流之间的并行执行为高性能计算提供了重要的支持。

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

#define N (1024*1024)
#define FULL_DATA_SIZE (N*20)


__global__ void kernel(int* a, int* b, int* c) {
	int idx = threadIdx.x + blockIdx.x * blockDim.x;
	if (idx < N) {
		int idx1 = (idx + 1) % 256;
		//因为一个块的大小为256，所以这里除以256，方便实现并行
		//因为并行是所有块同时运行线程也是
		//而idx是全局索引
		int idx2 = (idx + 2) % 256;
		float as = (a[idx] + a[idx1] + a[idx2]) / 3.0f;
		float bs = (b[idx] + b[idx1] + b[idx2]) / 3.0f;
		//这里as是将当前线程索引值以及后两个索引值做平均

		c[idx] = (as + bs) / 2;

		//因为这里只是从a，b中读出数据，运算完后统一放进c里面，所以没有读写冲突
	}
}

int main(void) {
	cudaDeviceProp prop;
	int whichDevice;
	CHECK(cudaGetDevice(&whichDevice));
	CHECK(cudaGetDeviceProperties(&prop, whichDevice));
	//判断，支持设备重叠功能
	if (!prop.deviceOverlap) {
		printf("设备不支持重叠，不能为cuda流加速\n");
		return 0;
	}
	//支持设别重叠功能的GPU能在执行一个CUDA C核函数的同时，还能在设备与主机之间执行复制操作。

	//创建事件和流，以及主机变量，设备变量
	cudaStream_t stream0,stream1;
	int* host_a, * host_b, * host_c;
	int* dev_a0, * dev_b0, * dev_c0;
	int* dev_a1, * dev_b1, * dev_c1;
	//计时模块初始化
	cudaEvent_t start, stop;
	float elapsedtime;
	CHECK(cudaEventCreate(&start));
	CHECK(cudaEventCreate(&stop));

	//初始化流
	CHECK(cudaStreamCreate(&stream0));
	CHECK(cudaStreamCreate(&stream1));

	//分配设备内存，只申请了1/20数据大小的内存
	CHECK(cudaMalloc((void**)&dev_a0, N * sizeof(int)));
	CHECK(cudaMalloc((void**)&dev_b0, N * sizeof(int)));
	CHECK(cudaMalloc((void**)&dev_c0, N * sizeof(int)));

	CHECK(cudaMalloc((void**)&dev_a1, N * sizeof(int)));
	CHECK(cudaMalloc((void**)&dev_b1, N * sizeof(int)));
	CHECK(cudaMalloc((void**)&dev_c1, N * sizeof(int)));

	// 在这里申请固定内存不仅仅是为了让复制操作执行得更快
	// 要以异步的方式在主机和设备之间复制数据必须是固定内存
	// 申请内存大小为数据大小
	CHECK(cudaHostAlloc((void**)&host_a, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault));
	CHECK(cudaHostAlloc((void**)&host_b, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault));
	CHECK(cudaHostAlloc((void**)&host_c, FULL_DATA_SIZE * sizeof(int), cudaHostAllocDefault));

	//填充申请的缓冲区host_a,host_b
	for (int i = 0; i < FULL_DATA_SIZE; i++) {
		host_a[i] = rand();
		host_b[i] = rand();
	}
	//开始计时
	CHECK(cudaEventRecord(start, 0));

	//循环整个数据，分成更小的块
	//我们不将输入缓冲区整体复制到GPU，而是将输入缓冲区划分成更小的块（分成20块），并在每个块上执行一个包含三个步骤的过程：
	//1.将一部分输入缓冲区复制到GPU ；
	// 2.在这部分缓冲区上运行核函数；
	// 3.然后将一部分输入缓冲区复制到GPU

	for (int i = 0; i < FULL_DATA_SIZE; i += N * 2) {
		//stream0
		//复制页锁定内存数据以异步方式复制到设备  async 异步操作
		CHECK(cudaMemcpyAsync(dev_a0, host_a + i, N * sizeof(int), cudaMemcpyHostToDevice, stream0));
		//host_a + i隔20个数据复制一个，将host的数据复制到device

		CHECK(cudaMemcpyAsync(dev_b0, host_b + i, N * sizeof(int), cudaMemcpyHostToDevice, stream0));

		//核函数带有流参数
		//刚好N个线程有N个数据，线程不需要多次工作
		//gridsize N/256
		//blocksize 256，
		//shared_size,0,共享内存大小，这块为0说明没有使用共享内存
		kernel << <N / 256, 256, 0, stream0 >> > (dev_a0, dev_b0, dev_c0);

		//将数据从设备复制到锁定内存
		CHECK(cudaMemcpyAsync(host_c + i, dev_c0, N * sizeof(int), cudaMemcpyDeviceToDevice, stream0));


		//stream1
		//复制页锁定内存数据以异步方式复制到设备  async 异步操作
		CHECK(cudaMemcpyAsync(dev_a1, host_a + i + N, N * sizeof(int), cudaMemcpyHostToDevice, stream1));
		//host_a + i隔20个数据复制一个，将host的数据复制到device

		CHECK(cudaMemcpyAsync(dev_b1, host_b + i + N, N * sizeof(int), cudaMemcpyHostToDevice, stream1));

		//核函数带有流参数
		//刚好N个线程有N个数据，线程不需要多次工作
		//gridsize N/256
		//blocksize 256，
		//shared_size,0,共享内存大小，这块为0说明没有使用共享内存
		kernel << <N / 256, 256, 0, stream1 >> > (dev_a1, dev_b1, dev_c1);

		//将数据从设备复制到锁定内存
		CHECK(cudaMemcpyAsync(host_c + i + N, dev_c1, N * sizeof(int), cudaMemcpyDeviceToDevice, stream1));


	}

	//
	CHECK(cudaStreamSynchronize(stream0));
	CHECK(cudaStreamSynchronize(stream1));

	//计时模块
	CHECK(cudaEventRecord(stop, 0));
	CHECK(cudaEventSynchronize(stop));
	CHECK(cudaEventElapsedTime(&elapsedtime, start, stop));
	printf("Time taken: %3.1f ms\n", elapsedtime);
	CHECK(cudaEventDestroy(start));
	CHECK(cudaEventDestroy(stop));

	//释放内存
	CHECK(cudaFreeHost(host_a));
	CHECK(cudaFreeHost(host_b));
	CHECK(cudaFreeHost(host_c));
	CHECK(cudaFree(dev_a0));
	CHECK(cudaFree(dev_b0));
	CHECK(cudaFree(dev_c0));
	CHECK(cudaFree(dev_a1));
	CHECK(cudaFree(dev_b1));
	CHECK(cudaFree(dev_c1));
	//清理流
	CHECK(cudaStreamDestroy(stream0));
	CHECK(cudaStreamDestroy(stream1));

	return 0;
}

在中间的for循环中 i 的增量变为N*2，即原来的两倍，意味着在数据传输与数据

cuda流0中dev_a0,dev_b0,dev_c0,取数据是取其中一半，cuda流1中dev_a1,dev_b1,dev_c1取剩下的一半，分别放进kernel中进行计算，最后复制会主机时，也是取一半。

3.主机与设备计算重叠

为了实现不同 CUDA 流之间的并发，主机在向某个 CUDA 流中发布一系列命令之后必须马上获得程序的控制权，不用等待该 CUDA 流中的命令在设备中执行完毕。这样，就可以通过主机产生多个相互独立的 CUDA 流。（本质是一种多线程编程问题）

为了检查一个 CUDA 流中的所有操作是否都在设备中执行完毕，

cudaError_t cudaStreamSynchronize(cudaStream_t stream);

cudaError_t cudaStreamQuery(cudaStream_t stream);

函数 cudaStreamSynchronize 会强制阻塞主机，直到 CUDA 流 stream 中的所有操作都执行完毕。

函数 cudaStreamQuery 不会阻塞主机，只是检查 CUDA 流 stream 中的所有操作是否都执行完毕。若是，返回 cudaSuccess，否则返回 cudaErrorNotReady。

虽然同一个 CUDA 流中的所有 CUDA 操作都是顺序执行的，但依然可以在默认流中重叠主机和设备的计算。

比如数组相加：

cudaMemcpy(d_x, h_x, M, cudaMemcpyHostToDevice);
cudaMemcpy(d_y, h_y, M, cudaMemcpyHostToDevice);
sum<<<grid_size, block_size>>>(d_x, d_y, d_z, N);
cudaMemcpy(h_z, d_z, M, cudaMemcpyDeviceToHost);

从设备的角度来看，以上 4 个 CUDA 操作语句将在默认的 CUDA 流中按代码出现的顺序依次执行。

从主机的角度来看，数据传输是同步的（synchronous），或者说是阻塞的（blocking），意思是主机发出命令

cudaMemcpy(d_x, h_x, M, cudaMemcpyHostToDevice);

之后，会等待该命令执行完毕，再往前走。

在进行数据传输时，主机是闲置的，不能进行其他操作。

不同的是，核函数的启动是异步的（asynchronous），或者说是非阻塞的（non-blocking），意思是主机发出命令

sum>>(d_x, d_y, d_z, N);

之后，不会等待该命令执行完毕，而会立刻得到程序的控制权。主机紧接着会发出从设备到主机传输数据的命令

cudaMemcpy(h_z, d_z, M, cudaMemcpyDeviceToHost);

然而，该命令不会被立即执行，因为这是默认流中的 CUDA 操作，必须等待前一个 CUDA 操作（即核函数的调用）执行完毕才会开始执行。

所以是命令是先发了，执行的时候还是得等待。

主机在发出核函数调用的命令之后，会立刻发出下一个命令。在上面的例子中，下一个命令是进行数据传输，但从设备的角度来看必须等待核函数执行完毕。如果下一个命令是主机中的某个计算任务，那么主机就会在设备执行核函数的同时去进行一些计算。这样，主机和设备就可以同时进行计算。设备完全不知道在它执行核函数时，主机偷偷地做了些计算（有点搞小三的意思）。

这种重叠主机和设备计算，指的是主机发给设备一个命令让设备去干，而主机不是在等他干完，而是偷偷干些其他的。

当主机和设备函数的计算量相当时，将主机函数放在设备函数之后可以达到主机函数与设备函数并发执行的效果，从而有效地隐藏主机函数的执行时间，提升整个程序的性能。在主机函数与设备函数的计算时间相当的情况下，可以获得接近 2 倍的加速比。所以，当一个主机函数与一个设备函数的计算相互独立时，应该将主机函数的调用放置在核函数的调用之后，而不是之前。

4.多流与多核函数执行

虽然在一个默认流中就可以实现主机计算和设备计算的并行，但是要实现多个核函数之间的并行必须使用多个 CUDA 流。这是因为，同一个 CUDA 流中的 CUDA 操作在设备中是顺序执行的，故同一个 CUDA 流中的核函数也必须在设备中顺序执行，虽然主机在发出每一个核函数调用的命令后都立刻重新获得程序控制权。

在主机端，CUDA操作的提交和执行通常是异步的。这意味着当主机代码提交CUDA操作到流中时，主机代码可以继续执行而不必等待CUDA操作完成。因此，主机端代码不会按照CUDA操作的提交顺序来顺序执行，而是可以继续执行其他操作。

核函数执行配置中的流参数：

my_kernel<<<N_grid, N_block>>>(函数参数);
my_kernel<<<N_grid, N_block, N_shared>>>(函数参数);
my_kernel<<<N_grid, N_block, N_shared, stream_id>>>(函数参数);

3中核函数的调用方式。stream_id是cuda流的编号。

如果用第一种调用方式，说明核函数没有使用动态共享内存，而且在默认流中执行；如果用第二种调用方式，说明核函数在默认流中执行，但使用了 N_shared 字节的动态共享内存；如果用第三种调用方式，则说明核函数在编号为stream_id 的 CUDA 流中执行，而且使用了 N_shared 字节的动态共享内存。在使用非空流但不使用动态共享内存的情况下，必须使用上述第三种调用方式，并将 N_shared 设置为零：

my_kernel<<<N_grid, N_block, 0, stream_id>>>(函数参数); // 正确
不能用如下调用方式：
my_kernel<<<N_grid, N_block, stream_id>>>(函数参数); // 错误

制约加速比的因素，GPU 的计算资源，单个 GPU 中能够并发执行的核函数个数的上限

并不是cuda流越多越好，流数到一定程度就限制了，与GPU架构有关。

用非默认 CUDA 流重叠核函数的执行与数据传递：

不可分页主机内存与异步的数据传输函数：要实现核函数执行与数据传输的并发（重叠），必须让这两个操作处于不同的非默认流，而且数据传输必须使用 cudaMemcpy 函数的异步版本，即 cudaMemcpyAsync 函数。异步传输由 GPU 中的 DMA（direct memory access）直接实现，不需要主机参与。如果用同步的数据传输函数，主机在向一个流发出数据传输的命令后，将无法立刻获得控制权，必须等待数据传输完毕。也就是说，主机无法同时去另一个流调用核函数。这样核函数与数据传输的重叠也就无法实现。

在使用异步的数据传输函数时，需要将主机内存定义为不可分页内存（non-pageable memory）或者固定内存（pinned memory）。不可分页内存是相对于可分页内存（pageable memory）的。操作系统有权在一个程序运行期间改变程序中使用的可分页主机内存的物理地址。相反，若主机中的内存声明为不可分页内存，则在程序运行期间，其物理地址将保不变。如果将可分页内存传给 cudaMemcpyAsync 函数，则会导致同步传输，达不到重叠核函数执行与数据传输的效果。主机内存为可分页内存时，数据传输过程在使用 GPU 中的 DMA 之前必须先将数据从可分页内存移动到不可分页内存，从而必须与主机同步。主机无法在发出数据传输的命令后立刻获得程序的控制权，从而无法实现不同 CUDA 流之间的并发。

不可分页主机内存的分配可以由以下两个 CUDA 运行时 API 函数中的任何一个实现： cudaError_t cudaMallocHost(void** ptr, size_t size); cudaError_t cudaHostAlloc(void** ptr, size_t size, size_t flags);注意，第二个函数的名字中没有字母 M。若函数 cudaHostAlloc 的第三个参数取默认值 cudaHostAllocDefault，则以上两个函数完全等价。

由以上函数分配的主机内存必须由如下函数释放：

cudaError_t cudaFreeHost(void* ptr); 如果不小心用了 free 函数释放不可分页主机内存，会出现运行错误。

假如在一段 CUDA 程序中，我们需要先从主机向设备传输一定数量的数据（我们将此 CUDA 操作简称为 H2D），然后在 GPU 中使用所传输的数据做一些计算（我们将此 CUDA 操作简称为 KER，意为核函数执行），最后将一些数据从设备传输至主机（我们将此 CUDA 操作简称为 D2H）。

H2D-》KER-》D2H；

如果仅使用一个 CUDA 流（如默认流），那么以上 3 个操作在设备中一定是顺序的：

Stream 0：H2D -> KER -> D2H

如果简单地将以上 3 个 CUDA 操作放入 3 个不同的流，相比仅使用一个 CUDA 流的情形依然不能得到加速，因为以上 3 个操作在逻辑上是有先后次序的。如果使用 3 个流，其执行流程可以理解如下：

Stream 1：H2D

Stream 2： -> KER

Stream 3： -> D2H

因为该方案不能带来性能提升，我们不讨论如何在 3 个流中保证这种执行次序。

要利用多个流提升性能，就必须创造出在逻辑上可以并发执行的 CUDA 操作。

一个方法是将以上 3 个 CUDA 操作都分成若干等份，然后在每个流中发布一个 CUDA 操作序列。例如，使用两个流时，我们将以上 3 个 CUDA 操作都分成两等份。在理想情况下，它们的执行流程可以如下：

Stream 1：H2D -> KER -> D2H

Stream 2： H2D -> KER -> D2H

这里的每个 CUDA 操作所处理的数据量只有使用一个 CUDA 流时的一半。我们注意到，两个流中的 H2D 操作不能并发地执行（受硬件资源的限制），但第二个流的 H2D 操作可以和第一个流的 KER 操作并发地执行，第二个流的 KER 操作也可以和第一个流的 D2H 操作并发地执行。如果 H2D、KER、和 D2H 这 3 个 CUDA 操作的执行时间都相同，那么就能有效地隐藏一个 CUDA 流中两个 CUDA 操作的执行时间，使得总的执行效率相比使用单个 CUDA 流的情形提升到 6/4 = 1.5 倍。

主机中数据存放的内存是不可分页的。如果为它们分配可分页内存，程序依然能够编译、运行，但计算时间会随 CUDA 流数目的增多而单调递增。这说明，将可分页内存变量传入异步传输函数时，异步的 cudaMemcpyAsync 函数将退化为同步的 cudaMemcpy 函数，导致同步传输的行为，从而达不到重叠核函数执行与数据传输的效果。

以上。总结一下（gpt）：

CUDA流是用于管理并发执行的一种机制，允许在NVIDIA GPU上并行执行多个任务，从而提高性能。以下是关于CUDA流的一些总结要点：

并行性和异步执行：CUDA流允许在同一个GPU上并行执行多个任务，而不需要等待前一个任务完成。这提高了GPU的利用率和性能，因为计算和数据传输可以重叠执行。
流的创建和销毁：您可以使用cudaStreamCreate函数创建CUDA流，并使用cudaStreamDestroy函数销毁它们。每个流代表一个独立的任务队列。
流的同步：使用cudaStreamSynchronize函数可以在主机端同步等待特定流中的操作完成。这对于确保在主机代码中正确处理CUDA操作非常重要。
流的优先级：CUDA流可以分配不同的优先级，以确保某些任务在其他任务之前执行。这可以通过cudaStreamCreateWithPriority函数来设置流的优先级。
多流管理：在一个应用程序中，可以创建多个CUDA流，每个流用于管理不同的任务。这有助于提高应用程序的并发性和性能。
流内执行顺序：在同一个CUDA流中，操作通常按照它们被提交到流中的顺序执行。这确保了在流内的操作按顺序执行。
流间执行顺序：在不同的CUDA流之间，操作的执行顺序是不固定的，它们可以并行执行。执行顺序可能受到GPU硬件资源和任务优先级等因素的影响。
异步传输：使用CUDA流，您可以执行异步数据传输操作，如cudaMemcpyAsync和cudaMemcpyPeerAsync，以在主机和设备之间传输数据并充分利用GPU并行性。
流的应用：CUDA流通常用于加速大规模并行计算任务，如深度学习、科学计算和图形渲染。它们有助于提高GPU的利用率和性能。