如何計算整個CUDA內核執行的執行線程數？

-1

我想對整個內核執行逐漸計算線程執行的次數。有沒有一個本地計數器或有沒有其他方法來做到這一點？我知道保留一個全局變量並且每個線程增量都不起作用，因爲全局內存中的一個變量不能保證線程的同步訪問。如何計算整個CUDA內核執行的執行線程數？

2013-05-28 erogol

您的意思是內核使用的線程數量？我認爲這是編譯時參數... – KiaMorot

「線程執行次數」是什麼意思？ – talonmies

對於整個內核的執行，任何線程的執行方式。例如，如果5個線程激活2個塊，每個塊被執行兩次，而預期的數量是20.獲得它？ – erogol

也許是這樣的：

__global__ void mykernel(int *current_thread_count, ...){ 
    atomicAdd(current_thread_count, 1); 
    // the rest of your kernel code 
    } 


int main() { 
    int tally, *dev_tally; 
    cudaMalloc((void **)&dev_tally, sizeof(int)); 
    tally = 0; 
    cudaMemcpy(dev_tally, &tally, sizeof(int), cudaMemcpyHostToDevice); 
    .... 
    // set up block and grid dimensions, etc. 
    dim3 grid(...); 
    dim3 block(...) 

    mykernel<<<grid, block>>>(dev_tally, ...); 
    cudaMemcpy(&tally, dev_tally, sizeof(int), cudaMemcpyDeviceToHost); 
    printf("total number of threads that executed was: %d\n", tally); 
    .... 

    return 0; 
    }

你可以閱讀更多關於原子功能here

部分由衆多的意見中表示的混亂的原因是，當mykernel完成後（假設它成功運行）每個人都期望tally最終的值等於grid.x*grid.y*grid.z*block.x*block.y*block.z

來源

2013-05-28 15:48:20

@GregSmith提供了一個更好的答案。 –

有很多方法可以測量線程級別的執行效率iciency。該答案提供了不同收集機制的列表。羅伯特克羅維拉的答案提供了一種手動儀器方法，可以準確地收集信息。類似的技術可以用來收集內核中的散度信息。

推出執行（靜態）

gridDim.x * gridDim.y * gridDim.z * blockDim.x * blockDim.y * blockDim.z

線程數線程數推出

gridDim.x * gridDim.y * gridDim.z * ROUNDUP（（blockDim.x * blockDim.y * blockDim.z），WARP_SIZE）

此數字包括THRE廣告在經紗的使用壽命內不活躍。

這可以使用PM計數器threads_launched收集。

經指令在被執行

計數器inst_executed計數退役執行的經紗的指令數/。

發行

經說明計數器計數inst_issued發行的指令數。 inst_issued> = inst_executed。爲了處理調度以縮小執行單元或處理共享內存和L1操作中的地址分歧，一些指令將被執行多次。

線程指令在被執行

計數器thread_inst_executed計數的執行線程的指令數。指標avg_threads_executed_per_instruction可以使用thread_inst_executed/inst_executed派生。該計數器的最大值是WARP_SIZE。

不預測關閉線程指令在被執行

計算能力2。0及以上的器件使用指令預測來禁止對warp中的線程進行回寫，以作爲短髮散指令序列的性能優化。

計數器not_predicated_off_thread_inst_executed統計所有線程執行的指令數。此計數器僅適用於計算能力3.0及以上的設備。

not_predicated_off_thread_inst_executed < = thread_inst_executed < = WARP_SIZE * inst_executed

這種關係將關閉稍有一些芯片由於thread_inst_executed和not_predicated_off_thread_inst_executed計數器小的錯誤。

廓

Nsight Visual Studio版本2.x的支持，收集上述計數器。

Nsight VSE 3.0支持一個新的指令計數實驗，它可以收集每個SASS指令的統計數據，並以表格形式或旁邊的高級來源，PTX或SASS代碼顯示數據。信息從SASS彙總到高級信息源。捲起的質量取決於編譯器輸出高質量符號信息的能力。建議您始終同時查看源和SASS。本實驗可以收集以下每條指令的統計數據：

a。 inst_executed b。 thread_inst_executed（或主動掩碼） c。 not_predicated_off_thread_inst_executed（主動謂詞掩碼） d。 e。直方圖active_mask e。 predicate_mask的直方圖

Visual Profiler 5.0可以準確地收集上述SM計數器。 nvprof可以收集並顯示每個SM的詳細信息。 Visual Profiler 5.x不支持收集Nsight VSE 3.0中提供的每條指令統計信息。較早版本的Visual Profiler和CUDA命令行分析器可以收集許多上述計數器，但結果可能不如5.0和更高版本的工具那麼精確。

來源

2013-05-29 02:08:25

謝謝你的回答。我有時想知道爲什麼CUDA分析器不提供可以根本不運行任何warp的循環數（例如「類似於'issued1''和'issued issued2''發佈的指令''），這是最有趣的事件進行優化，可能是Erogol之後的事情。 'sm_efficiency'包含了類似的信息，但是由於它似乎是從不同運行中採樣的事件派生的，所以很容易就會被忽略掉10％左右。 – tera

@tera我會在CUPTI/nvprof/VisualProfiler中爲計數器添加一個RFE。 Nsight VSE支持「PM計數器 - 費米/開普勒」原始計數器實驗中的計數器。問題效率實驗餅圖中還顯示了週期百分比（Issux 0指令/活動週期）。 –

謝謝Greg！最後是使用Windows的原因。 ;）非常感謝RFE將它帶入unixoides。 – tera

我不認爲有一種方法可以計算特定路徑分支中的線程數。對於直方圖來說，最好有以下幾種：

PS：直方圖是對每種顏色的像素進行計數。

for (i=0; i<256; i++) // 256 colors, 1 pixel = 1 thread if (threadidx.x == i) Histogramme[i] = CUDA_NbActiveThreadsInBranch() // Threads having i as color

來源

2015-03-17 18:40:16 Piotr

如何計算整個CUDA內核執行的執行線程數？

回答

相關問題