Thrust - 訪問鄰居

我想使用Thrust的流壓縮功能（copy_if）從元素中提取元素的索引，如果元素遵守一些約束。其中一個約束取決於相鄰元素的值（2D中8個，3D中26個）。我的問題是：我如何獲得Thrust中元素的鄰居？Thrust - 訪問鄰居

函子爲「copy_if」函數調用操作基本上看起來像：

__host__ __device__ bool operator()(float x) { 
    bool mark = x < 0.0f; 
    if (mark) { 
     if (left neighbor of x > 1.0f) return false; 
     if (right neighbor of x > 1.0f) return false; 
     if (top neighbor of x > 1.0f) return false; 
     //etc. 
    } 
    return mark; 
}

目前我使用一個變通方法，首先推出了CUDA核心（其中很容易訪問鄰居）適當地標記元素。之後，我將標記的元素傳遞給Thrust的copy_if以提取標記元素的索引。

我遇到了counting_iterator作爲直接使用threadIdx和blockIdx獲取處理元素的索引的一種替代方法。我嘗試了下面的解決方案，但在編譯時，它給了我一個「/usr/include/cuda/thrust/detail/device/cuda/copy_if.inl(151）：錯誤：不支持未對齊的內存訪問」。據我所知，我沒有試圖以不對齊的方式訪問內存。任何人都知道發生了什麼和/或如何解決這個問題？

struct IsEmpty2 { 
    float* xi; 

    IsEmpty2(float* pXi) { xi = pXi; } 

    __host__ __device__ bool operator()(thrust::tuple<float, int> t) { 
     bool mark = thrust::get<0>(t) < -0.01f; 
     if (mark) { 
      int countindex = thrust::get<1>(t); 
      if (xi[countindex] > 1.01f) return false; 
      //etc. 
     } 
     return mark; 
    } 
}; 


thrust::copy_if(indices.begin(), 
       indices.end(), 
       thrust::make_zip_iterator(thrust::make_tuple(xi, thrust::counting_iterator<int>())), 
       indicesEmptied.begin(), 
       IsEmpty2(rawXi));

來源

2012-10-04 Bart

您可以在函子類的構造函數中提供全局數組的地址。你也可以訪問threadIdx.x，並且可以在上面的__device__方法中使用共享內存（你應該在你的情況下使用共享內存）。 – phoad

將指向全局數組的指針傳遞給函數的構造函數是可能的，然後將其存儲爲成員變量。但那不會解決我的問題。我應該訪問全局數組的哪些元素？我只有浮動x到我的處置，而不是指向x的指針。至於使用共享內存，在我的情況下，這是不必要的。我只加載數據一次，並檢查它的特定值。 – Bart

Thrust中的Zip迭代器可以解決查找全局數組中float值的位置的問題。您可以使用threadIdx.x值，但需要稍有不同。關於使用共享內存，因爲您需要頂部，左側，右側，底部等，所以每個值所需的數據會有重疊，您可以使用局部性將這些數據從全局內存加載到寄存器一次，即共享內存（雖然全局內存緩存也可以）。 – phoad

@phoad：你說的是共享內存，在我發佈我的回覆之後，它給我留下了深刻的印象，後來認爲緩存可能會幫助我。但你用快速反應打敗了我。然而，if語句在所有情況下的執行率都不到5％，因此使用共享內存或依賴緩存可能會對性能產生微不足道的影響。

元組只支持10個值，所以這意味着我需要在3D情況下爲26個值的元組元組元組。使用元組和zip_iterator已經非常麻煩，所以我會通過這個選項（同樣來自代碼可讀性的立場）。我嘗試了直接在設備函數中使用threadIdx.x等的建議，但Thrust不喜歡這樣。我似乎正在得到一些無法解釋的結果，有時我最終會遇到Thrust錯誤。例如下面的程序產生一個「推力::系統:: SYSTEM_ERROR」與「未指定的發射失敗」，雖然它首先正確打印「處理10」到「處理41」：

struct printf_functor { 
    __host__ __device__ void operator()(int e) { 
     printf("Processing %d\n", threadIdx.x); 
    } 
}; 

int main() { 
    thrust::device_vector<int> dVec(32); 
    for (int i = 0; i < 32; ++i) 
     dVec[i] = i + 10; 

    thrust::for_each(dVec.begin(), dVec.end(), printf_functor()); 

    return 0; 
}

同樣適用於印刷blockIdx.x然後打印blockDim.x不會產生錯誤。我希望能有一個乾淨的解決方案，但我想我堅持使用當前的解決方案。

來源

2012-10-08 10:19:00 Bart

Thrust - 訪問鄰居

回答

相關問題