第12章 CUDA动态并行 - 核函数中启动核函数 (CUDA Dynamic Parallelism - Kernels Launching Kernels)

一、导读

12.1 本章学习目标

通过本章的学习，读者将能够：

理解CUDA动态并行（Dynamic Parallelism）的基本概念及其与传统主机启动模式的区别

掌握在设备端kernel中启动子kernel的语法和机制

深入理解动态并行的适用场景和性能考量

学会使用设备端同步原语协调父子kernel的执行

掌握递归算法在动态并行中的实现方法

了解动态并行对硬件资源的影响，包括内存占用和调度开销

能够分析何时使用动态并行能获得最佳收益，何时应避免使用

了解动态并行在实际应用中的典型案例，如树遍历、图搜索和自适应网格细化

理解命令处理器（Command Processor）的工作机制

掌握动态并行的调试技巧和常见问题排查方法

了解嵌套深度限制和资源分配策略

掌握动态并行与统一内存的交互行为

12.2 为什么本章如此重要

在CUDA 5.0之前，所有的kernel启动都必须由主机（CPU）发起。这种模式虽然简单明确，但在处理不规则或递归数据结构时显得笨拙。程序员需要在主机上管理复杂的任务队列，或者将递归算法改写为迭代形式，这增加了代码复杂度并可能降低可读性。

CUDA动态并行允许在设备端（GPU）的kernel中直接启动新的子kernel。这一特性带来了几个重要优势：

首先，动态并行简化了不规则算法的实现。对于树遍历、图搜索、自适应网格细化等应用，任务的产生是动态的、不可预测的。使用动态并行，子任务可以直接由父任务启动，无需返回主机进行调度。这不仅简化了代码，还减少了主机-设备通信的延迟。

其次，动态并行可以实现更细粒度的负载平衡。当某个线程或线程块发现额外的工作时，它可以立即启动新的子kernel来处理，而不需要等待全局同步。这使得工作可以更均匀地分布在GPU的各个SM上。

第三，动态并行支持自然的递归表达。许多算法（如快速排序、分治法、递归下降）天然就是递归的。动态并行允许这些算法以接近伪代码的形式实现，提高了代码的可维护性。

然而，动态并行并非银弹。它引入了额外的复杂性：设备端的kernel启动有显著的开销；子kernel的资源分配受到限制；调试变得更加困难。因此，理解动态并行的适用场景和性能特征至关重要。

动态并行的历史演进：

动态并行在CUDA 5.0（2012年）中首次引入，伴随着Kepler架构的发布。在此之前，所有kernel启动都必须由主机发起，限制了GPU处理不规则任务的能力。

随着架构的演进，动态并行的支持不断完善：

Kepler（GK110）：首次支持动态并行，命令处理器能力有限
Maxwell：改进了命令处理器性能，降低了启动开销
Pascal及以后：进一步优化了嵌套启动的开销，增加了嵌套深度支持
Ampere/Hopper：命令处理器更加高效，支持更深的嵌套和更多的并发子kernel

实际应用场景：

光线追踪：每个射线与场景求交后可能产生反射、折射等次级射线。动态并行可以自然地表达这种树状结构，每个次级射线由子kernel处理。

图算法：BFS的每一层可以由一个子kernel处理，避免不必要的全局同步。DFS可以使用递归来实现。

自适应网格细化：在科学模拟中，根据误差估计动态细化网格。细化产生的新网格单元可以由子kernel处理。

分治算法：快速排序、归并排序等可以用动态并行实现，每个子问题由子kernel处理。

稀疏数据处理：当处理稀疏矩阵或不规则数据结构时，动态并行可以根据非零元素的分布动态分配工作。

粒子系统：粒子碰撞检测可能产生可变数量的交互，动态并行可以高效处理这种情况。

12.3 核心问题引导

什么是CUDA动态并行？它与传统的host-launched kernel有何区别？

在设备端启动子kernel的语法是什么？需要什么编译选项？

动态并行的硬件支持要求是什么？哪些GPU架构支持？

子kernel如何与父kernel同步？cudaDeviceSynchronize()在设备端的行为是什么？

动态并行的开销来源有哪些？如何最小化这些开销？

递归算法如何在动态并行中实现？需要注意什么限制？

动态并行对寄存器使用和共享内存有什么影响？

在什么场景下动态并行能带来最大收益？什么场景下应该避免使用？

命令处理器的工作原理是什么？它如何调度子kernel？

如何处理动态并行中的错误传播？

二、核心概念详解

12.1 动态并行的基本概念

传统模式 vs 动态并行：

在传统CUDA编程模式中，kernel的启动流程是：

Host → Kernel Launch → Device Execution → Host Synchronization

所有kernel启动都由主机发起，设备端只能执行已启动的kernel，不能创建新的工作。

在动态并行模式中，流程变为：

Host → Parent Kernel → Child Kernel Launch → Device Execution
                ↓
          (on device)

设备端的kernel可以启动新的子kernel，形成层次化的执行结构。

关键特性：

嵌套启动：子kernel可以从父kernel中启动，理论上可以无限嵌套（受资源限制）。

隐式同步：父kernel在启动子kernel后，可以选择等待子kernel完成，或者继续执行其他工作。

独立调度：子kernel被提交到GPU的命令处理器（Command Processor），由硬件调度器独立调度执行。

资源共享：父子kernel共享同一设备的内存空间，可以直接访问彼此的数据（需要适当的同步）。

错误传播：子kernel中的错误可以传播到父kernel，但需要显式检查。

硬件要求：

动态并行需要Compute Capability 3.5或更高的GPU。具体来说：

Kepler架构（GK110）及以后支持动态并行
需要足够的硬件资源来支持设备端的kernel启动和管理
命令处理器的能力限制了并发子kernel的数量

编译要求：

使用动态并行需要特殊的编译选项：

bash

nvcc -arch=sm_35 -rdc=true -lcudadevrt program.cu

-rdc=true：启用可重定位设备代码
-lcudadevrt：链接设备运行时库
-arch=sm_35或更高：指定支持动态并行的架构

12.2 设备端Kernel启动语法

基本语法：

在设备端启动kernel的语法与主机端类似：

cpp

__global__ void childKernel(float* data, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        data[idx] *= 2.0f;
    }
}

__global__ void parentKernel(float* data, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    
    // Do some work
    if (idx < N) {
        data[idx] += 1.0f;
    }
    
    // Launch child kernel from device
    if (idx == 0) {  // Only one thread should launch
        childKernel<<<gridSize, blockSize>>>(data, N);
        
        // Wait for child to complete
        cudaDeviceSynchronize();
    }
}

注意事项：

启动线程的选择：通常只让一个线程（如threadIdx.x == 0 && blockIdx.x == 0）启动子kernel，避免重复启动。

配置参数：子kernel的grid和block大小可以是运行时计算的变量，提供了灵活性。

错误检查：设备端启动也可以使用cudaPeekAtLastError()检查错误。

资源限制：子kernel的资源（寄存器、共享内存）受SM剩余资源的限制。

带共享内存配置的启动：

cpp

// Launch with dynamic shared memory
childKernel<<<gridSize, blockSize, sharedMemSize>>>(args);

条件启动示例：

cpp

__global__ void conditionalLaunch(float* data, int* flags, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    
    if (idx < N && flags[idx]) {
        // This thread needs extra processing
        // Launch a specialized kernel for this element
        specializedKernel<<<1, 1>>>(data, idx);
    }
}

12.3 同步机制

`cudaDeviceSynchronize()`：

这是设备端最主要的同步原语。调用此函数会阻塞当前线程，直到设备上所有之前启动的kernel完成。

cpp

__global__ void parentWithSync(float* data, int N) {
    if (threadIdx.x == 0 && blockIdx.x == 0) {
        childKernel<<<1, 256>>>(data, N);
        
        // Wait for all previously launched kernels on this device
        cudaDeviceSynchronize();
        
        // Now safe to read results from child kernel
        float result = data[0];
    }
}

同步的范围：

cudaDeviceSynchronize()等待的是设备上所有之前启动的kernel，而不仅仅是当前父kernel启动的子kernel。这在多个子kernel并发执行时需要特别注意。

隐式同步点：

父kernel结束时，会自动等待其启动的所有子kernel完成
主机调用cudaDeviceSynchronize()或cudaMemcpy()时，会等待所有设备端工作完成
子kernel启动时，会等待父kernel中之前的所有操作完成

避免死锁：

确保不会出现循环依赖：父kernel等待子kernel，而子kernel又试图等待父kernel。

流与动态并行：

子kernel继承父kernel的流上下文。如果父kernel在特定流中执行，子kernel也在同一逻辑流中。

12.4 递归算法实现

动态并行使得递归算法的自然表达成为可能。

递归示例：二分搜索树遍历

cpp

struct TreeNode {
    int value;
    TreeNode* left;
    TreeNode* right;
};

__global__ void traverseKernel(TreeNode** nodes, int* results, 
                                int* outputIndex, int numNodes) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    
    if (tid < numNodes && nodes[tid] != NULL) {
        TreeNode* node = nodes[tid];
        results[tid] = node->value;
        
        // Prepare child nodes
        TreeNode* children[2] = {node->left, node->right};
        int childCount = 0;
        if (node->left) childCount++;
        if (node->right) childCount++;
        
        if (childCount > 0) {
            // Launch child kernel for next level
            TreeNode** d_children;
            cudaMalloc((void**)&d_children, childCount * sizeof(TreeNode*));
            cudaMemcpy(d_children, children, childCount * sizeof(TreeNode*),
                      cudaMemcpyDeviceToDevice);
            
            int childGridSize = 1;
            int childBlockSize = childCount;
            
            traverseKernel<<<childGridSize, childBlockSize>>>(
                d_children, results, outputIndex, childCount);
            
            cudaDeviceSynchronize();
            cudaFree(d_children);
        }
    }
}

注意事项：

递归深度限制：过深的递归可能导致栈溢出或资源耗尽。应设置最大深度限制。

内存分配：设备端的cudaMalloc和cudaFree开销较大，应尽量重用内存。

启动开销：每次kernel启动都有固定开销，递归产生的小kernel可能效率低下。

栈空间：递归调用会使用设备栈空间。默认栈空间可能不足以支持深层递归。

增加栈空间：

cpp

// On host side
cudaDeviceSetLimit(cudaLimitStackSize, 8192);  // Increase stack size

12.5 性能考量

开销来源：

启动延迟：设备端启动一个kernel的延迟约为几微秒，远高于线程切换的开销。

资源分配：每个新kernel需要分配执行资源，可能影响其他正在执行的kernel。

同步成本：cudaDeviceSynchronize()可能引入显著的等待时间。

命令处理器瓶颈：GPU的命令处理器处理能力有限，大量并发启动可能导致排队延迟。

内存分配开销：设备端的malloc/free比主机端慢。

优化策略：

批量启动：将多个小任务合并为一个较大的kernel启动，减少启动次数。

异步执行：如果子任务之间没有依赖，可以不立即同步，让它们并发执行。

阈值控制：只有当任务规模超过某个阈值时才使用动态并行，否则使用串行或线程级并行。

内存预分配：预先分配所需的内存，避免在递归过程中频繁分配和释放。

减少嵌套深度：尽量将深层递归转换为浅层迭代。

何时使用动态并行：

✅ 适合的场景：

任务结构是不规则的树或图
子任务的产生依赖于父任务的计算结果
任务规模足够大，启动开销可以忽略
代码简化带来的可维护性提升值得性能权衡
主机-设备通信延迟成为瓶颈

❌ 不适合的场景：

任务结构是规则的数组或矩阵
子任务非常小（如单个元素的处理）
对延迟极其敏感的实时应用
可以在编译时确定所有任务的情况
资源受限的环境

12.6 资源管理

寄存器使用：

设备端启动的kernel同样受寄存器限制。父kernel和子kernel的寄存器使用是独立的，但总的活跃线程数受SM资源限制。

共享内存：

子kernel可以使用共享内存，但其容量受SM限制。如果父kernel已经使用了大量共享内存，可能影响子kernel的occupancy。

栈空间：

递归调用会使用设备栈空间。默认栈空间可能不足以支持深层递归，可以通过cudaDeviceSetLimit(cudaLimitStackSize, size)增加。

本地内存：

如果寄存器不足，变量可能溢出到本地内存（全局内存的一部分），导致性能下降。应监控寄存器使用。

命令处理器资源：

命令处理器有有限的pending kernel队列。如果队列满，新的kernel启动会被阻塞。

12.7 错误处理

设备端错误检查：

cpp

__global__ void parentWithErrorCheck(float* data, int N) {
    if (threadIdx.x == 0 && blockIdx.x == 0) {
        childKernel<<<1, 256>>>(data, N);
        
        cudaError_t err = cudaGetLastError();
        if (err != cudaSuccess) {
            // Handle error
            printf("Child kernel launch failed: %s\n", cudaGetErrorString(err));
        }
        
        cudaDeviceSynchronize();
        
        err = cudaGetLastError();
        if (err != cudaSuccess) {
            printf("Child kernel execution failed: %s\n", cudaGetErrorString(err));
        }
    }
}

错误传播：

子kernel的错误不会自动传播到主机。父kernel需要显式检查并处理错误。

三、程序员的视角

3.1 在实际项目中的应用

光线追踪：在路径追踪中，每个射线与场景求交后可能产生反射、折射等次级射线。动态并行可以自然地表达这种树状结构，每个次级射线由子kernel处理。

图算法：BFS（广度优先搜索）的每一层可以由一个子kernel处理。当某一层的所有节点处理完毕后，启动下一层的kernel。

自适应网格细化：在科学模拟中，需要根据误差估计动态细化网格。细化产生的新网格单元可以由子kernel处理。

分治算法：快速排序、归并排序等分治算法可以用动态并行实现，每个子问题由子kernel处理。

稀疏矩阵运算：当处理稀疏矩阵时，非零元素的分布可能不均匀。动态并行可以根据实际分布动态分配工作。

3.2 常见陷阱与误区

过度使用动态并行：对于规则的数据并行任务，动态并行的开销可能超过收益。应优先使用传统的grid/block并行。

忘记同步：启动子kernel后忘记同步就读取结果，导致数据竞争和错误结果。

递归过深：深层递归可能导致栈溢出或资源耗尽。应设置最大深度限制或使用迭代替代。

重复启动：多个线程同时启动相同的子kernel，导致重复工作。应确保只有一个线程负责启动。

忽视启动开销：启动大量小kernel的累积开销可能很大。应批量处理小任务。

资源泄漏：设备端的malloc如果没有对应的free，会导致内存泄漏。

3.3 与其他层级的联系

与内存层次的联系：动态并行不改变内存层次结构，但子kernel可以访问父kernel写入全局内存的数据。

与同步机制的联系：设备端同步是动态并行的关键，需要仔细设计同步点以避免死锁和数据竞争。

与递归的联系：动态并行使得递归算法在GPU上的实现成为可能，但需要注意递归深度和资源限制。

与流的联系：子kernel继承父kernel的流上下文，可以与主机端的流操作协同工作。

四、代码实战

4.1 简单的动态并行示例

cpp

#include <cuda_runtime.h>
#include <stdio.h>

// Child kernel
__global__ void childKernel(float* data, int start, int end) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    int globalIdx = start + idx;
    
    if (globalIdx < end) {
        data[globalIdx] *= 2.0f;
    }
}

// Parent kernel that launches child
__global__ void parentKernel(float* data, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    
    // First pass: initial processing
    if (idx < N) {
        data[idx] += 1.0f;
    }
    
    // Second pass: launch child kernel for further processing
    if (idx == 0) {
        int childBlockSize = 256;
        int childGridSize = (N + childBlockSize - 1) / childBlockSize;
        
        childKernel<<<childGridSize, childBlockSize>>>(data, 0, N);
        
        // Wait for child to complete
        cudaDeviceSynchronize();
    }
}

int main() {
    int N = 1000;
    float *h_data = (float*)malloc(N * sizeof(float));
    for (int i = 0; i < N; i++) {
        h_data[i] = (float)i;
    }
    
    float *d_data;
    cudaMalloc((void**)&d_data, N * sizeof(float));
    cudaMemcpy(d_data, h_data, N * sizeof(float), cudaMemcpyHostToDevice);
    
    int blockSize = 256;
    int gridSize = (N + blockSize - 1) / blockSize;
    
    parentKernel<<<gridSize, blockSize>>>(d_data, N);
    cudaDeviceSynchronize();
    
    cudaMemcpy(h_data, d_data, N * sizeof(float), cudaMemcpyDeviceToHost);
    
    printf("First 10 elements after processing:\n");
    for (int i = 0; i < 10; i++) {
        printf("%.1f ", h_data[i]);  // Expected: (i + 1) * 2
    }
    printf("\n");
    
    cudaFree(d_data);
    free(h_data);
    
    return 0;
}

编译命令：

bash

nvcc -arch=sm_35 -rdc=true -lcudadevrt dynamic_parallel.cu -o dynamic_parallel

4.2 递归二分搜索

cpp

#include <cuda_runtime.h>
#include <stdio.h>

__global__ void binarySearchKernel(int* sortedArray, int target,
                                    int low, int high, int* result, int depth) {
    if (low > high) {
        return;
    }
    
    int mid = (low + high) / 2;
    
    if (sortedArray[mid] == target) {
        atomicMin(result, mid);  // Found
        return;
    }
    
    // Limit recursion depth
    if (depth > 20) {
        return;
    }
    
    // Launch sub-kernel for the appropriate half
    if (threadIdx.x == 0 && blockIdx.x == 0) {
        if (target < sortedArray[mid]) {
            binarySearchKernel<<<1, 1>>>(sortedArray, target, low, mid - 1, result, depth + 1);
        } else {
            binarySearchKernel<<<1, 1>>>(sortedArray, target, mid + 1, high, result, depth + 1);
        }
        
        cudaDeviceSynchronize();
    }
}

int main() {
    int h_array[] = {1, 3, 5, 7, 9, 11, 13, 15, 17, 19};
    int N = 10;
    int target = 7;
    
    int *d_array, *d_result;
    cudaMalloc((void**)&d_array, N * sizeof(int));
    cudaMalloc((void**)&d_result, sizeof(int));
    
    cudaMemcpy(d_array, h_array, N * sizeof(int), cudaMemcpyHostToDevice);
    
    int h_result = N;  // Initialize to invalid index
    cudaMemcpy(d_result, &h_result, sizeof(int), cudaMemcpyHostToDevice);
    
    binarySearchKernel<<<1, 1>>>(d_array, target, 0, N - 1, d_result, 0);
    cudaDeviceSynchronize();
    
    cudaMemcpy(&h_result, d_result, sizeof(int), cudaMemcpyDeviceToHost);
    
    if (h_result < N) {
        printf("Target %d found at index: %d\n", target, h_result);
    } else {
        printf("Target %d not found\n", target);
    }
    
    cudaFree(d_array);
    cudaFree(d_result);
    
    return 0;
}

4.3 树遍历示例

cpp

#include <cuda_runtime.h>
#include <stdio.h>

struct TreeNode {
    int value;
    TreeNode* left;
    TreeNode* right;
};

__global__ void treeSumKernel(TreeNode** nodes, int* partialSums, int numNodes, int depth) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    
    if (tid < numNodes && nodes[tid] != NULL) {
        TreeNode* node = nodes[tid];
        partialSums[tid] = node->value;
        
        // Collect children
        TreeNode* children[2];
        int childCount = 0;
        if (node->left) children[childCount++] = node->left;
        if (node->right) children[childCount++] = node->right;
        
        if (childCount > 0 && depth < 10) {
            TreeNode** d_children;
            cudaMalloc((void**)&d_children, childCount * sizeof(TreeNode*));
            cudaMemcpy(d_children, children, childCount * sizeof(TreeNode*),
                      cudaMemcpyDeviceToDevice);
            
            int* childSums;
            cudaMalloc((void**)&childSums, childCount * sizeof(int));
            
            treeSumKernel<<<1, childCount>>>(d_children, childSums, childCount, depth + 1);
            cudaDeviceSynchronize();
            
            // Accumulate children's sums
            for (int i = 0; i < childCount; i++) {
                partialSums[tid] += childSums[i];
            }
            
            cudaFree(d_children);
            cudaFree(childSums);
        }
    }
}

int main() {
    // Build a simple tree on host
    TreeNode nodes[7];
    for (int i = 0; i < 7; i++) {
        nodes[i].value = i + 1;
        nodes[i].left = (2*i+1 < 7) ? &nodes[2*i+1] : NULL;
        nodes[i].right = (2*i+2 < 7) ? &nodes[2*i+2] : NULL;
    }
    
    TreeNode* d_nodes[7];
    TreeNode* d_nodeArray;
    cudaMalloc((void**)&d_nodeArray, 7 * sizeof(TreeNode));
    cudaMemcpy(d_nodeArray, nodes, 7 * sizeof(TreeNode), cudaMemcpyHostToDevice);
    
    // Set up pointers
    for (int i = 0; i < 7; i++) {
        d_nodes[i] = d_nodeArray + i;
    }
    
    TreeNode** d_ptrArray;
    cudaMalloc((void**)&d_ptrArray, 7 * sizeof(TreeNode*));
    cudaMemcpy(d_ptrArray, d_nodes, 7 * sizeof(TreeNode*), cudaMemcpyHostToDevice);
    
    int* d_sums;
    cudaMalloc((void**)&d_sums, 7 * sizeof(int));
    
    treeSumKernel<<<1, 7>>>(d_ptrArray, d_sums, 7, 0);
    cudaDeviceSynchronize();
    
    int h_sums[7];
    cudaMemcpy(h_sums, d_sums, 7 * sizeof(int), cudaMemcpyDeviceToHost);
    
    printf("Root sum: %d\n", h_sums[0]);  // Should be 1+2+3+4+5+6+7 = 28
    
    cudaFree(d_nodeArray);
    cudaFree(d_ptrArray);
    cudaFree(d_sums);
    
    return 0;
}

五、深入思考

5.1 历史背景

动态并行在CUDA 5.0（2012年）中首次引入，伴随着Kepler架构的发布。在此之前，所有kernel启动都必须由主机发起，限制了GPU处理不规则任务的能力。

动态并行的引入反映了GPU计算从纯数据并行向更通用并行模型的演进。它使得GPU能够更好地处理递归、树结构和图等传统上难以并行化的数据结构。

5.2 前沿方向

硬件命令处理器改进：新一代GPU的命令处理器能力不断增强，可以处理更多的并发kernel启动。

细粒度同步：未来的架构可能提供更细粒度的同步原语，减少动态并行的等待开销。

与统一内存的集成：动态并行与统一内存结合，可以进一步简化不规则数据结构的处理。

机器学习应用：动态并行在决策树、图神经网络等机器学习模型中有潜在应用。

编译器优化：未来的编译器可能自动识别适合动态并行的代码模式，并进行转换。

5.3 延伸阅读

NVIDIA CUDA C Programming Guide - Dynamic Parallelism

"CUDA Dynamic Parallelism" by NVIDIA Developer Blog

Garland et al., "Exploring Interactive Ray Tracing with Dynamic Parallelism"

CUDA Samples - dynamicParallelism

"Efficient Implementation of Dynamic Parallelism in CUDA" - Research papers

六、实验笔记

实验1：动态并行开销测量

目的：量化设备端kernel启动的开销。

步骤：

启动一个空的子kernel，测量启动和同步的时间

比较不同数量的子kernel启动的累积开销

与主机启动相同kernel的开销对比

预期：设备端启动开销约为几微秒，对于小任务可能占主导。

实验2：递归深度对性能的影响

目的：研究递归深度与执行时间的关系。

步骤：

实现一个递归的二分搜索

测试不同深度的递归执行时间

找出性能开始下降的深度阈值

预期：浅层递归性能良好，深层递归由于栈空间和启动开销，性能下降。

实验3：树遍历性能对比

目的：比较动态并行树遍历与传统CPU实现的性能。

步骤：

实现GPU动态并行树遍历

实现CPU递归树遍历

对不同大小的树进行性能对比

预期：对于大树，动态并行可能有优势；对于小树，CPU可能更快。

七、本章小结

CUDA动态并行允许在设备端kernel中启动子kernel，为处理不规则和递归任务提供了强大的工具。它简化了树遍历、图搜索和自适应算法的实现，减少了主机-设备通信的需求。

动态并行的核心机制包括设备端kernel启动语法、cudaDeviceSynchronize()同步原语、以及硬件命令处理器的支持。正确使用这些机制可以实现高效的层次化并行。

然而，动态并行并非适用于所有场景。它的启动开销使得小任务的处理效率低下，深层递归可能导致资源耗尽。程序员需要仔细权衡动态并行的收益和成本。

在实际应用中，动态并行在光线追踪、图算法和自适应网格细化等领域展现了价值。对于这些应用，动态并行可以显著简化代码并提高性能。

八、习题

语法练习：编写一个parent kernel，启动两个不同的child kernel，分别处理数组的前半部分和后半部分。

同步分析：如果parent kernel启动子kernel后立即读取结果，而没有调用cudaDeviceSynchronize()，会发生什么？

递归限制：CUDA设备栈的默认大小是多少？如何增加它以支持更深的递归？

性能建模：估算启动100个子kernel（每个处理1000个元素）的总开销。与启动一个大kernel处理所有元素相比，差异有多大？

死锁检测：设计一个会导致死锁的动态并行代码示例。解释如何避免。

资源分析：如果parent kernel使用了大部分共享内存，子kernel的occupancy会受到什么影响？

应用场景：列举三个适合使用动态并行的应用，并解释原因。

替代方案：对于不适合动态并行的小任务，有什么替代的并行化策略？

错误处理：如何在设备端检查子kernel启动是否成功？

综合实现：实现一个使用动态并行的并行归并排序。比较其与传统实现的复杂度和性能。

命令处理器：研究命令处理器的工作原理。它如何调度多个并发的子kernel？

嵌套深度：CUDA支持的最大嵌套深度是多少？如何通过实验验证？