CUDA: comportement printf inattendu

Question

Je ne comprends pas le comportement que j'observe en utilisant printf dans un noyau CUDA. Quelqu'un peut-il nous éclairer là-dessus? Si c'est normal, pourquoi est-ce? Y at-il un moyen de m'assurer que je printf données avant ils sont modifiés dans le noyau (débogage)?

Voici le code:

~>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); 
} 

Et, comme prévu, je reçois pas de valeur supérieure à 100 (ligne 15 dans le fichier cpp: rand()% 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 

Maintenant, je décommenter la ligne 17 dans le fichier cu (iopData [tIdx] = 666): Je modifie toutes les valeurs à 666 (c'est-à-dire au-dessus de 100). Comme j'ai 4 données (dataSize = 4 dans le fichier cpp), une grille 2 X 2 et un __syncthreads() avant la modification des données dans le noyau CUDA, je ne devrais jamais imprimer aucune donnée modifiée, non? Cependant, je reçois ce (impression des données modifiées avec une valeur 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 

Je ne comprends pas pourquoi ces 666 apparaissent: pour moi, ils ne devraient pas apparaître! Si ce comportement est normal, pourquoi?

FH

Answer 1

C'est parce que vous lancez 2 threadblocks et les threadblocks pouvez exécuter dans un ordre quelconque, simultanément ou successivement.

Supposons que la ligne gênante ne soit pas mise en commentaire. Supposons maintenant que threadblock 1 s'exécute d'abord et se termine avant threadblock 0. Alors threadblock 0 s'exécute. Mais le threadblock 0 effectue l'impression et il imprime les 4 valeurs. Les valeurs précédemment définies par threadblock 1 à 666 sont imprimées par threadblock 0.

Cela ne pourrait pas arriver si threadblock 0 s'exécute en premier, donc je suppose que vous ne voyez jamais les 2 premières valeurs GPU répertoriées comme 666, seul le dernier 2 (émanant du bloc de threads 1). Vous ne le verrez jamais si vous lancez seulement 1 bloc, quel que soit le nombre de threads (au moins avec le code noyau affiché).

Vous pouvez également être confus en pensant que __syncthreads() est une synchronisation à l'échelle de l'appareil. Ce n'est pas. Il agit comme une barrière pour les threads dans un threadblock seulement. Il n'y a pas de synchronisation entre les threadblocs séparés.