NO IMAGE

CUDA陣列(CUDA Array)

參考:
https://blog.csdn.net/hhko12322/article/details/12004329
http://blog.sciencenet.cn/blog-398465-342089.html

引言

CUDA Array is used for the Texture memory.
本篇談一下不同維數的CUDA陣列的申請,賦值,複製和釋放。

CUDA array 在 cuda 中是一個特殊的型別,叫做 cudaArray,在 CUDA 中,他是專門給 texture 用的一種陣列;通過cudaMallocArray()cudaFreeArray()cudaMemcpyToArray()等函式對其進行管理。此外,由於 cudaArray 本身不是template 的型別,所以在通過cudaMallocArray()來申請CUDA Array時,要通過cudaChannelFormatDesc這個特別的方式,來定義CUDA Array的型別。

申請CUDA Array

如:申請一個一維CUDA Array,型別是float 大小是width×height

cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();
cudaArray* cuArray;
cudaMallocArray(&cuArray, &channelDesc, width, height);

cudaChannelFormatDesc 是一個用來描述 fetch 一個 texture 時,返回值的型別;
定義如下:

template<class T> struct cudaChannelFormatDesc cudaCreateChannelDesc< T>();

複製CUDA Array

而CUDA Array 需要主機端的線性記憶體賦值(copy):

cudaError_t cudaMemcpyToArray(struct cudaArray* dstArray,
size_t dstX, size_t dstY,
const void* src, size_t count,
enum cudaMemcpyKind kind);

cudaMemcpyKind是用來指定賦值的方向,有cudaMemcpyHostToHostcudaMemcpyHostToDevicecudaMemcpyDeviceToHostcudaMemcpyDeviceToDevice四種值。

繫結CUDA Array

CUDA Array是要被Texture使用的,需要繫結到Texture cudaBindTextureToArray()

template<class T, int dim, enum cudaTextureReadMode readMode>
cudaError_t cudaBindTextureToArray(
const struct texture<T, dim, readMode>& texRef,
const struct cudaArray* cuArray);

解繫結:cudaUnbindTexture()

取值

而在存取上,和 linear memory 的 texture 時的 tex1Dfetch() 不同,是要使用 tex1D()、tex2D() 這兩種,分別是用在 1D 和 2D 的 texture。其形式分別為:

template<class Type, enum cudaTextureReadMode readMode>
Type tex1D(texture<Type, 1, readMode> texRef, float x);
template<class Type, enum cudaTextureReadMode readMode>
Type tex2D(texture<Type, 2, readMode> texRef, float x, float y);

三維CUDA Array

上例描述了CUDA Array使用的全過程,其中CUDA Array是2維的,下面介紹3維的CUDA Array,為什麼不介紹一維的CUDA Array?因為Texture本身對區域性記憶體訪問有優化,並且紋理記憶體一般都是基於二維的(圖片都是二維的)。

在當前的CUDA版本中,3D的線性記憶體是無法直接繫結到texture memory,一維的可以,因此,需要將資料首先放進一個3D的CUDA array,然後將3D CUDA array繫結到texture memory上,訪問陣列元素時,通過取紋理的函式tex3D(tex,x,y,z)可以返回座標為(x,y,z)的元素。

  1. 建立CUDA 3D array
    在之前的CUDA版本中,extent.width與height,depth不同,其計數單位為bytes,所以在舊版本中必須使用array_width*sizeof(float),最新的3.1竟然悄悄的修改了。可以CUDA的文件一直是錯誤的,文件中記載width,height,depth均是in bytes,實際上賦值時使用元素個數即可。如果不直接賦值,還可以呼叫函式make_cudaExtent(extent,width,height,depth), 原理類似。

    cudaArray *d_u
    cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc< float>();
    cudaExtent extent;
    extent.width=array_width;
    extent.height=array_height;
    extent.depth=array_depth;
    cudaMalloc3DArray(&d_u,&channelDesc,extent);

  2. 複製資料至3D array
    首先解釋一下pitched pointer的工具原理,如果訪問陣列元素u[x][y][z],通過pitched pointer訪問則是u_p[x y*pitch z*pitch*height ]。 顯然,這裡pitch=width,因此當建立pitched pointer時我們需要將width和height作為引數傳遞給函式make_cudaPitchedPtr()。在這裡尤其要注意的是,pitched pointer指向的array與傳統的C語言陣列的儲存方式不同,C語言訪問元素u[x][y][z]是通過u[x*height*depth y*depth z]。因此為了正確讀取所需元素,我建議逆序建立pitched pointer:
    copyParams.srcPtr = make_cudaPitchedPtr((void*)u, array_depth*sizeof(float), array_depth, array_height);
    此時相當於陣列u[x][y][z]被轉置,在CUDA3D array中對應元素為u[z][y][x],CUDA文件與指南中並未提及這一點區別,這個問題當時也困擾我很久,費盡周折才搞清楚,希望以後的SDK sample能覆蓋這個注意點。

    cudaMemcpy3DParms copyParams = {0};
    copyParams.srcPtr = make_cudaPitchedPtr((void*)u, array_width*sizeof(float), array_width, array_height);
    copyParams.dstArray = d_u;
    copyParams.extent = extent;
    copyParams.kind = cudaMemcpyHostToDevice;
    cudaMemcpy3D(&copyParams);

  3. 繫結3D array至texture memory

    normalized 設定是否對紋理座標是否進行歸一化。如果normalized是一個非零值,那麼就會使用歸一化到[0,1)的座標進行定址,否則對尺寸為width, height, depth的紋理使用座標[0,width-1], [0,height-1], [0,depth-1]定址。例如,一個尺寸為64×32的紋理可以通過x維度範圍為[0,63],y維度範圍[0,31]的座標定址。如果採用歸一化方式對尺寸為64×32的紋理進行定址,在x和y維度上的座標就都是[0.0,1.0)。這樣就可以保證紋理的座標與紋理的尺寸無關。
    filterMode用於設定紋理的濾波模式,即如何根據座標計算返回的紋理值。濾波模式可以是cudaFilterModePoint或者cudaFilterModeLinear。濾波模式為CudaFilterModePoint時,返回值是與座標最接近的像元的值。CudaFilterModeLinear模式只能對返回值為浮點型的紋理使用,啟用這一種模式時將拾取紋理座標周圍的像元,然後根據座標與這些像元之間的距離進行插值計算。對一維紋理可以使用線性濾波,對二維紋理可以使用雙線性濾波。返回值會是對最接近紋理座標的兩個像元(對一維紋理),四個像元(對二維紋理)或者八個像元(對三維紋理)進行插值後得到的值。

    texture< float,3,cudaReadModeElementType> tex_u;
    tex_u.filterMode = cudaFilterModePoint;
    tex_u.normalized = false;
    tex_u.channelDesc = channelDesc;
    if (cudaBindTextureToArray(tex_u, d_u, channelDesc) != (unsigned int) CUDA_SUCCESS) {
    printf(“[ERROR] Could not bind texture un”);
    return;
    }

接下來會介紹,三通道的RGB影像怎麼處理,以及Texture Layer怎麼使用。