Cuda 中的纹理坐标

Question

我正在使用cuda查看DCT的实现：http://www.cse.nd.edu/courses/cse60881/www/source_code/dct8x8/dct8x8_kernel1.cu 有问题的部分在这里：

__shared__ float CurBlockLocal1[BLOCK_SIZE2];

__global__ void CUDAkernel1DCT(float *Dst, int ImgWidth, int OffsetXBlocks, int OffsetYBlocks)
{
    // Block index
    const int bx = blockIdx.x + OffsetXBlocks;
    const int by = blockIdx.y + OffsetYBlocks;

    // Thread index (current coefficient)
    const int tx = threadIdx.x;
    const int ty = threadIdx.y;

    // Texture coordinates
    const float tex_x = (float)( (bx << BLOCK_SIZE_LOG2) + tx ) + 0.5f;
    const float tex_y = (float)( (by << BLOCK_SIZE_LOG2) + ty ) + 0.5f;

    //copy current image pixel to the first block
    CurBlockLocal1[ (ty << BLOCK_SIZE_LOG2) + tx ] = tex2D(TexSrc, tex_x, tex_y);

    //synchronize threads to make sure the block is copied
    __syncthreads();

其中块大小为 8，因此 block_size_log2 为 3。

为什么要这样定义纹理坐标？为什么我们需要使用纹理坐标？Cuda中的“<<”是什么？

Answer 1

以相反的顺序回答您的问题：

1. 与标准 C 或 C++ 中一样，<< 运算符是按位左移运算符。这意味着a << b等于a * 2^b其中a和b都是正整数。因此，您所询问的代码基本上是二乘法的整数幂的简写。
2. 正如 Cuda 编程指南附录中所讨论的，纹理是使用以体素为中心的浮点坐标来索引的，这就是为什么您发布的代码中的读取参数在每个方向上偏移 0.5
3. 您所询问的代码看起来是为早期一代 CUDA 硬件编写的，该硬件的整数运算性能比浮点要慢得多。使用位移位代替二次乘法很可能是一种性能优化，并且可能对新一代 CUDA 硬件没有用处。

您询问的代码可能会写成

__shared__ float CurBlockLocal1[BLOCK_SIZE2];

__global__ void CUDAkernel1DCT(float *Dst, int ImgWidth, int OffsetXBlocks, int OffsetYBlocks)
{
    // Block index
    const int bx = blockIdx.x + OffsetXBlocks;
    const int by = blockIdx.y + OffsetYBlocks;

    // Thread index (current coefficient)
    const int tx = threadIdx.x;
    const int ty = threadIdx.y;

    // Texture coordinates
    const float tex_x = (float)( (bx * BLOCK_SIZE) + tx ) + 0.5f;
    const float tex_y = (float)( (by * BLOCK_SIZE) + ty ) + 0.5f;

    //copy current image pixel to the first block
    CurBlockLocal1[ (ty * BLOCK_SIZE) + tx ] = tex2D(TexSrc, tex_x, tex_y);

    ......
}