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

回答

相关问题