CUDA 11.6+设备函数中子核同步方案咨询

最优方案：放弃设备端嵌套核调用，主机端顺序执行 #

你的代码逻辑完全不需要依赖动态并行（设备端启动子核），直接在主机端顺序调用两个子核是最稳妥的方案，既避开了CUDA 11.6+对设备端同步的限制，还能减少动态并行带来的额外性能开销。

修改后的主机端调用逻辑示例：

// 计算第一个核的网格尺寸
int blocksize_multiple = (inputsize * outputsize + THREADS_PER_BLOCK - 1)/THREADS_PER_BLOCK;
NNFeedForwardNormalMultiple<<<blocksize_multiple, THREADS_PER_BLOCK>>>(values, weigths, result, inputsize, outputsize);

// 主机端同步，确保第一个核完全执行完毕
cudaDeviceSynchronize();

// 计算第二个核的网格尺寸
int blocksize_activate = (outputsize + THREADS_PER_BLOCK - 1)/THREADS_PER_BLOCK;
NNFeedForwardNormalActivate<<<blocksize_activate, THREADS_PER_BLOCK>>>(biases, result, outputsize);

特殊场景下的动态并行同步方案 #

如果你的业务逻辑必须在设备端触发核调用，可通过全局内存计数器+原子操作+内存栅栏实现子核同步：

1. 在全局内存中定义一个计数器，初始值设为子核的有效线程总数（即inputsize * outputsize）；
2. 子核每个有效线程执行完计算后，对计数器做原子减操作；
3. 父核循环等待计数器归0，配合内存栅栏保证变量可见性。

修改后的代码示例：

// 全局内存同步计数器，需在主机端初始化
__device__ volatile int kernel_completion_count = 0;

__global__ void NNFeedForwardNormalMultiple(double* __restrict__ values, double* __restrict__ weigths, double* result, int inputsize, int outputsize) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    int outputidx = idx / outputsize;
    int inputidx = idx % outputsize;

    if (outputidx >= outputsize || inputidx >= inputsize) {
        return;
    }

    atomicAdd(&result[outputidx], values[inputidx] * weigths[outputsize*outputidx + inputidx]);
    // 完成计算后原子减少计数器
    atomicAdd((int*)&kernel_completion_count, -1);
}

__global__ void NNFeedForwardNormalActivate(double* __restrict__ biases, double* result, int size) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;

    if (idx >= size) {
        return;
    }

    result[idx] = 1.0 / (1.0 + exp(-(result[idx] + biases[idx])));
}

__global__ void NNFeedForwardNormal(double* __restrict__ values, double* __restrict__ weigths, double* result, double* __restrict__ biases, int inputsize, int outputsize) {
    // 仅父核的一个线程初始化计数器
    if (threadIdx.x == 0 && blockIdx.x == 0) {
        kernel_completion_count = inputsize * outputsize;
        __threadfence(); // 确保计数器初始化对所有线程可见
    }
    __syncthreads(); // 父核所有线程等待初始化完成

    int blocksize = (inputsize * outputsize + THREADS_PER_BLOCK - 1)/THREADS_PER_BLOCK;
    NNFeedForwardNormalMultiple<<<blocksize, THREADS_PER_BLOCK>>>(values, weigths, result, inputsize, outputsize);

    // 循环等待子核全部完成
    while (kernel_completion_count > 0) {
        __threadfence(); // 保证读取到最新的计数器值
    }

    NNFeedForwardNormalActivate<<<(outputsize + THREADS_PER_BLOCK - 1)/THREADS_PER_BLOCK, THREADS_PER_BLOCK>>>(biases, result, outputsize);
}

注意：这种方式需严格保证计数器初始值等于子核有效线程数，否则会出现死锁；且父核线程会循环等待，资源占用较高，性能不如主机端直接调用。

内容的提问来源于stack exchange，提问作者ug0x01