CUDA全局函数printf缓冲规则疑问：特定输出为何始终首尾出现

你观察到的输出顺序，核心原因是CUDA设备端printf的缓冲机制与线程调度的不确定性，具体拆解如下：

1. CUDA设备printf的核心规则 #

• 设备端的printf不会实时输出到主机，而是先将内容写入缓冲区。缓冲区的刷新时机包括：缓冲区填满、kernel执行完毕、调用cudaDeviceSynchronize()等同步函数。
• 同一个线程内的多次printf输出顺序是严格保留的，但不同线程（甚至不同线程块）的输出顺序是未定义的，完全由CUDA调度器的线程执行顺序决定。

2. 你的代码与输出的具体分析 #

你的代码中，只有id=5的线程（对应blockIdx.x=0, threadIdx.x=1, blockIdx.y=0, threadIdx.y=1）会输出Hello from the start和Hello from the end，同时该线程还会输出中间的processing index 5日志。

从输出结果可以看出：

• 该线程的第一个printf（Hello from the start）是所有线程中最早被执行并写入缓冲区的，因此会最先出现在主机输出中。
• 之后CUDA调度器优先执行了其他线程块的所有线程，这些线程的printf内容陆续写入缓冲区并最终被刷新输出。
• 最后调度器才继续执行id=5线程所在的剩余代码，因此该线程的processing index 5和Hello from the end会在所有其他线程的输出之后出现。

由于同一个线程的输出顺序被保留，所以Hello from the start必然早于该线程的中间日志，而Hello from the end必然在最后。

代码与输出验证 #

原代码 #

#include <stdio.h>
#include <stdlib.h>
#include <vector>
#include <iostream>

#include "cuda_runtime.h"
#include "device_launch_parameters.h"


__global__ void add1InGPU( int *devArr, int n, int nx )
{
    int ix = blockIdx.x * blockDim.x + threadIdx.x;
    int iy = blockIdx.y * blockDim.y + threadIdx.y;
    int id = nx * iy + ix;

    int whateverNum = 5;

    if (id == whateverNum) printf("Hello from the start\n");
    
    if ( id < n ) {
        devArr[id] += 1; 
        printf("blockIdx.x: %d, threadIdx.x: %d, blockIdx.y: %d, threadIdx.y: %d, processing index %d\n", blockIdx.x, threadIdx.x, blockIdx.y, threadIdx.y, id );
    }
    else {
        printf("blockIdx.x: %d, threadIdx.x: %d, blockIdx.y: %d, threadIdx.y: %d, skipped\n", blockIdx.x, threadIdx.x, blockIdx.y, threadIdx.y);
    }

    if (id == whateverNum) printf("Hello from the end\n");
}

int main(void)
{
    
    int* d_arr = NULL;
    cudaMalloc(&d_arr, 16 * sizeof(int));
    cudaMemset(d_arr, 0, 16 * sizeof(int));

    int nElem = 16;
    dim3 block( 2, 2 );
    dim3 grid( 2, 2 );
    // dim3 grid( ( nElem + block.x - 1 ) / block.x );
    
    add1InGPU<<<grid, block>>> ( d_arr, nElem, 4 );
    cudaDeviceSynchronize();
    int *h_arr = (int*)malloc(16 * sizeof(int));
    cudaMemcpy(h_arr, d_arr, 10 * sizeof(int), cudaMemcpyDeviceToHost);
    cudaFree(d_arr);

    volatile int _ = 0;
    for (int i = 0; i < 10; i++) {
        _ += h_arr[i];
    }

    cudaDeviceReset();
    
    return 0;
}

输出结果 #

Hello from the start
blockIdx.x: 1, threadIdx.x: 0, blockIdx.y: 0, threadIdx.y: 0, processing index 2
blockIdx.x: 1, threadIdx.x: 1, blockIdx.y: 0, threadIdx.y: 0, processing index 3
blockIdx.x: 1, threadIdx.x: 0, blockIdx.y: 0, threadIdx.y: 1, processing index 6
blockIdx.x: 1, threadIdx.x: 1, blockIdx.y: 0, threadIdx.y: 1, processing index 7
blockIdx.x: 0, threadIdx.x: 0, blockIdx.y: 1, threadIdx.y: 0, processing index 8
blockIdx.x: 0, threadIdx.x: 1, blockIdx.y: 1, threadIdx.y: 0, processing index 9
blockIdx.x: 0, threadIdx.x: 0, blockIdx.y: 1, threadIdx.y: 1, processing index 12
blockIdx.x: 0, threadIdx.x: 1, blockIdx.y: 1, threadIdx.y: 1, processing index 13
blockIdx.x: 1, threadIdx.x: 0, blockIdx.y: 1, threadIdx.y: 0, processing index 10
blockIdx.x: 1, threadIdx.x: 1, blockIdx.y: 1, threadIdx.y: 0, processing index 11
blockIdx.x: 1, threadIdx.x: 0, blockIdx.y: 1, threadIdx.y: 1, processing index 14
blockIdx.x: 1, threadIdx.x: 1, blockIdx.y: 1, threadIdx.y: 1, processing index 15
blockIdx.x: 0, threadIdx.x: 0, blockIdx.y: 0, threadIdx.y: 0, processing index 0
blockIdx.x: 0, threadIdx.x: 1, blockIdx.y: 0, threadIdx.y: 0, processing index 1
blockIdx.x: 0, threadIdx.x: 0, blockIdx.y: 0, threadIdx.y: 1, processing index 4
blockIdx.x: 0, threadIdx.x: 1, blockIdx.y: 0, threadIdx.y: 1, processing index 5
Hello from the end

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