以下哪句对一个 CUDA 流的描述最恰当？

1.可从多个并行线程中读取的数据缓冲区
2.一种用于在 GPU 上并发运行任何操作的方法
3.用于协调来自多个 CPU 的指令的 CUDA 机制
4. 按发布顺序执行的一系列操作
答案：4

用于控制非默认流行为的两条规则是什么？

1在同一非默认流中发布的操作将并行执行
2同一流中的操作将依发布顺序执行
3在不同非默认流中发布的操作之间没有必然顺序
4不同非默认流中的操作将始终并行执行
答案：2、3

以下哪项可以在非默认流中执行？
可参考 。请选择所有适用项。

1cudaMalloc
2cudaMemcpy
3cudaMemcpyAsync
4核函数的启动
答案：3、4

关于默认流，下面哪些表述是正确的？请选择所有适用项。

1默认流中的操作无法与任何非默认流中的操作同时执行
2默认流可被用来实现内存复制和 GPU 计算的重叠
3默认情况下，核函数的启动及许多其它 CUDA Runtime调用均在默认流中运行
4默认流也称为“Stream 0”或“NULL stream”
答案：1、3、4

非默认流

要创建新的非默认流，请向 cudaStreamCreate 传递一个 cudaStream_t 指针：

cudaStream_t stream;
cudaStreamCreate(&stream);

在非默认流中启动核函数

cudaStream_t stream;
cudaStreamCreate(&stream);

kernel<<<grid, blocks, 0, stream>>>();

销毁非默认流

cudaStream_t stream;
cudaStreamCreate(&stream);

kernel<<<grid, blocks, 0, stream>>>();

cudaStreamDestroy(stream);

例子：

helpers.cuh

#pragma once

#include <iostream>
#include <cstdint>
#include <string>

uint64_t sdiv (uint64_t a, uint64_t b) {
    return (a+b-1)/b;
}

void check_last_error ( ) {

    cudaError_t err;
    if ((err = cudaGetLastError()) != cudaSuccess) {
        std::cout << "CUDA error: " << cudaGetErrorString(err) << " : "
                  << __FILE__ << ", line " << __LINE__ << std::endl;
            exit(1);
    }
}

class Timer {

    float time;
    const uint64_t gpu;
    cudaEvent_t ying, yang;

public:

    Timer (uint64_t gpu=0) : gpu(gpu) {
        cudaSetDevice(gpu);
        cudaEventCreate(&ying);
        cudaEventCreate(&yang);
    }

    ~Timer ( ) {
        cudaSetDevice(gpu);
        cudaEventDestroy(ying);
        cudaEventDestroy(yang);
    }

    void start ( ) {
        cudaSetDevice(gpu);
        cudaEventRecord(ying, 0);
    }

    void stop (std::string label) {
        cudaSetDevice(gpu);
        cudaEventRecord(yang, 0);
        cudaEventSynchronize(yang);
        cudaEventElapsedTime(&time, ying, yang);
        std::cout << "TIMING: " << time << " ms (" << label << ")" << std::endl;
    }
};

encryption.cuh

#pragma once

#include <cstdint>
#include <assert.h>

__host__ __device__ __forceinline__
uint32_t hash (uint32_t x) {

        x ^= x >> 16;
        x *= 0x85ebca6b;
        x ^= x >> 13;
        x *= 0xc2b2ae35;
        x ^= x >> 16;

        return x;
}

__host__ __device__ __forceinline__
uint64_t permute64(uint64_t x, uint64_t num_iters) {

    constexpr uint64_t mask = (1UL << 32)-1;

    for (uint64_t iter = 0; iter < num_iters; iter++) {
        const uint64_t upper = x >> 32;
        const uint64_t lower = x & mask;
        const uint64_t mixer = hash(upper);

        x = upper + ((lower^mixer&mask) << 32);
    }

    return x;
}

__host__ __device__ __forceinline__
uint64_t unpermute64(uint64_t x, uint64_t num_iters) {

    constexpr uint64_t mask = (1UL << 32)-1;
    
    for (uint64_t iter = 0; iter < num_iters; iter++) {
        const uint64_t upper = x & mask;
        const uint64_t lower = x >> 32;
        const uint64_t mixer = hash(upper);

        x = (upper << 32) + (lower^mixer&mask);
    }

    return x;
}

baseline.cu

#include <cstdint>
#include <iostream>
#include "helpers.cuh"
#include "encryption.cuh"

void encrypt_cpu(uint64_t * data, uint64_t num_entries, 
                 uint64_t num_iters, bool parallel=true) {

    #pragma omp parallel for if (parallel)
    for (uint64_t entry = 0; entry < num_entries; entry++)
        data[entry] = permute64(entry, num_iters);
}

__global__ 
void decrypt_gpu(uint64_t * data, uint64_t num_entries, 
                 uint64_t num_iters) {

    const uint64_t thrdID = blockIdx.x*blockDim.x+threadIdx.x;
    const uint64_t stride = blockDim.x*gridDim.x;

    for (uint64_t entry = thrdID; entry < num_entries; entry += stride)
        data[entry] = unpermute64(data[entry], num_iters);
}

bool check_result_cpu(uint64_t * data, uint64_t num_entries,
                      bool parallel=true) {

    uint64_t counter = 0;

    #pragma omp parallel for reduction(+: counter) if (parallel)
    for (uint64_t entry = 0; entry < num_entries; entry++)
        counter += data[entry] == entry;

    return counter == num_entries;
}

int main (int argc, char * argv[]) {

    Timer timer;
    Timer overall;

    const uint64_t num_entries = 1UL << 26;
    const uint64_t num_iters = 1UL << 10;
    const bool openmp = true;

    timer.start();
    uint64_t * data_cpu, * data_gpu;
    cudaMallocHost(&data_cpu, sizeof(uint64_t)*num_entries);
    cudaMalloc    (&data_gpu, sizeof(uint64_t)*num_entries);
    timer.stop("allocate memory");
    check_last_error();

    timer.start();
    cudaStream_t stream;
    cudaStreamCreate(&stream);
    
    encrypt_cpu(data_cpu, num_entries, num_iters, openmp);
    timer.stop("encrypt data on CPU");

    overall.start();
    timer.start();
    cudaMemcpy(data_gpu, data_cpu, 
               sizeof(uint64_t)*num_entries, cudaMemcpyHostToDevice);
    timer.stop("copy data from CPU to GPU");
    check_last_error();

    timer.start();
    decrypt_gpu<<<80*32, 64, 0, stream>>>(data_gpu, num_entries, num_iters);
    
    cudaStreamDestroy(stream);
    
    timer.stop("decrypt data on GPU");
    check_last_error();

    timer.start();
    cudaMemcpy(data_cpu, data_gpu, 
               sizeof(uint64_t)*num_entries, cudaMemcpyDeviceToHost);
    timer.stop("copy data from GPU to CPU");
    overall.stop("total time on GPU");
    check_last_error();

    timer.start();
    const bool success = check_result_cpu(data_cpu, num_entries, openmp);
    std::cout << "STATUS: test " 
              << ( success ? "passed" : "failed")
              << std::endl;
    timer.stop("checking result on CPU");

    timer.start();
    cudaFreeHost(data_cpu);
    cudaFree    (data_gpu);
    timer.stop("free memory");
    check_last_error();
}

Makefile

CUDACXX=nvcc
CUDACXXFLAGS=-arch=sm_70 -O3
CXXFLAGS=-march=native -fopenmp
NSYS=nsys profile
NSYSFLAGS=--stats=true --force-overwrite=true

all: baseline

baseline: baseline.cu
	$(CUDACXX) $(CUDACXXFLAGS) -Xcompiler="$(CXXFLAGS)" baseline.cu -o baseline

profile: baseline
	$(NSYS) $(NSYSFLAGS) -o baseline-report ./baseline

clean:
	rm -f baseline baseline_solution *.qdrep *.sqlite

非默认流中内存复制

为了要异步复制数据，CUDA 需对其位置作出假设。典型的主机内存使用 分页技术，这样除了 RAM 之外，数据还可存储在某个备份存储设备上（如物理磁盘）。

固定（或锁页）内存会绕过主机操作系统分页，在 RAM 中存储所分配的内存。在非默认流中异步传输内存时，必须使用锁页（或固定）内存。

固定内存会阻止将数据存储在某个备份存储设备上，因此是一个受限资源，请务必当心不要过度使用它。

固定主机内存通过 cudaMallocHost 进行分配：

const uint64_t num_entries = 1UL << 26;
uint64_t *data_cpu;
cudaMallocHost(&data_cpu, sizeof(uint64_t)*num_entries);

非默认流中主机到设备的内存传输

通过使用类似于 cudaMemcpy 的 cudaMemcpyAsync，您可在非默认流中将固定主机内存传输到 GPU 显存，但需提供第 5 个流标识符参数：

cudaStream_t stream;
cudaStreamCreate(&stream);

const uint64_t num_entries = 1UL << 26;

uint64_t *data_cpu, *data_gpu;

cudaMallocHost(&data_cpu, sizeof(uint64_t)*num_entries);
cudaMalloc(&data_gpu, sizeof(uint64_t)*num_entries);

cudaMemcpyAsync(data_gpu, 
                data_cpu, 
                sizeof(uint64_t)*num_entries, 
                cudaMemcpyHostToDevice, 
                stream);

非默认流中设备到主机的内存传输

通过使用 cudaMemcpyAsync，您也可在非默认流中将 GPU 显存传输到固定主机内存：

// Assume data is already present on the GPU, and that `data_cpu` is pinned.

cudaMemcpyAsync(data_cpu, 
                data_gpu, 
                sizeof(uint64_t)*num_entries, 
                cudaMemcpyDeviceToHost, 
                stream);

与所有现代 GPU 一样，具有 2 个或更多复制引擎的 GPU 设备可以同时在不同的非默认流中执行主机到设备和设备到主机的内存传输。

流同步

使用cudaStreamSynchronize可导致主机代码阻塞，直到给定的流完成其操作为止。 当需要保证完成流工作时，例如，当主机代码需要等待非默认流中的异步内存传输完成时，应使用流同步：

// Assume data is already present on the GPU, and that `data_cpu` is pinned.

cudaMemcpyAsync(data_cpu, 
                data_gpu, 
                sizeof(uint64_t)*num_entries, 
                cudaMemcpyDeviceToHost, 
                stream);

// Block until work (in this case memory transfer to host) in `stream` is complete.
cudaStreamSynchronize(stream);

// `data_cpu` transfer to host via `stream` is now guaranteed to be complete.
checkResultCpu(data_cpu);

例子：

#include <cstdint>
#include <iostream>
#include "helpers.cuh"
#include "encryption.cuh"

void encrypt_cpu(uint64_t * data, uint64_t num_entries, 
                 uint64_t num_iters, bool parallel=true) {

    #pragma omp parallel for if (parallel)
    for (uint64_t entry = 0; entry < num_entries; entry++)
        data[entry] = permute64(entry, num_iters);
}

__global__ 
void decrypt_gpu(uint64_t * data, uint64_t num_entries, 
                 uint64_t num_iters) {

    const uint64_t thrdID = blockIdx.x*blockDim.x+threadIdx.x;
    const uint64_t stride = blockDim.x*gridDim.x;

    for (uint64_t entry = thrdID; entry < num_entries; entry += stride)
        data[entry] = unpermute64(data[entry], num_iters);
}

bool check_result_cpu(uint64_t * data, uint64_t num_entries,
                      bool parallel=true) {

    uint64_t counter = 0;

    #pragma omp parallel for reduction(+: counter) if (parallel)
    for (uint64_t entry = 0; entry < num_entries; entry++)
        counter += data[entry] == entry;

    return counter == num_entries;
}

int main (int argc, char * argv[]) {

    Timer timer;
    Timer overall;

    const uint64_t num_entries = 1UL << 26;
    const uint64_t num_iters = 1UL << 10;
    const bool openmp = true;

    timer.start();
    uint64_t * data_cpu, * data_gpu;
    cudaMallocHost(&data_cpu, sizeof(uint64_t)*num_entries);
    cudaMalloc    (&data_gpu, sizeof(uint64_t)*num_entries);
    timer.stop("allocate memory");
    check_last_error();

    timer.start();
    encrypt_cpu(data_cpu, num_entries, num_iters, openmp);
    timer.stop("encrypt data on CPU");

    overall.start();
    timer.start();
    
    cudaStream_t stream;
    cudaStreamCreate(&stream);
    
    
    //cudaMemcpy(data_gpu, data_cpu, 
    //           sizeof(uint64_t)*num_entries, cudaMemcpyHostToDevice);
    cudaMemcpyAsync(data_gpu, data_cpu, 
               sizeof(uint64_t)*num_entries, cudaMemcpyHostToDevice, stream);
    
    
    timer.stop("copy data from CPU to GPU");
    check_last_error();

    timer.start();
    decrypt_gpu<<<80*32, 64>>>(data_gpu, num_entries, num_iters);
    timer.stop("decrypt data on GPU");
    check_last_error();

    timer.start();
    // cudaMemcpy(data_cpu, data_gpu, 
    //           sizeof(uint64_t)*num_entries, cudaMemcpyDeviceToHost);
    
    cudaMemcpyAsync(data_cpu, data_gpu, 
               sizeof(uint64_t)*num_entries, cudaMemcpyDeviceToHost, stream);
               

    cudaStreamSynchronize(stream);
               
               
    timer.stop("copy data from GPU to CPU");
    overall.stop("total time on GPU");
    check_last_error();

    timer.start();
    const bool success = check_result_cpu(data_cpu, num_entries, openmp);
    std::cout << "STATUS: test " 
              << ( success ? "passed" : "failed")
              << std::endl;
    timer.stop("checking result on CPU");
    
    cudaStreamDestroy(stream);

    timer.start();
    cudaFreeHost(data_cpu);
    cudaFree    (data_gpu);
    timer.stop("free memory");
    check_last_error();
}