Was genau sind die von NVPROF berichteten Transaktionsmetriken?

Question

Ich versuche herauszufinden, was genau die von "nvprof" gemeldeten Messwerte sind. Genauer gesagt kann ich nicht herausfinden, welche Transaktionen Systemspeicher und Gerätespeicher lesen und schreiben. Ich habe einen sehr einfachen Code geschrieben, um das herauszufinden.

#define TYPE float 
#define BDIMX 16 
#define BDIMY 16 
#include <cuda.h> 
#include <cstdio> 
#include <iostream> 
__global__ void kernel(TYPE *g_output, TYPE *g_input, const int dimx, const int dimy) 
{ 
__shared__ float s_data[BDIMY][BDIMX]; 
    int ix = blockIdx.x * blockDim.x + threadIdx.x; 
    int iy = blockIdx.y * blockDim.y + threadIdx.y; 
    int in_idx = iy * dimx + ix; // index for reading input 
    int tx = threadIdx.x; // thread’s x-index into corresponding shared memory tile 
    int ty = threadIdx.y; // thread’s y-index into corresponding shared memory tile 
    s_data[ty][tx] = g_input[in_idx]; 
    __syncthreads(); 
    g_output[in_idx] = s_data[ty][tx] * 1.3; 
    } 


int main(){ 
    int size_x = 16, size_y = 16; 
    dim3 numTB; 
    numTB.x = (int)ceil((double)(size_x)/(double)BDIMX) ; 
    numTB.y = (int)ceil((double)(size_y)/(double)BDIMY) ; 
    dim3 tbSize; 
    tbSize.x = BDIMX; 
    tbSize.y = BDIMY; 
    float* a,* a_out; 
    float *a_d = (float *) malloc(size_x * size_y * sizeof(TYPE)); 
    cudaMalloc((void**)&a,  size_x * size_y * sizeof(TYPE)); 
    cudaMalloc((void**)&a_out, size_x * size_y * sizeof(TYPE)); 
    for(int index = 0; index < size_x * size_y; index++){ 
     a_d[index] = index; 
    } 
    cudaMemcpy(a, a_d, size_x * size_y * sizeof(TYPE), cudaMemcpyHostToDevice); 
    kernel <<<numTB, tbSize>>>(a_out, a, size_x, size_y); 
    cudaDeviceSynchronize(); 
    return 0; 
} 

Dann starte ich nvprof --metrics alle für die Ausgabe, um alle Metriken zu sehen. Dies ist der Teil, der mich interessiert:

  Metric Name      Metric Description   Min   Max   Avg 
Device "Tesla K40c (0)" 
    Kernel: kernel(float*, float*, int, int) 
    local_load_transactions     Local Load Transactions   0   0   0 
    local_store_transactions     Local Store Transactions   0   0   0 
    shared_load_transactions     Shared Load Transactions   8   8   8 
    shared_store_transactions     Shared Store Transactions   8   8   8 
      gld_transactions     Global Load Transactions   8   8   8 
      gst_transactions     Global Store Transactions   8   8   8 
    sysmem_read_transactions   System Memory Read Transactions   0   0   0 
    sysmem_write_transactions   System Memory Write Transactions   4   4   4 
    tex_cache_transactions    Texture Cache Transactions   0   0   0 
    dram_read_transactions   Device Memory Read Transactions   0   0   0 
    dram_write_transactions   Device Memory Write Transactions   40   40   40 
     l2_read_transactions      L2 Read Transactions   70   70   70 
     l2_write_transactions      L2 Write Transactions   46   46   46 

Ich verstehe die gemeinsamen und globalen Zugriffe. Die globalen Zugriffe sind vereinigt und da es 8 Warps gibt, gibt es 8 Transaktionen. Aber ich kann nicht die Systemspeicher- und Gerätespeicher-Schreibtransaktionsnummern herausfinden.

Answer 1

Es hilft, wenn man ein Modell der GPU Speicherhierarchie mit beiden logischen und physikalischen Räumen, wie die here hat.

Mit Bezug auf die „Übersicht Registerkarte“ Diagramm:

1. gld_transactions bezieht sich auf Transaktionen aus der Kette ausgegeben, um den globalen logischen Raum Targeting. Im Diagramm wäre dies die Zeile von der "Kernel" -Box links zur "globalen" Box rechts davon und die logische Datenbewegungsrichtung wäre von rechts nach links.

2. gst_transactions beziehen sich auf die gleiche Zeile wie oben, aber logisch von links nach rechts. Beachten Sie, dass diese globale Transaktion logisch in einem Cache treffen und danach nicht mehr weiter gehen kann. Vom Standpunkt der Metriken beziehen sich diese Transaktionstypen nur auf die angegebene Linie im Diagramm.

3. dram_write_transactions beziehen sich auf die Linie auf dem Diagramm, das von links nach rechts in dieser Zeile Durchflußeinrichtung Speicher auf der rechten Seite mit der L2-Cache und die logischen Daten verbindet. Da die L2-Cacheline 32 Byte groß ist (während die L1-Cachezeile und die Größe einer globalen Transaktion 128 Byte beträgt), betragen die Speichertransaktionen des Geräts ebenfalls 32 Byte, nicht 128 Byte. Also eine globale Schreibtransaktion, die L1 durchläuft (es ist ein Write-Through-Cache, falls aktiviert) und L2 wird 4 dram_write-Transaktionen generieren. Dies sollte 32 der 40 Transaktionen erklären.

4. Systemspeicher-Transaktionen Ziel Null-Kopie Host-Speicher. Du scheinst das nicht zu haben, also kann ich das nicht erklären.

Beachten Sie, dass in einigen Fällen für einige Metriken, auf einigen GPUs kann der Profiler einige „Ungenauigkeit“ haben, wenn sehr kleine Anzahl von threadblocks starten. Zum Beispiel werden einige Metriken auf einer SM-Basis abgetastet und skaliert. (Gerätespeichertransaktionen sind jedoch nicht in dieser Kategorie). Wenn Sie bei jedem SM unterschiedliche Arbeit geleistet haben (möglicherweise aufgrund einer sehr geringen Anzahl von gestarteten Threadblocks), kann die Skalierung irreführend/weniger genau sein. Wenn Sie eine größere Anzahl von Threadblocks starten, werden diese normalerweise unbedeutend.