CUDA：创建和使用Stream

CUDA中的流（stream）表示一个GPU操作队列，该队列中的操作将以添加到流中的先后顺序而依次执行。可以将一个流看做是GPU上的一个任务，不同的流（即不同的任务）可以并行执行。这种并行性允许在同一设备上的多个操作重叠执行，从而提高了GPU的利用率和程序的执行效率。

cudaSetStream(): 用于设置当前线程的执行流。
cudaStreamCreate(): 用于创建执行流。
cudaStreamSynchronize(): 用于等待执行流中的所有操作完成。
cudaStreamDestroy(): 用于销毁执行流。

使用 CUDA 流的基本步骤如下：
创建流：通过调用 cudaStreamCreate() 创建一个或多个流对象。
指定流：在启动核函数、执行内存传输或其他 CUDA 操作时，通过传递流对象来指定应该使用哪个流。
同步流：可以使用 cudaStreamSynchronize() 函数等待特定流中的所有操作完成。这会阻塞主机执行，直到 GPU 完成该流中的所有工作。
销毁流：当不再需要流时，应该使用 cudaStreamDestroy() 函数来释放与流相关联的资源。

下面是一个简单的例子，展示了如何在CUDA中使用流来并行执行数据复制和核函数运算：

#include <stdio.h>  
#include <cuda_runtime.h>  
  
// 假设我们有一个简单的CUDA核函数  
__global__ void simpleKernel(int *d_array, int arraySize) {  
    int index = threadIdx.x + blockIdx.x * blockDim.x;  
    if (index < arraySize) {  
        d_array[index] = index * index; // 计算平方  
    }  
}  
  
int main() {  
    const int arraySize = 100;  
    int h_array[arraySize]; // 主机上的数组  
    int *d_array1, *d_array2; // 设备上的两个数组  
  
    // 创建两个CUDA流  
    cudaStream_t stream1, stream2;  
    cudaStreamCreate(&stream1);  
    cudaStreamCreate(&stream2);  
  
    // 分配设备内存  
    cudaMalloc((void**)&d_array1, arraySize * sizeof(int));  
    cudaMalloc((void**)&d_array2, arraySize * sizeof(int));  
  
    // 在第一个流中复制数据到设备并运行核函数  
    cudaMemcpyAsync(d_array1, h_array, arraySize * sizeof(int), cudaMemcpyHostToDevice, stream1);  
    simpleKernel<<<1, 256, 0, stream1>>>(d_array1, arraySize);  
  
    // 在第二个流中复制不同的数据到设备并运行核函数  
    cudaMemcpyAsync(d_array2, h_array + arraySize, arraySize * sizeof(int), cudaMemcpyHostToDevice, stream2);  
    simpleKernel<<<1, 256, 0, stream2>>>(d_array2, arraySize);  
  
    // 等待两个流中的所有操作完成  
    cudaStreamSynchronize(stream1);  
    cudaStreamSynchronize(stream2);  
  
    // 在主流中从设备复制结果回主机  
    cudaMemcpy(h_array, d_array1, arraySize * sizeof(int), cudaMemcpyDeviceToHost);  
    cudaMemcpy(h_array + arraySize, d_array2, arraySize * sizeof(int), cudaMemcpyDeviceToHost);  
  
    // 释放资源和销毁流  
    cudaFree(d_array1);  
    cudaFree(d_array2);  
    cudaStreamDestroy(stream1);  
    cudaStreamDestroy(stream2);  
  
    // 输出结果  
    for (int i = 0; i < 2 * arraySize; i++) {  
        printf("%d ", h_array[i]);  
    }  
    printf("\n");  
  
    return 0;  
}

在这个例子中，我们创建了两个CUDA流（stream1 和 stream2）。每个流中，我们异步地执行了一个数据复制操作（使用 cudaMemcpyAsync）和一个核函数运算（使用 simpleKernel）。这两个操作在各自的流中并行执行，互不影响。
我们使用 cudaStreamSynchronize 来等待每个流中的所有操作完成，然后我们再从设备复制结果回主机。最后，我们释放了分配的资源并销毁了流。
请注意，这个例子仅用于演示CUDA流的使用方式。在实际应用中，你可能需要更复杂的流管理策略，以确保操作的正确顺序和性能优化。此外，由于流是异步的，需要特别注意数据依赖和同步问题，以避免出现错误的结果。

CUDA流的应用

使用CUDA流（Streams）与不使用流在性能上的差异取决于多种因素，包括GPU的架构、并行任务的数量和性质、内存访问模式等。在某些情况下，使用流可以显著提高性能，尤其是当多个独立的任务可以并行执行，并且这些任务的执行顺序不相互依赖时。

下面是一个简单的例子来说明使用流和不使用流在性能上可能的差异：
不使用流的例子
假设我们有两个独立的CUDA任务：Task A 和 Task B。这两个任务分别执行数据复制和核函数运算，并且它们之间没有数据依赖关系。如果我们不使用流，代码可能看起来像这样：

// 假设我们有相应的数据结构和核函数  
int *d_array1, *d_array2;  
cudaMalloc(&d_array1, ...);  
cudaMalloc(&d_array2, ...);  
  
// 执行Task A  
cudaMemcpy(d_array1, h_array1, ...);  
simpleKernelA<<<1, 256>>>(d_array1, ...);  
  
// 等待Task A完成  
cudaDeviceSynchronize();  
  
// 执行Task B  
cudaMemcpy(d_array2, h_array2, ...);  
simpleKernelB<<<1, 256>>>(d_array2, ...);  
  
// 等待Task B完成  
cudaDeviceSynchronize();  
  
// 从设备复制结果回主机  
cudaMemcpy(h_result1, d_array1, ...);  
cudaMemcpy(h_result2, d_array2, ...);  
  
cudaFree(d_array1);  
cudaFree(d_array2);

在这种情况下，Task A 和 Task B 是顺序执行的。即使它们可以在不同的CUDA核心上并行执行，它们也必须一个接一个地完成。这意味着在Task A完成之前，Task B不能开始。这可能会导致GPU资源的浪费，尤其是在Task A和Task B执行时间相差很大的情况下。
使用流的例子
如果我们使用两个独立的流来执行这两个任务，代码可能如下所示：

cudaStream_t stream1, stream2;  
cudaStreamCreate(&stream1);  
cudaStreamCreate(&stream2);  
  
// 在stream1中执行Task A  
cudaMemcpyAsync(d_array1, h_array1, ..., stream1);  
simpleKernelA<<<1, 256, 0, stream1>>>(d_array1, ...);  
  
// 在stream2中执行Task B  
cudaMemcpyAsync(d_array2, h_array2, ..., stream2);  
simpleKernelB<<<1, 256, 0, stream2>>>(d_array2, ...);  
  
// 等待两个流都完成  
cudaStreamSynchronize(stream1);  
cudaStreamSynchronize(stream2);  
  
// 从设备复制结果回主机  
cudaMemcpy(h_result1, d_array1, ...);  
cudaMemcpy(h_result2, d_array2, ...);  
  
cudaFree(d_array1);  
cudaFree(d_array2);  
cudaStreamDestroy(stream1);  
cudaStreamDestroy(stream2);

在这种情况下，Task A 和 Task B 可以并行执行，因为它们在不同的流中。这意味着只要GPU资源允许，Task B可以在Task A还在执行时开始。这可以提高GPU的利用率，并可能减少总体执行时间。
性能差异
性能差异的大小取决于多个因素，包括GPU的工作负载、内存带宽的利用率、CUDA核心的使用率等。如果GPU资源充足且任务之间没有数据依赖关系，使用流通常可以提高性能。然而，如果任务之间存在数据依赖关系或者GPU资源已经饱和，那么使用流可能不会带来明显的性能提升，甚至可能导致性能下降（因为额外的流管理可能会引入额外的开销）。
总的来说，是否使用流需要根据具体的应用场景和任务性质来决定。在开发CUDA程序时，建议使用性能分析工具（如NVIDIA Nsight或Visual Profiler）来评估不同实现之间的性能差异，并据此做出决策。