CUDA 并发流

缺乏、安全感 2022-12-23 03:59 70阅读 0赞

### 文章目录 ###

*   *  创建，使用和销毁非默认CUDA流
     *   *  例子1：
         *  例子2:
         *  例子3 ： 多流向量加法

![在这里插入图片描述][watermark_type_ZmFuZ3poZW5naGVpdGk_shadow_10_text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L1RIX05VTQ_size_16_color_FFFFFF_t_70_pic_center]

默认流会阻塞非默认流的执行，非默认流之间不会相互阻塞。

## 创建，使用和销毁非默认CUDA流 ##

以下代码段演示了如何创建，利用和销毁非默认CUDA流。注意到，要在非默认CUDA流中启动CUDA核函数，必须将流作为执行配置的第4个可选参数传递给该核函数。到目前为止，仅利用了执行配置的前两个参数：

cudaStream_t stream;   // CUDA流的类型为 `cudaStream_t`
    cudaStreamCreate(&stream); // 注意，必须将一个指针传递给 `cudaCreateStream`
    
    someKernel<<<number_of_blocks, threads_per_block, 0, stream>>>();   // `stream` 作为第4个EC参数传递
    
    cudaStreamDestroy(stream); // 注意，将值（而不是指针）传递给 `cudaDestroyStream`

但值得一提的是，执行配置的第3个可选参数。此参数允许程序员提供共享内存（当前将不涉及的高级主题）中为每个内核启动动态分配的字节数。每个块分配给共享内存的默认字节数为“0”。

### 例子1： ###

#include <stdio.h>
    #include <unistd.h>
    
    __global__ void printNumber(int number)
    {
      printf("%d\n", number);
    }
    
    int main()
    {
      for (int i = 0; i < 5; ++i)
      {
        cudaStream_t stream;
        cudaStreamCreate(&stream);
        printNumber<<<1, 1, 0, stream>>>(i);
        cudaStreamDestroy(stream);
      }
      cudaDeviceSynchronize();
    }

### 例子2: ###

#include <stdio.h>
    
    __global__
    void initWith(float num, float *a, int N)
    {
    
      int index = threadIdx.x + blockIdx.x * blockDim.x;
      int stride = blockDim.x * gridDim.x;
    
      for(int i = index; i < N; i += stride)
      {
        a[i] = num;
      }
    }
    
    __global__
    void addVectorsInto(float *result, float *a, float *b, int N)
    {
      int index = threadIdx.x + blockIdx.x * blockDim.x;
      int stride = blockDim.x * gridDim.x;
    
      for(int i = index; i < N; i += stride)
      {
        result[i] = a[i] + b[i];
      }
    }
    
    void checkElementsAre(float target, float *vector, int N)
    {
      for(int i = 0; i < N; i++)
      {
        if(vector[i] != target)
        {
          printf("FAIL: vector[%d] - %0.0f does not equal %0.0f\n", i, vector[i], target);
          exit(1);
        }
      }
      printf("Success! All values calculated correctly.\n");
    }
    
    int main()
    {
      int deviceId;
      int numberOfSMs;
    
      cudaGetDevice(&deviceId);
      cudaDeviceGetAttribute(&numberOfSMs, cudaDevAttrMultiProcessorCount, deviceId);
    
      const int N = 2<<24;
      size_t size = N * sizeof(float);
    
      float *a;
      float *b;
      float *c;
    
      cudaMallocManaged(&a, size);
      cudaMallocManaged(&b, size);
      cudaMallocManaged(&c, size);
    
      cudaMemPrefetchAsync(a, size, deviceId);
      cudaMemPrefetchAsync(b, size, deviceId);
      cudaMemPrefetchAsync(c, size, deviceId);
    
      size_t threadsPerBlock;
      size_t numberOfBlocks;
    
      threadsPerBlock = 256;
      numberOfBlocks = 32 * numberOfSMs;
    
      cudaError_t addVectorsErr;
      cudaError_t asyncErr;
    
      /*
       * Create 3 streams to run initialize the 3 data vectors in parallel.
       */
    
      cudaStream_t stream1, stream2, stream3;
      cudaStreamCreate(&stream1);
      cudaStreamCreate(&stream2);
      cudaStreamCreate(&stream3);
    
      /*
       * Give each `initWith` launch its own non-standard stream.
       */
    
      initWith<<<numberOfBlocks, threadsPerBlock, 0, stream1>>>(3, a, N);
      initWith<<<numberOfBlocks, threadsPerBlock, 0, stream2>>>(4, b, N);
      initWith<<<numberOfBlocks, threadsPerBlock, 0, stream3>>>(0, c, N);
    
      addVectorsInto<<<numberOfBlocks, threadsPerBlock>>>(c, a, b, N);
    
      addVectorsErr = cudaGetLastError();
      if(addVectorsErr != cudaSuccess) printf("Error: %s\n", cudaGetErrorString(addVectorsErr));
    
      asyncErr = cudaDeviceSynchronize();
      if(asyncErr != cudaSuccess) printf("Error: %s\n", cudaGetErrorString(asyncErr));
    
      cudaMemPrefetchAsync(c, size, cudaCpuDeviceId);
    
      checkElementsAre(7, c, N);
    
      /*
       * Destroy streams when they are no longer needed.
       */
    
      cudaStreamDestroy(stream1);
      cudaStreamDestroy(stream2);
      cudaStreamDestroy(stream3);
    
      cudaFree(a);
      cudaFree(b);
      cudaFree(c);
    }

### 例子3 ： 多流向量加法 ###

#include <stdio.h>
    
    __global__
    void initWith(float num, float *a, int N)
    {
    
      int index = threadIdx.x + blockIdx.x * blockDim.x;
      int stride = blockDim.x * gridDim.x;
    
      for(int i = index; i < N; i += stride)
      {
        a[i] = num;
      }
    }
    
    __global__
    void addVectorsInto(float *result, float *a, float *b, int N)
    {
      int index = threadIdx.x + blockIdx.x * blockDim.x;
      int stride = blockDim.x * gridDim.x;
    
      for(int i = index; i < N; i += stride)
      {
        result[i] = a[i] + b[i];
      }
    }
    
    void checkElementsAre(float target, float *vector, int N)
    {
      for(int i = 0; i < N; i++)
      {
        if(vector[i] != target)
        {
          printf("FAIL: vector[%d] - %0.0f does not equal %0.0f\n", i, vector[i], target);
          exit(1);
        }
      }
      printf("Success! All values calculated correctly.\n");
    }
    
    int main()
    {
      int deviceId;
      int numberOfSMs;
    
      cudaGetDevice(&deviceId);
      cudaDeviceGetAttribute(&numberOfSMs, cudaDevAttrMultiProcessorCount, deviceId);
    
      const int N = 2<<24;
      size_t size = N * sizeof(float);
    
      float *a;
      float *b;
      float *c;
      float *h_c;
    
      cudaMalloc(&a, size);
      cudaMalloc(&b, size);
      cudaMalloc(&c, size);
      cudaMallocHost(&h_c, size);
    
      size_t threadsPerBlock;
      size_t numberOfBlocks;
    
      threadsPerBlock = 256;
      numberOfBlocks = 32 * numberOfSMs;
    
      cudaError_t addVectorsErr;
      cudaError_t asyncErr;
    
      /*
       * Create 3 streams to run initialize the 3 data vectors in parallel.
       */
    
      cudaStream_t stream1, stream2, stream3;
      cudaStreamCreate(&stream1);
      cudaStreamCreate(&stream2);
      cudaStreamCreate(&stream3);
    
      /*
       * Give each `initWith` launch its own non-standard stream.
       */
    
      initWith<<<numberOfBlocks, threadsPerBlock, 0, stream1>>>(3, a, N);
      initWith<<<numberOfBlocks, threadsPerBlock, 0, stream2>>>(4, b, N);
      initWith<<<numberOfBlocks, threadsPerBlock, 0, stream3>>>(0, c, N);
    
      for (int i = 0; i < 4; ++i)
      {
        cudaStream_t stream;
        cudaStreamCreate(&stream);
    
        addVectorsInto<<<numberOfBlocks/4, threadsPerBlock, 0, stream>>>(&c[i*N/4], &a[i*N/4], &b[i*N/4], N/4);
        cudaMemcpyAsync(&h_c[i*N/4], &c[i*N/4], size/4, cudaMemcpyDeviceToHost, stream);
        cudaStreamDestroy(stream);
      }
    
      addVectorsErr = cudaGetLastError();
      if(addVectorsErr != cudaSuccess) printf("Error: %s\n", cudaGetErrorString(addVectorsErr));
    
      asyncErr = cudaDeviceSynchronize();
      if(asyncErr != cudaSuccess) printf("Error: %s\n", cudaGetErrorString(asyncErr));
    
      checkElementsAre(7, h_c, N);
    
      /*
       * Destroy streams when they are no longer needed.
       */
    
      cudaStreamDestroy(stream1);
      cudaStreamDestroy(stream2);
      cudaStreamDestroy(stream3);
    
      cudaFree(a);
      cudaFree(b);
      cudaFree(c);
      cudaFreeHost(h_c);
    }

[watermark_type_ZmFuZ3poZW5naGVpdGk_shadow_10_text_aHR0cHM6Ly9ibG9nLmNzZG4ubmV0L1RIX05VTQ_size_16_color_FFFFFF_t_70_pic_center]: /images/20221120/c4a0f084f26a483baf4204b653445a89.png