希爾利斯和斯蒂爾：內核功能

Question

有人可以幫助瞭解如何希爾利斯&斯蒂爾：爲每個線程執行核函數的工作？

__global__ void scan(float *g_odata, float *g_idata, int n) 
{ 
    extern __shared__ float temp[]; // allocated on invocation 
    int thid = threadIdx.x; 
    int pout = 0, pin = 1; 
    // load input into shared memory. 
    // This is exclusive scan, so shift right by one and set first elt to 0 
    temp[pout*n + thid] = (thid > 0) ? g_idata[thid-1] : 0; 
    __syncthreads(); 
    for (int offset = 1; offset < n; offset *= 2) 
    { 
     pout = 1 - pout; // swap double buffer indices 
     pin = 1 - pout; 
     if (thid >= offset) 
     temp[pout*n+thid] += temp[pin*n+thid - offset]; 
     else 
     temp[pout*n+thid] = temp[pin*n+thid]; 
    __syncthreads(); 
    } 
    g_odata[thid] = temp[pout*n+thid1]; // write output 
} 

由於現在我明白了以下幾點：首先，我們有pout=0, pin=1 and thid = [1,bockDim.x]。 所以直到第一個同步，我們有一個簡單的右移，e.g如果我們有數組[1 | 2 | 5 | 7 ]新數組是[0 |1 | 2 | 5 | 7 ]。

我認爲for loop的多個實例的執行，每個實例每個thId。例如，如果thId=0我們將有以下執行：

1. thid=0

   • offset=1
   • 噘嘴= 1-0 = 1（使用噘嘴初始化之初功能）
   • 銷= 1 - 1 = 0; （使用POUT，這只是計算，E.I 1）
   • 溫度[4] =溫度[0]（else語句）

   • [0 | 1 | 2 | 5 | 0]

   • offset=2

   • POUT = 1-1 = 0（使用從在循環前一步驟POUT）
   • 銷= 1 - 0 = 1; （剛剛計算值）
   • 溫度[0] =溫度[4]（else語句）
   • [0 | 1 | 2 | 5 | 0]

   的撅嘴和腳變量改變基於這些信息在for循環中，而不是
   考慮在開始這些變量的初始化。用相同的方法
   我想象執行thid=1。

2. thid=1

   • offset=1
   • POUT = 1 - 0 = 1（使用POUT初始化的在函數的開頭）
   • 銷= 1 - 1 = 0
   • 溫度[4 + 1] =溫度[0 + 1-1]（if語句）???? 臨時存儲器出界 ????

任何人都可以提供如何，它是執行的直觀的例子？另外當最後的代碼語句將要執行pout的哪個值將被使用？

Answer 1

如果您指的是並行掃描算法，您可以在這裏看到一個直觀的解釋。

http://en.wikipedia.org/wiki/Prefix_sum

而且我相信這些鏈接也有幫助。

http://nvlabs.github.io/cub/structcub_1_1_device_scan.html

http://nvlabs.github.io/cub/classcub_1_1_block_scan.html

http://nvlabs.github.io/cub/namespacecub.html#abec44bba36037c547e7e84906d0d23aba0fa6cac57b7df2f475a67af053b9371c