優化Bit-Wise操縱內核

我有以下代碼，它逐步遍歷一串位並將它們重新排列爲20個字節的塊。我使用32 * 8塊，每塊有40個線程。然而這個過程在我的GT630M上需要36ms。我可以做什麼進一步的優化？特別是關於刪除最內層循環中的if-else。優化Bit-Wise操縱內核

__global__ void test(unsigned char *data) 
{ 
    __shared__ unsigned char dataBlock[20]; 
    __shared__ int count; 
    count = 0; 

    unsigned char temp = 0x00; 

    for(count=0; count<(streamSize/8); count++) 
    { 
     for(int i=0; i<8; i++) 
     { 
      if(blockIdx.y >= i) 
       temp |= (*(data + threadIdx.x*(blockIdx.x + gridDim.x*(i+count)))&(0x01<<blockIdx.y))>>(blockIdx.y - i); 
      else 
       temp |= (*(data + threadIdx.x*(blockIdx.x + gridDim.x*(i+count)))&(0x01<<blockIdx.y))<<(i - blockIdx.y); 
     } 
     dataBlock[threadIdx.x] = temp; 
      //do something 

    } 

}

來源

2012-07-15 gamerx

您可以從*始終*開始，使用大小爲整經倍數的塊大小。每塊使用40個線程會浪費GPU上所有可用週期的37％。 – talonmies 2012-07-15 15:19:58

此外，您可能希望同時啓動多個線程以隱藏更新。 – 2012-07-15 21:19:50

你的外循環在每次迭代時寫入'dataBlock [threadIdx.x] = temp;'所以你覆蓋了'streamSize/8'的同一個位置。將該線移到循環外部。 – harrism 2012-09-12 00:24:22

目前還不清楚是什麼你的代碼是試圖完成，但一對夫婦明顯的機會：

1）如果可能的話，而不是使用無符號的字符32位字。

2）使用的是32

3）條件代碼可以按預期不會盡可能多的成本你的倍數的塊大小。您可以通過使用--cubin --gpu-architecture sm_xx（其中xx是目標硬件的SM版本）進行編譯來檢查，並在生成的Cubin文件上使用cuobjdump --dump-sass來查看生成的程序集。您可能必須修改源代碼以將公共子表達式放到單獨的變量中，和/或使用三元運算符？：提示編譯器使用預測。

來源

2012-07-17 05:57:03 ArchaeaSoftware

理想情況下，每塊至少使用64個線程。由於每個SM最多可以有8個駐留塊，因此sm_30和更早版本設備上的32個線程塊將限制佔用率（因此可以覆蓋內存延遲）。每個塊有32個線程，因此每個SM最多有8個線程，而如果使用64線程塊，則有16個線程。 – harrism 2012-09-12 00:22:39

優化Bit-Wise操縱內核

回答

相關問題