detach vs kernel launch

Author: archibate

cudaguide/1_moderncuda/README.md

@@ -355,3 +355,97 @@ int main() {
     host_lambda(1);
 }
 ```
+
+### 第一个 CUDA 程序
+
+```cuda
+#include <cuda_runtime.h>
+#include <nvfunctional>
+#include "cudapp.cuh" // 小彭老师现代 CUDA 框架，更符合现代 C++ 风格，减少官方 C 风格接口的繁琐
+
+using namespace cudapp;
+
+__global__ void kernel(int x) {
+    printf("线程编号 (%d, %d)\n", blockIdx.x, threadIdx.x);
+}
+
+int main() {
+    // 三箭头语法糖启动内核
+    // kernel<<<blockDim, gridDim>>>(...)
+    kernel<<<3, 4>>>();
+
+    // 强制同步：等待此前启动过的所有内核执行完成
+    cudaDeviceSynchronize();
+
+    return 0;
+}
+```
+
+重要知识点：CUDA 为了极致性能，会使用**异步**的方式启动内核。
+
+使用三箭头语法调用内核函数，只是把内核提交到 GPU 上去而已，不代表内核已经在 GPU 上执行完毕。
+
+而 `cudaDeviceSynchronize()` 的作用就是等待**此前提交的所有内核**执行完毕，然后才能继续往下执行 `return 0`。
+
+> {{ icon.warn }} 如果不在程序退出前强制同步，则内核可能未执行！printf 语句会不生效！
+
+### 类比
+
+为了方便理解，可以把 `kernel<<<3, 4>>>()` 看作是启动了一个**后台线程**，而且还是以 `detach` 的方式，启动后就在后台默默运行，不会阻塞启动了内核的 CPU 线程。
+
+| GPU 操作 | CPU 类比 |
+|-|-|
+| 启动内核 `kernel<<<3, 4>>>()` | `std::async` 或 `std::thread` |
+| `cudaDeviceSynchronize` | `future.wait()` 或 `thread.join()` |
+
+### 小彭老师为你准备的 CUDA 框架
+
+```cuda
+#include <cuda_runtime.h>
+#include <nvfunctional>
+#include "cudapp.cuh" // 小彭老师现代 CUDA 框架，更符合现代 C++ 风格，减少官方 C 风格接口的繁琐
+
+using namespace cudapp;
+
+__global__ void kernel(int x) {
+    printf("内核参数 x = %d\n", x);
+    printf("线程编号 (%d, %d)\n", blockIdx.x, threadIdx.x);
+}
+
+int main() {
+    // 启动内核的3种方式
+    // 1. 官方三箭头语法糖（常用）
+    // kernel<<<blockDim, gridDim, dynamicSmemBytes, stream>>>(...)
+    int x = 42;
+    kernel<<<3, 4, 0, 0>>>(x);
+
+    // 2. cudaLaunchKernel
+    void *args[] = {&x};
+    CHECK_CUDA(cudaLaunchKernel(kernel, dim3(3), dim3(4), args, 0, 0));
+
+    // 3. cudaLaunchKernelEx
+    cudaLaunchConfig_t cfg{};
+    cfg.blockDim = dim3(3);
+    cfg.gridDim = dim3(4);
+    cfg.dynamicSmemBytes = 0;
+    cfg.stream = 0;
+    cfg.attrs = nullptr;
+    cfg.numAttrs = 0;
+    CHECK_CUDA(cudaLaunchKernelEx(&cfg, kernel, x));
+
+    const char *name;
+    CHECK_CUDA(cudaFuncGetName(&name, kernel));
+    printf("内核名字：%s\n", name);
+
+    // 1. 强制同步：等待此前启动过的所有内核执行完成
+    CHECK_CUDA(cudaDeviceSynchronize());
+
+    // 2. 仅同步 0 号流（null-stream）
+    CHECK_CUDA(cudaStreamSynchronize(0));
+
+    // 3. 仅同步 0 号流，但使用小彭老师现代 CUDA 框架
+    CudaStream::nullStream().join();
+
+    return 0;
+}
+```

cudaguide/1_moderncuda/main.cu

@@ -11,7 +11,8 @@ __global__ void kernel(int x) {
 
 int main() {
     // 启动内核的3种方式
-    // 1. 三箭头语法糖（常用）
+    // 1. 官方三箭头语法糖（常用）
+    // kernel<<<blockDim, gridDim, dynamicSmemBytes, stream>>>(...)
     int x = 42;
     kernel<<<3, 4, 0, 0>>>(x);