CUDA atomare Betriebsleistung in verschiedenen Szenarien

Question

Als ich auf SO this question stieß, war ich neugierig, die Antwort zu wissen. Also habe ich unten ein Stück Code geschrieben, um die atomare Betriebsleistung in verschiedenen Szenarien zu testen. Das Betriebssystem ist Ubuntu 12.04 mit CUDA 5.5 und das Gerät ist GeForce GTX780 (Kepler-Architektur). Ich habe den Code mit -O3 Flag und für CC = 3,5 kompiliert.

#include <stdio.h> 

static void HandleError(cudaError_t err, const char *file, int line) { 
    if (err != cudaSuccess) { 
     printf("%s in %s at line %d\n", cudaGetErrorString(err), file, line); 
     exit(EXIT_FAILURE); 
    } 
} 
#define HANDLE_ERROR(err) (HandleError(err, __FILE__, __LINE__)) 

#define BLOCK_SIZE 256 
#define RESTRICTION_SIZE 32 

__global__ void CoalescedAtomicOnGlobalMem(int* data, int nElem) 
{ 
    unsigned int tid = (blockIdx.x * blockDim.x) + threadIdx.x; 
    for (unsigned int i = tid; i < nElem; i += blockDim.x*gridDim.x){ 
     atomicAdd(data+i, 6); //arbitrary number to add 
    } 
} 

__global__ void AddressRestrictedAtomicOnGlobalMem(int* data, int nElem) 
{ 
    unsigned int tid = (blockIdx.x * blockDim.x) + threadIdx.x; 
    for (unsigned int i = tid; i < nElem; i += blockDim.x*gridDim.x){ 
     atomicAdd(data+(i&(RESTRICTION_SIZE-1)), 6); //arbitrary number to add 
    } 
} 

__global__ void WarpRestrictedAtomicOnGlobalMem(int* data, int nElem) 
{ 
    unsigned int tid = (blockIdx.x * blockDim.x) + threadIdx.x; 
    for (unsigned int i = tid; i < nElem; i += blockDim.x*gridDim.x){ 
     atomicAdd(data+(i>>5), 6); //arbitrary number to add 
    } 
} 

__global__ void SameAddressAtomicOnGlobalMem(int* data, int nElem) 
{ 
    unsigned int tid = (blockIdx.x * blockDim.x) + threadIdx.x; 
    for (unsigned int i = tid; i < nElem; i += blockDim.x*gridDim.x){ 
     atomicAdd(data, 6); //arbitrary number to add 
    } 
} 

__global__ void CoalescedAtomicOnSharedMem(int* data, int nElem) 
{ 
    __shared__ int smem_data[BLOCK_SIZE]; 
    unsigned int tid = (blockIdx.x * blockDim.x) + threadIdx.x; 
    for (unsigned int i = tid; i < nElem; i += blockDim.x*gridDim.x){ 
     atomicAdd(smem_data+threadIdx.x, data[i]); 
    } 
} 

__global__ void AddressRestrictedAtomicOnSharedMem(int* data, int nElem) 
{ 
    __shared__ int smem_data[BLOCK_SIZE]; 
    unsigned int tid = (blockIdx.x * blockDim.x) + threadIdx.x; 
    for (unsigned int i = tid; i < nElem; i += blockDim.x*gridDim.x){ 
     atomicAdd(smem_data+(threadIdx.x&(RESTRICTION_SIZE-1)), data[i&(RESTRICTION_SIZE-1)]); 
    } 
} 

__global__ void WarpRestrictedAtomicOnSharedMem(int* data, int nElem) 
{ 
    __shared__ int smem_data[BLOCK_SIZE]; 
    unsigned int tid = (blockIdx.x * blockDim.x) + threadIdx.x; 
    for (unsigned int i = tid; i < nElem; i += blockDim.x*gridDim.x){ 
     atomicAdd(smem_data+(threadIdx.x>>5), data[i>>5]); 

    } 
} 

__global__ void SameAddressAtomicOnSharedMem(int* data, int nElem) 
{ 
    __shared__ int smem_data[BLOCK_SIZE]; 
    unsigned int tid = (blockIdx.x * blockDim.x) + threadIdx.x; 
    for (unsigned int i = tid; i < nElem; i += blockDim.x*gridDim.x){ 
     atomicAdd(smem_data, data[0]); 
    } 
} 

int main(void) 
{ 

    const int n = 2 << 24; 
    int* data = new int[n]; 

    int i; 
    for(i=0; i<n; i++) { 
     data[i] = i%1024+1; 
    } 

    int* dev_data; 
    HANDLE_ERROR(cudaMalloc((void **)&dev_data, sizeof(int) * size_t(n))); 
    HANDLE_ERROR(cudaMemset(dev_data, 0, sizeof(int) * size_t(n))); 
    HANDLE_ERROR(cudaMemcpy(dev_data, data, n * sizeof(int), cudaMemcpyHostToDevice)); 

    for(int i=0; i<50; i++) 
    { 
     dim3 blocksize(BLOCK_SIZE); 
     dim3 griddize((12*2048)/BLOCK_SIZE); //12 SMX ON GTX780 each can have 2048 threads 
     CoalescedAtomicOnGlobalMem<<<griddize, blocksize>>>(dev_data, n); 
     HANDLE_ERROR(cudaPeekAtLastError()); 
    } 
    HANDLE_ERROR(cudaDeviceSynchronize()); 

    for(int i=0; i<50; i++) 
    { 
     dim3 blocksize(BLOCK_SIZE); 
     dim3 griddize((12*2048)/BLOCK_SIZE); //12 SMX ON GTX780 each can have 2048 threads 
     AddressRestrictedAtomicOnGlobalMem<<<griddize, blocksize>>>(dev_data, n); 
     HANDLE_ERROR(cudaPeekAtLastError()); 
    } 
    HANDLE_ERROR(cudaDeviceSynchronize()); 

    for(int i=0; i<50; i++) 
    { 
     dim3 blocksize(BLOCK_SIZE); 
     dim3 griddize((12*2048)/BLOCK_SIZE); //12 SMX ON GTX780 each can have 2048 threads 
     WarpRestrictedAtomicOnGlobalMem<<<griddize, blocksize>>>(dev_data, n); 
     HANDLE_ERROR(cudaPeekAtLastError()); 
    } 
    HANDLE_ERROR(cudaDeviceSynchronize()); 

    for(int i=0; i<50; i++) 
    { 
     dim3 blocksize(BLOCK_SIZE); 
     dim3 griddize((12*2048)/BLOCK_SIZE); //12 SMX ON GTX780 each can have 2048 threads 
     SameAddressAtomicOnGlobalMem<<<griddize, blocksize>>>(dev_data, n); 
     HANDLE_ERROR(cudaPeekAtLastError()); 
    } 
    HANDLE_ERROR(cudaDeviceSynchronize()); 

    for(int i=0; i<50; i++) 
    { 
     dim3 blocksize(BLOCK_SIZE); 
     dim3 griddize((12*2048)/BLOCK_SIZE); //12 SMX ON GTX780 each can have 2048 threads 
     CoalescedAtomicOnSharedMem<<<griddize, blocksize>>>(dev_data, n); 
     HANDLE_ERROR(cudaPeekAtLastError()); 
    } 
    HANDLE_ERROR(cudaDeviceSynchronize()); 

    for(int i=0; i<50; i++) 
    { 
     dim3 blocksize(BLOCK_SIZE); 
     dim3 griddize((12*2048)/BLOCK_SIZE); //12 SMX ON GTX780 each can have 2048 threads 
     AddressRestrictedAtomicOnSharedMem<<<griddize, blocksize>>>(dev_data, n); 
     HANDLE_ERROR(cudaPeekAtLastError()); 
    } 
    HANDLE_ERROR(cudaDeviceSynchronize()); 

    for(int i=0; i<50; i++) 
    { 
     dim3 blocksize(BLOCK_SIZE); 
     dim3 griddize((12*2048)/BLOCK_SIZE); //12 SMX ON GTX780 each can have 2048 threads 
     WarpRestrictedAtomicOnSharedMem<<<griddize, blocksize>>>(dev_data, n); 
     HANDLE_ERROR(cudaPeekAtLastError()); 
    } 
    HANDLE_ERROR(cudaDeviceSynchronize()); 

    for(int i=0; i<50; i++) 
    { 
     dim3 blocksize(BLOCK_SIZE); 
     dim3 griddize((12*2048)/BLOCK_SIZE); //12 SMX ON GTX780 each can have 2048 threads 
     SameAddressAtomicOnSharedMem<<<griddize, blocksize>>>(dev_data, n); 
     HANDLE_ERROR(cudaPeekAtLastError()); 
    } 
    HANDLE_ERROR(cudaDeviceSynchronize()); 

    HANDLE_ERROR(cudaDeviceReset()); 
    printf("Program finished without error.\n"); 
    return 0; 
} 

Grundsätzlich in obigen Code gibt es 8 Kerne, in denen alle Gewinde atomicAdd auf alle Daten zu tun.

1. Koaleszierte atomare Addition auf dem globalen Speicher.
2. Atomare Addition auf einem eingeschränkten Adressraum im globalen Speicher. (32 im Code)
3. Atomare Addition für Kettspuren an der gleichen Adresse im globalen Speicher.
4. Atomare Addition aller Threads an der gleichen Adresse im globalen Speicher.

Die Einträge 5 bis 8 können durch Ersetzen von global mit den oben genannten Elementen gefunden werden. Die gewählte Blockgröße ist 256.

Ich benutzte nvprof, um das Programm zu profilieren. Die Ausgabe lautet:

Time(%)  Time  Calls  Avg  Min  Max Name 
44.33% 2.35113s  50 47.023ms 46.987ms 47.062ms SameAddressAtomicOnSharedMem(int*, int) 
31.89% 1.69104s  50 33.821ms 33.818ms 33.826ms SameAddressAtomicOnGlobalMem(int*, int) 
10.10% 535.88ms  50 10.718ms 10.707ms 10.738ms WarpRestrictedAtomicOnSharedMem(int*, int) 
3.96% 209.95ms  50 4.1990ms 4.1895ms 4.2103ms AddressRestrictedAtomicOnSharedMem(int*, int) 
3.95% 209.47ms  50 4.1895ms 4.1893ms 4.1900ms AddressRestrictedAtomicOnGlobalMem(int*, int) 
3.33% 176.48ms  50 3.5296ms 3.5050ms 3.5498ms WarpRestrictedAtomicOnGlobalMem(int*, int) 
1.08% 57.428ms  50 1.1486ms 1.1460ms 1.1510ms CoalescedAtomicOnGlobalMem(int*, int) 
0.84% 44.784ms  50 895.68us 888.65us 905.77us CoalescedAtomicOnSharedMem(int*, int) 
0.51% 26.805ms   1 26.805ms 26.805ms 26.805ms [CUDA memcpy HtoD] 
0.01% 543.61us   1 543.61us 543.61us 543.61us [CUDA memset] 

Offensichtlich koalesziert konfliktfreie atomare Operationen die beste Leistung hatten, und gleiche Adresse hatte das das Schlimmste. Eine Sache, die ich nicht erklären konnte, war, dass die gleiche atomare Adresse im Shared Memory (innerhalb eines Blocks) langsamer ist als im globalen Speicher (der zwischen allen Threads üblich ist).
Wenn alle Warpspuren in shared memory auf die gleiche Stelle zugreifen, ist die Leistung sehr schlecht, aber es ist (überraschend) nicht der Fall, wenn sie es auf globalen Speicher ausführen. Ich kann nicht erklären warum. Ein weiterer Verwirrungsfall ist, dass atomare und globale Adressen schlechter funktionieren, als wenn alle Threads innerhalb des Warp es an derselben Adresse ausführen, während es scheint, dass Speicherkonflikte im ersten Fall niedriger sind.

Wie auch immer ich würde mich freuen, wenn jemand oben Profiling Ergebnisse erklären könnte.

Answer 1

Als eine zukunftsgerichtete Aussage können meine Kommentare hier teilweise architekturspezifisch sein. Aber für die vorhandenen Architekturen (bis zu cc 3.5, AFAIK) werden Shared-Memory-Atomics über eine (vom Assembler erzeugte) Codefolge implementiert. Diese Codesequenz, die im gemeinsam genutzten Speicher ausgeführt wird, unterliegt der Serialisierung, wenn mehrere Threads um den Zugriff auf die gleiche Bank/den gleichen Speicherort konkurrieren.

Die RMW-Operation selbst ist in dem Sinne atomar, dass kein anderer Thread die Operation stören kann (dh falsche Ergebnisse erzeugen), aber wenn die Threads eine atomare Operation an einem einzigen gemeinsamen Speicherort ausführen, entsteht die Konkurrenz Serialisierung, was die mit Atomics verbundene Verzögerung verschärft.

Nick Zitat aus the CUDA Handbook:

Im Gegensatz zu globalen Speichern, der atomics mit einzelnen Anweisungen implementiert (entweder GATOM oder GRED, je nachdem, ob der Rückgabewert verwendet wird), wird gemeinsam genutzten Speicher atomics mit expliziter Sperre umgesetzt// semantisch freizugeben, und der Compiler gibt Code aus, der bewirkt, dass jeder Thread diese Sperrvorgänge wiederholt, bis der Thread seine atomare Operation ausgeführt hat.

und:

Achten Sie darauf, Konflikte zu vermeiden oder die Schleife 8-2 in Listing kann bis zu 32 mal durchlaufen werden.

Ich schlage vor, lesen Sie den gesamten Abschnitt 8.1.5, zumindest.