Координаты текстуры в Cuda - PullRequest
Новая
       34

Координаты текстуры в Cuda

1 голос
/ 26 марта 2012

Я смотрю на эту реализацию DCT с использованием cuda: 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? </p>

1 Ответ

4 голосов
/ 26 марта 2012

Чтобы ответить на ваши вопросы в обратном порядке:

  1. Как и в стандарте C или C ++, оператор << - это <a href="http://en.wikipedia.org/wiki/Bitwise_operation#Shifts_in_C.2C_C.2B.2B.2C_C.23" rel="nofollow"> битовый оператор сдвига влево .Это означает, что 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);

    ......
}
...