CUDA 5.0重播開銷

我是新手CUDA程序員。我最近更多地瞭解到在較低的入住率下取得更好的表現。下面的代碼片段，我需要幫助理解一些事情有關回放開銷和指令級ParallellismCUDA 5.0重播開銷

__global__ void myKernel(double *d_dst, double *d_a1, double *d_a2, size_t SIZE) 
{ 

    int tId = threadIdx.x + blockDim.x * blockIdx.x; 

    d_dst[tId]   = d_a1[tId] * d_a2[tId]; 
    d_dst[tId + SIZE]  = d_a1[tId + SIZE] * d_a2[tId + SIZE]; 
    d_dst[tId + SIZE * 2] = d_a1[tId + SIZE * 2] * d_a2[tId + SIZE * 2]; 
    d_dst[tId + SIZE * 3] = d_a1[tId + SIZE * 3] * d_a2[tId + SIZE * 3]; 
}

這是我簡單的內核，它只是乘以2個二維陣列，形成第三個二維數組（從邏輯的角度）這些陣列都放置在設備內存中作爲平面1D陣列。

下面我將介紹另一塊代碼片斷：

void doCompute() { 

    double *h_a1; 
    double *h_a2; 

    size_t SIZE = pow(31, 3) + 1; 

    // Imagine h_a1, h_a2 as 2D arrays 
    // with 4 rows and SIZE Columns 
    // For convenience created as 1D arrays 

    h_a1 = (double *) malloc(SIZE * 4 * sizeof(double)); 
    h_a2 = (double *) malloc(SIZE * 4 * sizeof(double)); 

    memset(h_a1, 5.0, SIZE * 4 * sizeof(double)); 
    memset(h_a2, 5.0, SIZE * 4 * sizeof(double)); 

    double *d_dst; 
    double *d_a1; 
    double *d_a2; 

    cudaMalloc(&d_dst, SIZE * 4 * sizeof(double)); 
    cudaMalloc(&d_a1, SIZE * 4 * sizeof(double)); 
    cudaMalloc(&d_a2, SIZE * 4 * sizeof(double)); 

    cudaMemcpy(d_a1, h_a1, SIZE * 4 * sizeof(double), cudaMemcpyHostToDevice); 
    cudaMemcpy(d_a2, h_a2, SIZE * 4 * sizeof(double), cudaMemcpyHostToDevice); 

    int BLOC_SIZE = 32; 
    int GRID_SIZE = (SIZE + BLOC_SIZE - 1)/BLOC_SIZE; 

    myKernel <<< GRID_SIZE, BLOC_SIZE >>> (d_dst, d_a1, d_a2, SIZE); 
}

Q1）我在這裏打破任何凝聚的內存訪問模式？ Q2）我可以說對內存的訪問，它們在內核中的編碼方式也是指令級並行性的例子嗎？如果是，我使用ILP2還是ILP4？和爲什麼？如果我所做的都是正確的，那麼爲什麼nvvp分析器會給我下面的消息？

Total Replay Overhead: 4.6% 
Global Cache Replay Overhead: 30.3%

我該如何減少它們或修復它們？

乾杯，

來源

2013-07-11 Psypher

如果您選擇'SIZE'作爲2的非功率值，將無法提供最佳的融合可能性。但緩存應該有所幫助。我認爲，視覺輪廓儀應報告全球內存使用效率的50％至60％。你的內核代碼應該爲ILP提供一些機會，因爲每行代碼都不依賴於以前的代碼。但ILP也將取決於您使用的特定GPU。如果您將「SIZE」設置爲32或16的倍數，您可能會得到不同/更好的結果。 –

@Robert嗨，我試圖解決的問題的性質不允許我將SIZE精確到2的冪。所以我即興創作，我爲它們添加了額外的內存位置，使SIZE成爲32的倍數。同樣在這種情況下，我添加了1 分析器顯示加載和存儲效率爲100％。但是，這是否意味着它是一個合併的內存訪問？ ILP是否也適用於內存訪問？ warp中的線程正在跳躍，但它們不會破壞128字節的緩存線？或者他們？我正在使用雙打而不是浮動？最後如果一切都在這裏，爲什麼然後開銷？ – Psypher

是的，我的錯誤是，SIZE值可以被32整除，並且應該足以實現良好的合併，因此VP報告的全局內存負載效率爲100％。 ILP和合並實際上是兩個不同的概念。沒有ILP機會的代碼仍然可以實現100％的帶寬利用率，100％的合併，並且利用100％的可用內存帶寬。你的代碼應該100％合併。 –

編譯器有來調度可能ILP開發的指令的能力有限。 GPU本身也必須具有ILP能力，其程度因GPU生成而異。是的，任何不可用的資源都會導致變形停頓，典型的變化是內存中需要的數據。您詢問的重播數量定義爲here。

因此，例如，全局高速緩存重播開銷將由高速緩存未命中觸發，並且您的代碼將有一些高速緩存未命中。即使您擁有100％的聯合訪問和（接近）100％帶寬利用率，緩存未命中也是可能的。

來源

2013-07-11 19:25:50

CUDA 5.0重播開銷

回答

相關問題