CUDA流、纹理绑定和异步内存拷贝

Question

CUDA流、纹理绑定和异步内存拷贝

3

最近我在CUDA中编写一些信号处理代码，通过使用1D纹理和调整我的访问模式，我成功地将性能提升了10倍。(之前我尝试过从全局内存到共享内存的事务对齐预取，但后来发生的非均匀访问模式弄乱了warp→shared cache bank关联(我想)).

现在我面临的问题是，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文档没有提供关于这个问题的信息。如果必须解开这个问题以允许并发访问，似乎我需要创建一些纹理引用，并使用一个切换语句来根据作为内核启动参数传递的流编号进行选择。

不幸的是，CUDA不允许在设备端放置纹理数组，即以下内容无法工作：

texture<...> texarray[N_CUDA_STREAMS];

分层纹理不是一个选择，因为我拥有的数据量只适用于一个普通的1D纹理，而不绑定到CUDA数组中（请参见CUDA 4.2 C编程指南中的表F-2）。

- datenwolf

我不确定，但我认为这样重新绑定纹理会引起问题。然而，OpenCL允许创建纹理数组，因此如果您无法通过CUDA解决问题，可以考虑切换到OpenCL，通常非常简单。 - aland

@aland：你知道有没有类似于CUFFT性能的OpenCL替代品吗？ - datenwolf

我不知道有什么比较成熟的库，但是网络上有很多代码，所以你可能可以找到适合你需求的东西。 - aland

1个回答

网页内容由stack overflow 提供, 点击上面的

可以查看英文原文，
原文链接

- tera · Accepted Answer

实际上，当你在不同的流中使用纹理时，你不能解绑它。

由于隐藏异步内存复制所需的流数量并不需要很大（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]);

    // ...
}