CUDA流，紋理綁定和異步memcpy

在CUDA中編寫一些信號處理我最近在優化它方面取得了巨大的進步。通過使用1D紋理和調整我的訪問模式，我設法獲得了10倍的性能提升。（我以前嘗試從全局到事務對齊的預取到共享內存中，但後來發生的非均勻訪問模式混淆了warp→shared cache bank association（我認爲））。CUDA流，紋理綁定和異步memcpy

所以，現在我面臨的問題是，CUDA紋理和綁定如何與異步memcpy交互。

考慮以下內核

texture<...> mytexture; 

__global__ void mykernel(float *pOut) 
{ 
    pOut[threadIdx.x] = tex1Dfetch(texture, threadIdx.x); 
}

內核在多個流

extern void *sourcedata; 

#define N_CUDA_STREAMS ... 

cudaStream stream[N_CUDA_STREAMS]; 
void *d_pOut[N_CUDA_STREAMS]; 
void *d_texData[N_CUDA_STREAMS]; 

for(int k_stream = 0; k_stream < N_CUDA_STREAMS; k_stream++) { 
    cudaStreamCreate(stream[k_stream]); 

    cudaMalloc(&d_pOut[k_stream], ...); 
    cudaMalloc(&d_texData[k_stream], ...); 
} 

/* ... */ 

for(int i_datablock; i_datablock < n_datablocks; i_datablock++) { 
    int const k_stream = i_datablock % N_CUDA_STREAMS; 
    cudaMemcpyAsync(d_texData[k_stream], (char*)sourcedata + i_datablock * blocksize, ..., stream[k_stream]); 

    cudaBindTexture(0, &mytexture, d_texData[k_stream], ...); 

    mykernel<<<..., stream[k_stream]>>>(d_pOut); 
}

現在我想知道的是，由於只有一個紋理參考推出，當我綁定會發生什麼緩衝區到紋理而其他流的內核訪問該紋理？ cudaBindStream不需要流參數，所以我擔心通過將紋理綁定到另一個設備指針，而運行內核異步訪問所述紋理時，我會將其訪問轉移到其他數據。

CUDA文檔沒有提供任何有關這方面的信息。如果必須解決這個問題以允許併發訪問，那麼似乎我必須創建一些紋理引用，並根據作爲內核啓動參數傳遞的流號，使用switch語句在它們之間進行選擇。

不幸的是CUDA不允許把紋理陣列在設備側，即下不起作用：

texture<...> texarray[N_CUDA_STREAMS];

分層紋理不是一種選擇，因爲數據量我只適用於未綁定到CUDA數組的簡單1D紋理（參見CUDA 4.2 C編程指南中的表F-2）。

來源

2012-09-13 datenwolf

我不確定，但我認爲這種紋理重新綁定會導致問題。但是，OpenCL允許創建紋理數組，所以如果您無法解決您的CUDA問題，則可以考慮切換到OpenCL，它通常非常簡單。 – aland

@aland：你是否知道CUFFT的OpenCL對象具有相似的性能？ – datenwolf

我不知道任何或多或少已建立的圖書館，但互聯網上有很多代碼，所以你可能會找到適合你需要的東西。 – aland

事實上，你不能在不同的流中使用它時解除紋理綁定。

由於流的數量並不需要很大隱藏異步memcpys（2就已經這樣做），你可以使用C++模板給每個流自身的質地：

texture<float, 1, cudaReadModeElementType> mytexture1; 
texture<float, 1, cudaReadModeElementType> mytexture2; 

template<int TexSel> __device__ float myTex1Dfetch(int x); 

template<> __device__ float myTex1Dfetch<1>(int x) { return tex1Dfetch(mytexture1, x); } 
template<> __device__ float myTex1Dfetch<2>(int x) { return tex1Dfetch(mytexture2, x); } 


template<int TexSel> __global__ void mykernel(float *pOut) 
{ 
    pOut[threadIdx.x] = myTex1Dfetch<TexSel>(threadIdx.x); 
} 


int main(void) 
{ 
    float *out_d[2]; 

    // ... 

    mykernel<1><<<blocks, threads, stream[0]>>>(out_d[0]); 
    mykernel<2><<<blocks, threads, stream[1]>>>(out_d[1]); 

    // ... 
}

來源

2012-09-13 18:42:36 tera

如何使用模板紋理參考？ – datenwolf

我正在想像（完全未經測試！） – tera

您的評論似乎缺少一些東西。 – datenwolf

CUDA流，紋理綁定和異步memcpy

回答

相關問題