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 文档没有说明这一点。如果必须解开它以允许并发访问，似乎我必须创建许多纹理引用并使用 switch 语句在它们之间进行选择，基于作为内核启动参数传递的流号。

不幸的是，CUDA 不允许在设备端放置纹理数组，即以下方法不起作用：

texture<...> texarray[N_CUDA_STREAMS];

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

4

1 回答 1

5

事实上，你不能在不同的流中使用纹理时解除绑定。

由于流的数量不需要很大来隐藏异步 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-13T18:42:36.090 回答