解读CUDA C Programming Guide 第三章第2节（共6节）（3/5）

最新推荐文章于 2023-04-05 19:57:04 发布

AliceWanderAI

最新推荐文章于 2023-04-05 19:57:04 发布

阅读量331

点赞数

本文链接：https://blog.csdn.net/NXHYD/article/details/104882843

版权

本系列为《解读CUDA C Programming Guide》.

本书旨在介绍进行CUDA并行优化的C编程指导。共5章，内容分别是：

Introduction
Programming Model
Programming Interface
Hardware Implementation
Performance Guidelines

本文简单解读第三章：Programming Interface.

本章主要内容包括：

Compilation with NVCC
CUDA C Runtime
Versioning and Compatibility
Compute Modes
Mode Switches
Tesla Compute Cluster Mode for Windows

CUDA C Runtime

Runtime被应用于cudart库，其通过cudart.lib和libcudart.a静态链接到应用程序，或通过cudart.dll和libcudart.so动态地链接到应用程序。需要cudart.dll或cudart.so进行动态链接的应用程序通常会将它们作为应用程序安装包的一部分包含在内。只有在链接到CUDA Runtime时相同实例的组件之间传递CUDA Runtime符号的地址才是安全的。其所有入口点都以cuda开头。

Initialization

运行时没有显式的初始化函数；它会在首次调用运行时功能时进行初始化。在定时运行时函数调用以及解释从第一次调用到运行时的错误代码时，需要牢记这一点。

在初始化期间，运行时会为系统中的每个设备创建CUDA context. 这是 Initialization阶段所解决的主要问题。

那什么是CUDA context呢?

cuda context 非常重要，它作为一个容器，管理了所有对象的生命周期，大多数的CUDA函数调用需要context。这些对象如下：

所有分配内存
Modules，类似于动态链接库，以.cubin和.ptx结尾【在jcuda中要使用】
CUDA streams，管理执行单元的并发性
CUDA events
texture和surface引用
kernel里面使用到的本地内存（设备内存）
用于调试、分析和同步的内部资源
用于分页复制的固定缓冲区

此context是此设备的主要context，并且在应用程序的所有主机线程之间共享。作为此context创建的一部分，如有必要，设备代码会及时进行编译并加载到设备内存中。这一切都是在后台进行的，运行时不会向应用程序公开主要context.

当主机线程调用cudaDeviceReset（）时，这会破坏有主机线程运行的设备的主要context。任何主机线程（有当前设备）进行的下一个运行时函数调用，都将为此设备创建一个新的主要context.

Device Memory

如Heterogeneous Programming异构编程中所提到的，CUDA编程模型假定一个由主机和设备组成的系统，每个主机和设备都有各自独立的内存。内核在设备内存之外运行，因此运行时提供了分配，取消分配和复制设备内存以及在主机内存和设备内存之间传输数据的功能。

设备内存可以分配为线性内存或CUDA阵列。

线性内存在于设备的40位地址空间中，因此单独分配的实体可以通过指针相互引用。

CUDA阵列是为纹理获取而优化的不透明的内存布局。

线性内存通常使用cudaMalloc（）分配，并使用cudaFree（）释放，并且主机内存和设备内存之间的数据传输通常使用cudaMemcpy（）完成。在内核的向量加法代码示例中，需要将向量从主机存储器复制到设备存储器：

// Device code
__global__ void VecAdd(float* A, float* B, float* C, int N)
{
    int i = blockDim.x * blockIdx.x + threadIdx.x;
    if (i < N)
    C[i] = A[i] + B[i];
}

// Host code
int main()
{
    int N = ...;
    size_t size = N * sizeof(float);
    // Allocate input vectors h_A and h_B in host memory
    float* h_A = (float*)malloc(size);
    float* h_B = (float*)malloc(size);
    // Initialize input vectors
    ...
    // Allocate vectors in device memory
    float* d_A;
    cudaMalloc(&d_A, size);
    float* d_B;
    cudaMalloc(&d_B, size);
    float* d_C;
    cudaMalloc(&d_C, size);
    // Copy vectors from host memory to device memory
    cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);
    // Invoke kernel
    int threadsPerBlock = 256;
    int blocksPerGrid =
    (N + threadsPerBlock - 1) / threadsPerBlock;
    VecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);
    // Copy result from device memory to host memory
    // h_C contains the result in host memory
    cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);
    // Free device memory
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);
    // Free host memory
    ...
}

线性内存也可以通过cudaMallocPitch（）和cudaMalloc3D（）进行分配。建议将这些功能用于2D或3D阵列的分配，因为它确保适当填充分配以满足设备内存访问中描述的对齐要求，从而确保在访问行地址或在2D阵列与其他区域之间执行复制时确保最佳性能。设备内存（使用cudaMemcpy2D（）和cudaMemcpy3D（）函数）。返回的音高（或步幅）必须用于访问数组元素。

以下代码示例分配了一个宽度x高度的2D浮点值数组，并显示了如何在设备代码中循环遍历数组元素：

// Host code
int width = 64, height = 64;
float* devPtr;
size_t pitch;
cudaMallocPitch(&devPtr, &pitch,
 width * sizeof(float), height);
 
MyKernel<<<100, 512>>>(devPtr, pitch, width, height);
// Device code
__global__ void MyKernel(float* devPtr,
 size_t pitch, int width, int height)
{
	for (int r = 0; r < height; ++r) 
	{
		float* row = (float*)((char*)devPtr + r * pitch);
		for (int c = 0; c < width; ++c) 
		{
			float element = row[c];
		}
	}
}

以下代码示例分配了一个x高度x深度3d浮点值数组，并显示了如何在设备代码中循环遍历数组元素：

// Host code
int width = 64, height = 64, depth = 64;
cudaExtent extent = make_cudaExtent(width * sizeof(float),
 height, depth);
cudaPitchedPtr devPitchedPtr;

cudaMalloc3D(&devPitchedPtr, extent);
MyKernel<<<100, 512>>>(devPitchedPtr, width, height, depth);
// Device code
__global__ void MyKernel(cudaPitchedPtr devPitchedPtr,
 int width, int height, int depth)
{
	 char* devPtr = devPitchedPtr.ptr;
	 size_t pitch = devPitchedPtr.pitch;
	 size_t slicePitch = pitch * height;
	 for (int z = 0; z < depth; ++z) 
	 {
		 char* slice = devPtr + z * slicePitch;
		 for (int y = 0; y < height; ++y) 
		 {
			 float* row = (float*)(slice + y * pitch);
			 for (int x = 0; x < width; ++x) 
			 {
				float element = row[x];
			 }
		 }
	 }
}

参考手册列出了用于在使用cudaMalloc（）分配的线性内存和使用cudaMalloc（）分配的线性内存之间复制内存的所有各种函数. cudaMallocPitch（）或cudaMalloc3D（），CUDA数组以及为在全局或恒定内存空间中声明的变量分配的内存。

以下代码示例说明了通过运行时API访问全局变量的各种方法：

__constant__ float constData[256];
float data[256];
cudaMemcpyToSymbol(constData, data, sizeof(data));
cudaMemcpyFromSymbol(data, constData, sizeof(data));
__device__ float devData;
float value = 3.14f;
cudaMemcpyToSymbol(devData, &value, sizeof(float));
__device__ float* devPointer;
float* ptr;
cudaMalloc(&ptr, 256 * sizeof(float));
cudaMemcpyToSymbol(devPointer, &ptr, sizeof(ptr));

cudaGetSymbolAddress（）用于检索指向为全局内存空间中声明的变量分配的内存的地址。分配的内存大小通过cudaGetSymbolSize（）获得。

Shared Memory

共享内存使用__shared__内存空间说明符。

如线程层次结构中提到并在共享内存中详细介绍的那样，共享内存预计比全局内存快得多。因此，如以下矩阵乘法示例所示，应利用任何机会用共享内存访问替换全局内存访问。

下面的代码示例是矩阵乘法的直接实现，没有利用共享内存。每个线程读取A的一行和B的一列，并计算C的对应元素，如图9所示。因此，A是从全局内存中读取B.width倍，而B则是A.height倍。

// Matrices are stored in row-major order:
// M(row, col) = *(M.elements + row * M.width + col)
typedef struct {
	 int width;
	 int height;
	 float* elements;
} Matrix;

// Thread block size
#define BLOCK_SIZE 16

// Forward declaration of the matrix multiplication kernel
__global__ void MatMulKernel(const Matrix, const Matrix, Matrix);

// Matrix multiplication - Host code
// Matrix dimensions are assumed to be multiples of BLOCK_SIZE
void MatMul(const Matrix A, const Matrix B, Matrix C)
{
	 // Load A and B to device memory
	 Matrix d_A;
	 d_A.width = A.width; d_A.height = A.height;
	 size_t size = A.width * A.height * sizeof(float);
	 cudaMalloc(&d_A.elements, size);
	 cudaMemcpy(d_A.elements, A.elements, size,
	 cudaMemcpyHostToDevice);
	 Matrix d_B;
	 d_B.width = B.width; d_B.height = B.height;
	 size = B.width * B.height * sizeof(float);
	 cudaMalloc(&d_B.elements, size);
	 cudaMemcpy(d_B.elements, B.elements, size,
	 cudaMemcpyHostToDevice);
	 // Allocate C in device memory
	 Matrix d_C;
	 d_C.width = C.width; d_C.height = C.height;
	 size = C.width * C.height * sizeof(float);
	 cudaMalloc(&d_C.elements, size);
	 // Invoke kernel
	 dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);
	 dim3 dimGrid(B.width / dimBlock.x, A.height / dimBlock.y);
	 MatMulKernel<<<dimGrid, dimBlock>>>(d_A, d_B, d_C);
	 // Read C from device memory
	 cudaMemcpy(C.elements, Cd.elements, size,
	 cudaMemcpyDeviceToHost);
	 // Free device memory
	 cudaFree(d_A.elements);
	 cudaFree(d_B.elements);
	 cudaFree(d_C.elements);
}
// Matrix multiplication kernel called by MatMul()
__global__ void MatMulKernel(Matrix A, Matrix B, Matrix C)
{
	 // Each thread computes one element of C
	 // by accumulating results into Cvalue
	 float Cvalue = 0;
	 int row = blockIdx.y * blockDim.y + threadIdx.y;
	 int col = blockIdx.x * blockDim.x + threadIdx.x;
	 for (int e = 0; e < A.width; ++e)
	 Cvalue += A.elements[row * A.width + e]
 
	 * B.elements[e * B.width + col];
	 C.elements[row * C.width + col] = Cvalue;
}

下面的代码示例是确实利用共享内存的矩阵乘法的实现。在该实现中，每个线程块负责计算C的一个平方子矩阵Csub，并且该块内的每个线程负责计算Csub的一个元素。如图10所示，Csub等于两个矩形矩阵的乘积：维A的子矩阵（A.width，block_size）与Csub具有相同的行索引，维B的子矩阵（block_size，A.width）具有与Csub相同的列索引。为了适合设备的资源，这两个矩形矩阵根据需要划分为多个具有block_size尺寸的正方形矩阵，并将Csub计算为这些正方形矩阵的乘积之和。这些产品中的每一个产品都是通过以下方式执行的：首先从一个全局线程将两个相应的平方矩阵加载到共享内存，然后通过一个线程加载每个矩阵的一个元素，然后让每个线程计算该乘积的一个元素。每个线程将每个乘积的结果累加到一个寄存器中，完成后将结果写入全局内存。

过以这种方式阻塞计算，我们利用了快速共享内存的优势，并节省了大量全局内存带宽，因为仅从全局内存中读取A（B.width / block_size）次，而读取B（A.height / block_size）次。

前一个代码示例中的Matrix类型增加了跨度字段，因此子矩阵可以有效地表示为相同类型。 __device__函数用于获取和设置元素以及从矩阵构建任何子矩阵。

// Matrices are stored in row-major order:
// M(row, col) = *(M.elements + row * M.stride + col)
typedef struct {
	 int width;
	 int height;
	 int stride;
	 float* elements;
} Matrix;
// Get a matrix element
__device__ float GetElement(const Matrix A, int row, int col)
{
	return A.elements[row * A.stride + col];
}
// Set a matrix element
__device__ void SetElement(Matrix A, int row, int col,
 float value)
{
	A.elements[row * A.stride + col] = value;
}
// Get the BLOCK_SIZExBLOCK_SIZE sub-matrix Asub of A that is
// located col sub-matrices to the right and row sub-matrices down
// from the upper-left corner of A
__device__ Matrix GetSubMatrix(Matrix A, int row, int col)
{
	 Matrix Asub;
	 Asub.width = BLOCK_SIZE;
	 Asub.height = BLOCK_SIZE;
	 Asub.stride = A.stride;
	 Asub.elements = &A.elements[A.stride * BLOCK_SIZE * row
	 + BLOCK_SIZE * col];
	 return Asub;
}
// Thread block size
#define BLOCK_SIZE 16
// Forward declaration of the matrix multiplication kernel
__global__ void MatMulKernel(const Matrix, const Matrix, Matrix);
// Matrix multiplication - Host code
// Matrix dimensions are assumed to be multiples of BLOCK_SIZE
void MatMul(const Matrix A, const Matrix B, Matrix C)
{
	 // Load A and B to device memory
	 Matrix d_A;
	 d_A.width = d_A.stride = A.width; d_A.height = A.height;
	 size_t size = A.width * A.height * sizeof(float);
	 cudaMalloc(&d_A.elements, size);
	 cudaMemcpy(d_A.elements, A.elements, size,
	 cudaMemcpyHostToDevice);
	 Matrix d_B;
	 d_B.width = d_B.stride = B.width; d_B.height = B.height;
	 size = B.width * B.height * sizeof(float);
	  cudaMalloc(&d_B.elements, size);
	 cudaMemcpy(d_B.elements, B.elements, size,
	 cudaMemcpyHostToDevice);
	 // Allocate C in device memory
	 Matrix d_C;
	 d_C.width = d_C.stride = C.width; d_C.height = C.height;
	 size = C.width * C.height * sizeof(float);
	 cudaMalloc(&d_C.elements, size);
	 // Invoke kernel
	 dim3 dimBlock(BLOCK_SIZE, BLOCK_SIZE);
	 dim3 dimGrid(B.width / dimBlock.x, A.height / dimBlock.y);
	 MatMulKernel<<<dimGrid, dimBlock>>>(d_A, d_B, d_C);
	 // Read C from device memory
	 cudaMemcpy(C.elements, d_C.elements, size,
	 cudaMemcpyDeviceToHost);
	 // Free device memory
	 cudaFree(d_A.elements);
	 cudaFree(d_B.elements);
	 cudaFree(d_C.elements);
}
// Matrix multiplication kernel called by MatMul()
__global__ void MatMulKernel(Matrix A, Matrix B, Matrix C)
{
	 // Block row and column
	 int blockRow = blockIdx.y;
	 int blockCol = blockIdx.x;
	 // Each thread block computes one sub-matrix Csub of C
	 Matrix Csub = GetSubMatrix(C, blockRow, blockCol);
	 // Each thread computes one element of Csub
	 // by accumulating results into Cvalue
	 float Cvalue = 0;
	 // Thread row and column within Csub
	 int row = threadIdx.y;
	 int col = threadIdx.x;
	 // Loop over all the sub-matrices of A and B that are
	 // required to compute Csub
	 // Multiply each pair of sub-matrices together
	 // and accumulate the results
	for (int m = 0; m < (A.width / BLOCK_SIZE); ++m) 
	{
		 // Get sub-matrix Asub of A
		 Matrix Asub = GetSubMatrix(A, blockRow, m);
		 // Get sub-matrix Bsub of B
		 Matrix Bsub = GetSubMatrix(B, m, blockCol);
		 // Shared memory used to store Asub and Bsub respectively
		 __shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
		 __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];
		 // Load Asub and Bsub from device memory to shared memory
		 // Each thread loads one element of each sub-matrix
		 As[row][col] = GetElement(Asub, row, col);
		 Bs[row][col] = GetElement(Bsub, row, col);
		 // Synchronize to make sure the sub-matrices are loaded
		 // before starting the computation
		 __syncthreads();
		  // Multiply Asub and Bsub together
		 for (int e = 0; e < BLOCK_SIZE; ++e)
		 Cvalue += As[row][e] * Bs[e][col];
		 // Synchronize to make sure that the preceding
		 // computation is done before loading two new
		 // sub-matrices of A and B in the next iteration
		 __syncthreads();
	}
	// Write Csub to device memory
	// Each thread writes one element
	SetElement(Csub, row, col, Cvalue);
}

Page-Locked Host Memory

运行时提供允许使用页面锁定（也称为固定）主机内存的功能（与malloc（）分配的常规可分页主机内存相对）：

cudaHostAlloc（）和cudaFreeHost（）分配并释放页面锁定的主机内存；
cudaHostRegister（）页面锁定由malloc（）分配的内存范围；

使用页面锁定的主机内存有几个好处：

如异步并发执行中所述，可以在某些设备的内核执行的同时执行页面锁定主机内存和设备内存之间的复制；
在某些设备上，可以将页面锁定的主机内存映射到设备的地址空间，而无需像“映射内存”中所述将其复制到设备内存或从设备内存复制。
在具有前端总线的系统上，如果主机内存被分配为页面锁定，则主机内存和设备内存之间的带宽会更高；如果另外将其分配为写合并，则主机内存与设备内存之间的带宽会更高，如写入合并内存中所述。

分页锁定的主机内存是一种稀缺资源，因此，分页锁定的内存中的分配将在可分页内存中的分配开始很久之前就开始失败。另外，通过减少操作系统可用于分页的物理内存量，消耗过多的页面锁定内存会降低整体系统性能。

分页锁定的主机内存未在非I / O一致的Tegra设备上缓存。另外，非I / O相干Tegra设备也不支持cudaHostRegister（）。

Portable Memory

可以将页面锁定的内存块与系统中的任何设备结合使用，但是默认情况下，使用上述页面锁定的内存的好处仅是与分配块时最新的设备一起使用。为了使这些优势可用于所有设备，需要通过将标志cudaHostAllocPortable传递给cudaHostAlloc（）来分配块，或者通过将标志cudaHostRegisterPortable传递给cudaHostRegister（）来进行页面锁定。

Write-Combining Memory

默认情况下，页面锁定主机内存分配为可缓存。可以通过将标志cudaHostAllocWriteCombined传递给cudaHostAlloc（）来选择将其分配为写合并。写合并内存释放了主机的L1和L2缓存资源，使更多的缓存可供应用程序的其余部分使用。此外，在通过PCI Express总线进行传输时，不会监听写合并内存，这可以将传输性能提高多达40％。

从主机读取写入合并内存的速度过慢，因此，写入合并内存通常应用于主机仅写入的内存。

Mapped Memory

通过将标志cudaHostAllocMapped传递给cudaHostAlloc（）或将标志cudaHostRegisterMapped传递给cudaHostRegister（），也可以将页面锁定主机内存块映射到设备的地址空间。因此，这样的块通常具有两个地址：一个位于主机内存中，由cudaHostAlloc（）或malloc（）返回，另一个位于设备内存中，可以使用cudaHostGetDevicePointer（）进行检索，然后用于从内核内部访问该块。唯一的例外是使用cudaHostAlloc（）分配的指针，以及主机和设备使用统一的地址空间时。

直接从内核内部访问主机内存有几个优点：

无需在设备存储器中分配块，也无需在该块与主机存储器中的块之间复制数据；数据传输由内核根据需要隐式执行；

无需使用流将数据传输与内核执行重叠；源自内核的数据传输会自动与内核执行重叠。

但是，由于映射的页面锁定内存在主机和设备之间共享，应用程序必须使用流或事件同步内存访问，以避免任何潜在的写后读，写后读或写后写危险。

为了能够检索指向任何映射的页面锁定内存的设备指针，必须在执行任何其他CUDA调用之前通过使用带有cudaDeviceMapHost标志的cudaSetDeviceFlags（）来启用页面锁定内存映射。否则，cudaHostGetDevicePointer（）将返回错误。

如果设备不支持映射的页面锁定主机内存，则cudaHostGetDevicePointer（）也会返回错误。应用程序可以通过检查canMapHostMemory设备属性（来查询此功能，对于支持映射的页面锁定主机内存的设备，该属性等于1。

Asynchronous Concurrent Execution

CUDA将以下操作公开为可以相互并发运行的独立任务：

主机上的计算；
设备上的计算；
内存从主机传输到设备；
内存从设备传输到主机；
内存在给定设备的内存中传输；
设备之间的内存传输。

Concurrent Execution between Host and Device

通过异步库函数促进并发主机执行，这些函数在设备完成请求的任务之前将控制权返回给主机线程。使用异步调用，可以在适当的设备资源可用时将许多设备操作排在一起，由CUDA驱动程序执行。这减轻了主机线程管理设备的大部分责任，使它可以自由执行其他任务。以下设备操作相对于主机是异步的：

内核启动；
内存复制在单个设备的内存中；
从主机到设备的内存副本，大小为64 KB或更小；
由带有Async后缀的功能执行的内存副本；
内存设置函数调用。

程序员可以通过将CUDA_LAUNCH_BLOCKING环境变量设置为1，为系统上运行的所有CUDA应用程序全局禁用内核启动的异步性。此功能仅用于调试目的，不应用作使生产软件可靠运行的方式。

如果通过分析器收集硬件计数器，则内核启动是同步的，除非启用了并发内核分析。如果异步内存副本涉及未分页锁定的主机内存，则它们也将是同步的。

Concurrent Kernel Execution

一些具有2.x和更高版本的计算能力的设备可以同时执行多个内核。应用程序可以通过检查并发内核设备属性来查询此功能，对于支持该功能的设备，该属性等于1。

设备可以同时执行的最大内核启动次数取决于其计算能力，并在表Table 14中列出。

来自一个CUDA context的内核不能与来自另一个CUDA context的内核同时执行。使用许多纹理或大量本地内存的内核不太可能与其他内核同时执行。

Overlap of Data Transfer and Kernel Execution

一些设备可以在执行内核的同时向GPU或从GPU执行异步内存复制。应用程序可以通过检查asyncEngineCount设备属性（请参阅设备枚举）来查询此功能，对于支持该属性的设备，该属性大于零。如果副本中包含主机内存，则必须对其进行页面锁定。

还可以在内核执行的同时（在支持并发内核设备属性的设备上）和/或在设备之间进行复制（对于支持asyncEngineCount属性的设备）同时执行设备内复制。设备内复制是使用标准内存复制功能启动的，目标和源地址位于同一设备上。

Concurrent Data Transfers

某些具有2.x和更高版本的计算能力的设备可以使设备之间的副本重叠。应用程序可以通过检查asyncEngineCount设备属性来查询此功能，对于支持该属性的设备，该属性等于2。为了重叠，传输中涉及的所有主机内存都必须页锁定。

Streams

应用程序通过流管理上述并行操作。流是按顺序执行的一系列命令（可能由不同的主机线程发出）。另一方面，不同的流可能会相对于彼此无序或同时执行它们的命令。这种行为无法得到保证，因此不应依赖于它的正确性（例如，内核间通信未定义）。

Creation and Destruction

通过创建流对象并将其指定为一系列内核启动和主机<->设备内存副本的流参数来定义流。下面的代码示例创建两个流，并在页面锁定的内存中分配一个float数组hostPtr。

cudaStream_t stream[2];
for (int i = 0; i < 2; ++i)
 cudaStreamCreate(&stream[i]);
float* hostPtr;
cudaMallocHost(&hostPtr, 2 * size);

以下代码示例将这些流中的每一个定义为从主机到设备的一个内存副本，从内核到设备的一个启动以及从设备到主机的一个内存副本的序列：

for (int i = 0; i < 2; ++i) {
     cudaMemcpyAsync(inputDevPtr + i * size, hostPtr + i * size,
     size, cudaMemcpyHostToDevice, stream[i]);
     MyKernel <<<100, 512, 0, stream[i]>>>
     (outputDevPtr + i * size, inputDevPtr + i * size, size);
     cudaMemcpyAsync(hostPtr + i * size, outputDevPtr + i * size,
     size, cudaMemcpyDeviceToHost, stream[i]);
}

每个流将其输入数组hostPtr的一部分复制到设备内存中的数组inputDevPtr，通过调用MyKernel（）处理设备上的inputDevPtr，并将结果outputDevPtr复制回到hostPtr的相同部分。重叠行为描述了此示例中的流如何根据设备的功能重叠。请注意，hostPtr必须指向页面锁定的主机内存，以便发生任何重叠.

通过调用cudaStreamDestroy（）释放流。

for (int i = 0; i < 2; ++i)
     cudaStreamDestroy(stream[i]);

如果在调用cudaStreamDestroy（）时设备仍在流中工作，则该函数将立即返回，并且一旦设备完成了流中的所有工作，与流相关的资源就会自动释放。

AliceWanderAI

关注

0
点赞
踩
0

收藏

觉得还不错? 一键收藏
0
评论
复制链接

分享到 QQ

分享到新浪微博

扫一扫