// nvcc -G stream.cu
// 关闭编译器优化
// Single GPU
 
#include <algorithm>
#include <iostream>
#include <cstdint>
 
// 一些简单的耗时操作
__global__ void someWork(int32_t length) {
    float res = 1.0f;
    for(int32_t i = 0; i < length; i++) {
        res = sinf(res) + cosf(res);
    }
}
 
int main() {
    // 0对应default stream
 
    // 然后单独创建两个新的Stream
    cudaStream_t s1, s2;
    // 使用普通方式创建的Stream 默认是和Default Stream同步的（不并行）
    cudaStreamCreate(&s1);
    // 明确指出此Stream不和Default Stream同步
    cudaStreamCreateWithFlags(&s2, cudaStreamNonBlocking);
 
    someWork<<<1, 32>>>(99999); // 默认在default stream上执行
    someWork<<<1, 32>>>(99999);
    someWork<<<1, 32, 0, s1>>>(99999);
    someWork<<<1, 32, 0, s1>>>(99999);
    someWork<<<1, 32, 0, s2>>>(99999);
    someWork<<<1, 32, 0, s2>>>(99999);
 
    std::cout << "kernel called" << std::endl;
 
    cudaStreamSynchronize(0); // 等待 default stream 结束
    // 同步自定义Stream 然后销毁
    cudaStreamSynchronize(s1);
    cudaStreamDestroy(s1);
    cudaStreamSynchronize(s2);
    cudaStreamDestroy(s2);
 
    // 这里只是写出来用于对比 其实没必要调用
    cudaDeviceSynchronize(); // 等待 device 上所有的 stream 都结束
 
    std::cout << "all finished" << std::endl;
 
    return 0;
}

一个Stream和Default Stream是同步的（线性执行），另一个和Default Stream是并行的

TODO

每张 GPU（每个设备）都有自己独立的默认 stream，也就是 device 0 的 default stream 和 device 1 的 default stream 互不相干

1. kernel会被放在stream中执行，异步的CUDA API好像也会被放到stream中，比如cudaMemcpyAsync
2. 同步的API呢？好像没有放到stream中的必要？