CUDA流,纹理绑定和异步memcpy

Question

在CUDA中编写一些信号处理我最近在优化它方面取得了巨大进步.通过使用1D纹理并调整我的访问模式,我设法获得了10倍的性能提升.(我之前尝试过将事务对齐从全局到共享内存的预取,但后来发生的非均匀访问模式搞砸了warp→共享缓存库关联(我认为)).

所以现在我正面临着这个问题,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).

Answer 1

实际上,您仍然无法解除纹理,同时仍在不同的流中使用它.

由于流的数量不需要很大来隐藏异步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]);

    // ...
}