Texturkoordinaten in Cuda

Question

Im bei dieser Implementierung von DCT suchen CUDA: http://www.cse.nd.edu/courses/cse60881/www/source_code/dct8x8/dct8x8_kernel1.cu Der Teil in Frage hier ist:

__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(); 

wo Blockgröße 8 ist so block_size_log2 3. ist

Warum das ist Texturkoordinaten definiert wie es ist? Warum müssen wir Texturkoordinaten verwenden? Was ist das "< <" in Cuda?

Answer 1

Um Ihre Fragen in umgekehrter Reihenfolge zu beantworten:

1. Wie in Standard-C oder C++, der < < Operator ist die bitwise left shift operator. Dies bedeutet, dass a << b äquivalent zu a * 2^b ist, wobei a und b beide positive ganze Zahlen sind. Der Code, nach dem Sie fragen, ist im Grunde eine Abkürzung für die ganzzahlige Potenz zweier Multiplikationen.
2. Wie in den Anhängen des Cuda-Programmierleitfadens beschrieben, werden Texturen mithilfe von Voxel-zentrierten Fließkomma-Koordinaten indiziert. Daher werden die Leseargumente in dem von Ihnen geposteten Code in jeder Richtung um 0,5 verschoben.
3. Der Code Sie Fragen über Looks, die für eine frühe Generation von CUDA-Hardware geschrieben wurden, die eine beträchtlich langsamere ganzzahlige arithmetische Leistung als Fließkomma hatte. Die Verwendung der Bitverschiebung anstelle der Potenz von zwei Multiplikationen ist meistens eine Leistungsoptimierung und kann bei neueren CUDA-Hardware-Generationen keinen Nutzen bringen.

Der Code Sie wahrscheinlich gefragt haben, wie

__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); 

    ...... 
} 

geschrieben werden können