CUDA：意想不到的printf行爲

Question

我不明白我在CUDA內核中使用printf觀察到的行爲。 有人可以對此有所瞭解嗎？如果這是正常的，爲什麼呢？ 有沒有辦法確保我的printf數據之前它們在內核（調試）內被修改？

下面是代碼：

~>more * 
:::::::::::::: 
Makefile 
:::::::::::::: 
all: 
    nvcc -o WTF.cu.o -arch=sm_21 -c WTF.cu 
    g++ -o WTF.exe -I/usr/local/cuda/include WTF.cpp WTF.cu.o -L/usr/local/cuda/lib64 -lcuda -lcudart 
:::::::::::::: 
WTF.cpp 
:::::::::::::: 
#include <iostream> // cout 
#include <cstdlib> // rand, srand 

#include <cuda_runtime_api.h> // cudaXXX 
void PrintOnGPU (unsigned int const iDataSize, int * const iopData); 

using namespace std; 

int main() 
{ 
    // Allocate and initialize CPU data 
    unsigned int dataSize = 4; 
    srand (time (NULL)); // Random seed 
    int * pCPUData = (int *) malloc (sizeof (int) * dataSize); 
    for (unsigned int i = 0; i < dataSize; i++) { pCPUData[i] = rand() % 100; cout << "CPU : " << pCPUData[i] << endl; } 

    // Print from GPU 
    int * pGPUData = NULL; 
    cudaMalloc ((void **) &pGPUData, dataSize * sizeof (int)); 
    cudaMemcpy (pGPUData, pCPUData, dataSize * sizeof (int), cudaMemcpyHostToDevice); 
    PrintOnGPU (dataSize, pGPUData); 

    // Get out 
    cudaFree (pGPUData); 
    if (pCPUData) { free (pCPUData); pCPUData = NULL; } 
    return 0; 
} 
:::::::::::::: 
WTF.cu 
:::::::::::::: 
#include "stdio.h" 

__global__ void WTF (unsigned int const iDataSize, int * const iopData) 
{ 
    if (iDataSize == 0 || !iopData) return; 

    // Don't modify : just print 
    unsigned long long int tIdx = blockIdx.x * blockDim.x + threadIdx.x; // 1D grid 
    if (tIdx == 0) 
    { 
    for (unsigned int i = 0; i < iDataSize; i++) 
     printf ("GPU : %i \n", iopData[i]); 
    } 
    __syncthreads(); 

    // Modify 
    // iopData[tIdx] = 666; // WTF ?... 
} 

void PrintOnGPU (unsigned int const iDataSize, int * const iopData) 
{ 
    WTF<<<2,2>>> (iDataSize, iopData); 
} 

而且，正如預期的，我沒有得到任何值在100以上（在CPP文件中第15行：蘭特（）％100）：

~>make; ./WTF.exe 
nvcc -o WTF.cu.o -arch=sm_21 -c WTF.cu 
g++ -o WTF.exe -I/usr/local/cuda/include WTF.cpp WTF.cu.o -L/usr/local/cuda/lib64 -lcuda -lcudart 
CPU : 38 
CPU : 73 
CPU : 28 
CPU : 82 
GPU : 38 
GPU : 73 
GPU : 28 
GPU : 82 

現在我在cu文件中取消註釋第17行（iopData [tIdx] = 666）：我將所有值修改爲666（即100以上）。由於我在CUDA內核中有4個數據（dataSize = 4，在cpp文件中），一個2×2網格和一個__syncthreads（），所以我不應該打印任何修改過的數據，對不對？不過，我得到這個（打印修改的數據與值666）：

~>make; ./WTF.exe 
nvcc -o WTF.cu.o -arch=sm_21 -c WTF.cu 
g++ -o WTF.exe -I/usr/local/cuda/include WTF.cpp WTF.cu.o -L/usr/local/cuda/lib64 -lcuda -lcudart 
CPU : 29 
CPU : 72 
CPU : 66 
CPU : 90 
GPU : 29 
GPU : 72 
GPU : 666 
GPU : 666 

我不明白爲什麼這些666顯示：對我來說，他們不應該出現？！如果這種行爲是正常的，那爲什麼呢？

FH

Answer 1

這是因爲你正在推出2個threadblocks，而那些threadblocks可以以任何順序執行，同時或順序。

假設您有麻煩的行未註釋。現在假設線程塊1先運行並在線程塊0之前完成。然後線程塊0運行。但threadblock 0正在進行打印，並且它正在打印所有4個值。所以之前由線程塊1設置爲666的值是通過線程塊0打印出來的。

如果線程塊0先運行，則不會發生這種情況，相應地，我的猜測是，您從未看到列爲666的前2個GPU值， 2（從線程塊1發出）。不管線程的數量如何（至少在發佈的內核代碼中），如果您只啓動1個塊，您也永遠不會看到它。

您也可能會覺得__syncthreads()是一個設備範圍的同步時感到困惑。不是這樣。它僅作爲線程塊中的線程的屏障。單獨的線程塊之間沒有同步。