CUDA Texture Part.2 Linear Memory

CUDA Texture Part.2 Linear Memory

CUDA Texture Part.2 Linear Memory

在前一篇已經大概介紹過 texture 的基本概念了，在這一篇，就來大概看一下，一般使用 linear memory 的 texture 的方法吧。

在宣告的部分，前面已經提過了，大致上就是：

<int, 1, > texRef;

的形式。這樣，就可以宣告出一個一維的整數 texture 了；而由於他的讀取模式是給定為 cudaReadModeElementType，所以之後由這個 texture 取出的值，也都會是整數（如果是 cudaReadModeNormalizedFloat，就會是浮點數）。

而宣告出 texture reference 的物件後，接下來就是要和現有的變數做連結了～

Bind Texture

使用 linear memory 時，是要將在 global memory 中用 cudaMalloc() 定義出來的記憶體空間，透過 cudaBindTexture() 這個函式，來將 texture reference 連結到變數。其函式的型態為：

template<class T, int Dim, enum readMode>
cudaError_t ( size_t* offset,
                         const struct <T, dim, readMode>& texRef,
                         const void* devPtr,
                         size_t size = UINT_MAX);

其中，offset 算是比較進階的設定，在這邊先略過，不過一般是都給 0。而 texRef 就是要 bind 的 texture，devPtr 則是要 bind 到 texRef 的資料；size 就是 devPtr 的記憶體空間大小，不過一般應該是可以省略不給。

而除了 cudaBindTexture() 外，當然也有相反的 cudaUnbindTexture()。他的動作就是解除 texture reference 和變數的關係。

template<class T, int Dim, enum readMode>
cudaError_t (
                const struct <T, dim, readMode>& texRef );

 

 

在 kernel 中使用

在 kernel function 中，要存取使用 linear memory 的 texture reference，要透過 tex1Dfetch() 這個函式；其型態為：

template<class Type>
Type tex1Dfetch( <Type, 1, ReadMode> texRef, int x);

也就是只要給他要讀取的 texture reference，以及要讀取的位置 x，就可以取得資料的值了～

 

 

 

簡單的範例

上面已經把要使用 linear memory 的 CUDA texture 時所需要的基本功能都介紹過了，接下來就來給個簡單的範例吧～程式可以到這個連結下載。

Heresy 這邊用的例子，是把兩張圖做 alpha blending，也就是把兩張圖各自設定透明度，疊在一起看；不過為了簡化，所以沒有去讀圖檔，而是直接建立測試的資料。下面就是 main() 的部分

void main( int argc, char** argv )
{
    int width   = 1920,
        height  = 1200,
        channel = 3;

    
    unsigned char *aImg1 = new unsigned char[ width*height*channel ],
                  *aImg2 = new unsigned char[ width*height*channel ],
                  *aRS1  = new unsigned char[ width*height*channel ],
                  *aRS2  = new unsigned char[ width*height*channel ];
    for( int i = 0; i < width * height * channel; ++ i )
    {
        aImg1[i] = 0;
        aImg2[i] = 255;
    }

    
    Blend_CPU( aImg1, aImg2, aRS1, width, height, channel );
	
    
    Blend_GPU( aImg1, aImg2, aRS2, width, height, channel );

    
    for( int i = 0; i < width * height * channel; ++ i )
        if( aRS1[i] != aRS2[i] )
        {
            printf(  );
            break;
        }
}

在這裡，是建立了四個大小是 1920*1200*3 的 unsigned char 一維陣列來充當圖檔；其中，aImg1 和 aImg2 是當作來源，分別全部填入 0 和 255，而 aRS1 和 aRS2 則分別拿來儲存用 CPU 以及 GPU 計算後的結果。其中 Blend_CPU() 就是用 CPU 計算的函式、Blend_GPU() 則是用 GPU 計算的函式；在計算完後，結果會分別存在 aRS1 和 aRS2 中。最後「check」的部分，就是在驗證 CPU 和 GPU 計算的結果是否一致了～

而其中，Blend_CPU() 的內容是：

void Blend_CPU( unsigned char* aImg1, unsigned char* aImg2,
                unsigned char* aRS,
                int width, int height, int channel )
{
    for( int i = 0; i < width * height * channel; ++ i )
        aRS[i] = (unsigned char)( 0.5 * aImg1[i] + 0.5 * aImg2[i] );
}

可以看到這邊的程式非常簡單，就是用一個 for 迴圈，把整個陣列掃一遍，並把 aImg1 和 aImg2 的值都乘上 0.5 後加起來。

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

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

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

 void Blending_Texture( unsigned char* aRS, int size )
{
    int index =  *  + ;

    if( index < size )
        aRS[index] = 0.5 * ( rT1, index )
                   + 0.5 * ( 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 );

    
    unsigned char	*dev_A,	*dev_B,	*dev_C;
    ( (void**)&dev_A, data_size );
    ( (void**)&dev_B, data_size );
    ( (void**)&dev_C, data_size );

    
    ( dev_A, aImg1, data_size,  );
    ( dev_B, aImg2, data_size,  );

    
    (0, rT1, dev_A );
    (0, rT2, dev_B );

    
    Blending_Texture( dev_C, size );

    
    ( aRS, dev_C, data_size,  );

    
    (rT1);
    (rT2);

    (dev_A);
    (dev_B);
    (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() 這個函式，他的步驟如下：

1. 先把所需要的記憶體大小計算出來
2. [part1] 宣告 dev_A、dev_B、dev_C，並指派記憶體空間；此時，dev_A、dev_B、dev_C 就是使用 global memory 的變數。
3. [part2]  透過 cudaMemcpy() 把資料由 host memory（aImg1、aImg2） 複製到 device memory（dev_A、dev_B）。
4. [part2a] 透過 cudaBindTexture() 將 rT1、rT2 和 dev_A、dev_B 做聯繫。而此時，rT1、rT2  就算是使用 texture memory 的變數。
5. [part3] 呼叫 kernel function：Blending_Texture() 來進行計算了。
6. [part4] 將結果由 device memory（dev_C）複製回 host memory（aRS）。
7. [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 該做的事了。

posted on 2010-06-17 20:20 GXW 阅读( ...) 评论( ...) 编辑 收藏

转载于:https://www.cnblogs.com/Fancyboy2004/archive/2010/06/17/1759751.html