CUDA Texture Part.2 Linear Memory-技术开发专区

CUDA Texture Part.2 Linear Memory

作者：heresy live空间编辑：覃里 2009-11-27 17:18 来源：IT168�

　　接下來，就是 GPU 程式 Blend_GPU() 所在的的 .cu 檔了～

#define BLOCK_DIM 512
texture<unsigned char, 1, cudaReadModeElementType> rT1;
texture<unsigned char, 1, cudaReadModeElementType> rT2;

extern "C" void Blend_GPU( unsigned char* aImg1, unsigned char* aImg2,                unsigned char* aRS,                int width, int height, int channel );

__global__ void Blending_Texture( unsigned char* aRS, int size )
{
    int index = blockIdx.x * blockDim.x + threadIdx.x;

    if( index < size )
        aRS[index] = 0.5 * tex1Dfetch( rT1, index )                   + 0.5 * tex1Dfetch( rT2, index );
}

void Blend_GPU( unsigned char* aImg1, unsigned char* aImg2,                unsigned char* aRS,                int width, int height, int channel )
{
    int size = height * width * channel;
    int data_size = size * sizeof( unsigned char );

    // part1, allocate data on device
    unsigned char    *dev_A,    *dev_B,    *dev_C;
    cudaMalloc( (void**)&dev_A, data_size );
    cudaMalloc( (void**)&dev_B, data_size );
    cudaMalloc( (void**)&dev_C, data_size );

    // part2, copy memory to device
    cudaMemcpy( dev_A, aImg1, data_size, cudaMemcpyHostToDevice );
    cudaMemcpy( dev_B, aImg2, data_size, cudaMemcpyHostToDevice );

    // part2a, bind texture
    cudaBindTexture(0, rT1, dev_A );
    cudaBindTexture(0, rT2, dev_B );

    // part3, run kernel
    Blending_Texture<<< ceil((float)size/BLOCK_DIM), BLOCK_DIM >>>                    ( dev_C, size );

    // part4, copy data from device
    cudaMemcpy( aRS, dev_C, data_size, cudaMemcpyDeviceToHost );

    // part5, release data
    cudaUnbindTexture(rT1);
    cudaUnbindTexture(rT2);

    cudaFree(dev_A);
    cudaFree(dev_B);
    cudaFree(dev_C);
}

　　第一行的所定義的 BLOCK_DIM 是定義成每一個 thread block 的大小為 512 個，而如果要執行的 thread 超過這個數值的話，就再切成數個 block 來做;也就是 part3 所指定的執行參數：「<<< ceil((float)size / BLOCK_DIM), BLOCK_DIM >>>」。

　　第二行和第三行是宣告出兩個 CUDA 的 1D texture rT1、rT2 出來，準備之後拿來當輸入用的兩個陣列用;而由於 texture 不能寫入，所以輸出的陣列也就沒必要轉換成 texture 來使用了。而應該是由於目前 CUDA 版本(1.1)的限制，texture reference 只能在 file-scope 宣告成為 global 變數，在 kernel function 中使用。

　　接下來先看 Blend_GPU() 這個函式，他的步驟如下：

先把所需要的記憶體大小計算出來
[part1] 宣告 dev_A、dev_B、dev_C，並指派記憶體空間;此時，dev_A、dev_B、dev_C 就是使用 global memory 的變數。
[part2] 透過 cudaMemcpy() 把資料由 host memory(aImg1、aImg2) 複製到 device memory(dev_A、dev_B)。
[part2a] 透過 cudaBindTexture() 將 rT1、rT2 和 dev_A、dev_B 做聯繫。而此時，rT1、rT2 就算是使用 texture memory 的變數。
[part3] 呼叫 kernel function：Blending_Texture() 來進行計算了。
[part4] 將結果由 device memory(dev_C)複製回 host memory(aRS)。
[part4] 透過 cudaUnbindTexture() 將 rT1、rT2 和 dev_A、dev_B 間的聯繫解除，並使用 cudaFree() 將 device memory 釋放掉。

　　而最後就是這份程式的 kernel function：Blending_Texture() 了～

　　在一開始，還是先利用 CUDA 自動提供的變數 blockIdx、blockDim、threadIdx 來計算出 index 值，並判斷該 thread 是否超出要處理的大小。而之後，就透過 tex1Dfetch() 這個函式，來個別取出 rT1 和 rT2 在 index 的值，並將計算後的結果，存入 aRS[index] 中。如此，就完成了 kernel function 該做的事了。

第1页：使用 linear memory 的 texture 的方法第2页：簡單的範例第3页：GPU 程式 Blend_GPU() 所在的的 .cu 檔

关注我们