CUDA 紋理對(duì)象實(shí)現(xiàn)二維卷積

? ? ? ?在筆記 2.0.2?中我們用普通的方法實(shí)現(xiàn)了二維卷積運(yùn)算，本次主要通過(guò)紋理對(duì)象調(diào)用紋理內(nèi)存實(shí)現(xiàn)二維卷積，并對(duì)筆記 1.0.4-3 中紋理內(nèi)存的使用細(xì)則做更詳細(xì)的說(shuō)明。

1 申請(qǐng)數(shù)組

? ? ? ?在 2.0.3 的方法中，我們對(duì)卷積前后的數(shù)據(jù)分別申請(qǐng)了 host 端和 device 端數(shù)據(jù)，采用紋理內(nèi)存時(shí)，不再分配卷積前數(shù)據(jù)的 device 端內(nèi)存，而是將 host 端數(shù)據(jù)傳入 CUDA 數(shù)組。申請(qǐng)數(shù)組的方式如下：

cudaChannelFormatDesc channelDesc?

? ? ? ? ? ? ? ? ? ? = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindUnsigned);

//這個(gè)函數(shù)中注意兩個(gè)參數(shù)，第一個(gè)32是指數(shù)據(jù)大小，常用的int 和float都是32位的，最后一個(gè)參數(shù)有3種：cudaChannelFormatKindUnsigned→無(wú)符號(hào)整型；cudaChannelFormatKindsigned→帶符號(hào)整型；cudaChannelFormatKindFloat→浮點(diǎn)型。

cudaArray_t cuArray;

cudaMallocArray(&cuArray, &channelDesc, width, height);? ? ?//width為列數(shù)，height為行數(shù)

2 拷貝數(shù)據(jù)

? ? ? ?數(shù)據(jù)拷貝有許多函數(shù)，同時(shí)不同的函數(shù)也可以實(shí)現(xiàn)相同的拷貝需求，參考下方網(wǎng)址：(https://developer.download.nvidia.cn/compute/DevZone/docs/html/C/doc/html/group__CUDART__MEMORY.html)，本次拷貝方式如下：

const size_t spitch = width * sizeof(int);

cudaMemcpy2DToArray

? ? ? ??(cuArray, o1, o2, data, spitch, width * sizeof(int), height,?cudaMemcpyHostToDevice);

? ? ? ?cudaArray為目標(biāo)數(shù)組，o1為width方向起始偏移量，o2為height方向起始偏移量，data為源主機(jī)端數(shù)據(jù)，最后一個(gè)為不同的拷貝類(lèi)型。

3 給出紋理對(duì)象的屬性參數(shù)并創(chuàng)建紋理對(duì)象

? ? ? ?具體說(shuō)明參考學(xué)習(xí)筆記 1.0.4_3。但需要注意以下事項(xiàng)，如果你在紋理內(nèi)存中使用源數(shù)據(jù)，texDesc.filterMode 要使用 cudaFilterModePoint，而不能使用線性濾波；想保持源數(shù)據(jù)數(shù)值 texDesc.readMode 使用 cudaReadModeElementType，否則可以選擇進(jìn)行歸一化。索引的坐標(biāo)由texDesc.normalizedCoords 決定是否歸一化，如果歸一化 (值為1) ，在核函數(shù)中訪問(wèn)時(shí)要使用歸一化之后的坐標(biāo)訪問(wèn)，否則值為0。

? ? ? ?另外，對(duì)于二維卷積而言，使用紋理可以省去擴(kuò)邊操作，紋理可以對(duì)超出索引范圍的數(shù)據(jù)進(jìn)行內(nèi)部處理，我們只需要選擇超出范圍后賦值 0 即可。

4 核函數(shù)訪問(wèn)并將計(jì)算結(jié)果傳回 host 端

(注意：在VS中tex2D函數(shù)會(huì)出現(xiàn)紅色波浪線報(bào)錯(cuò)，直接忽略即可，如果其余部分正確，是可以成功編譯并運(yùn)行的，不要被某些網(wǎng)絡(luò)回答誤導(dǎo))

? ? ? ?在核函數(shù)中使用tex2D等函數(shù)拾取紋理時(shí)，要指明數(shù)據(jù)類(lèi)型，同時(shí)避免出現(xiàn)賦值兩邊的數(shù)據(jù)類(lèi)型矛盾，如：tex2D <int> (? ) 和 tex2D <float>?(??) 。函數(shù)中共包含3個(gè)參數(shù)，第一個(gè)為紋理名稱(chēng)，后面為x和y坐標(biāo) (三維手動(dòng)加1即可)。

5 銷(xiāo)毀紋理釋放CUDA數(shù)組及其余內(nèi)存

完整代碼

#include <stdio.h>

#include <stdlib.h>

#include <math.h>

#include <string.h>

#include "cuda_runtime.h"

#include "device_launch_parameters.h"

const int m = 32, n = 32, p = 3, q = 3;

__constant__ int N[p * q];

__global__ void conv2(int* s, cudaTextureObject_t texObj, int bdy_x, int bdy_y)

{

int i = blockIdx.x * blockDim.x + threadIdx.x;

int j = blockDim.y * blockIdx.y + threadIdx.y;

int i2, j2, r1, r2;

int mids;

if (i < m && j < n) {

? ? for (i2 = 0; i2 < p; i2++) {

? ? ? ? for (j2 = 0; j2 < q; j2++) {

? ? ? ? ? ? r1 = i2 - bdy_x;

? ? ? ? ? ? r2 = j2 - bdy_y;

? ? ? ? ? ? mids += N[i2 * q + j2] * tex2D<int>(texObj, i+r1, j+r2);}}

? ? s[i * n + j] = mids; }

}

int main(int argc, char* argv[])

{

int i, j;

int* a, * d_s;

int B1[p][q], B2[p][q];

int bdy_x = floor(p / 2), bdy_y = floor(q / 2);

cudaMallocHost(&a, m * n * sizeof(int));

cudaMalloc(&d_s, m * n * sizeof(int));

//給出內(nèi)核

for (i = 0; i < p; i++) {

? ? for (j = 0; j < q; j++){

? ? ? ? B1[i][j] = i + j; } }

//給出源數(shù)據(jù)

for (i = 0; i < m; i++) {

? ? for (j = 0; j < n; j++) {

? ? ? ? a[i * n + j] = i+j; } }

//內(nèi)核翻轉(zhuǎn)并拷貝至常量?jī)?nèi)存

for (i = 0; i < p; i++) {

? ? for (j = 0; j < q; j++) {

? ? ? ? B2[i][j] = B1[p - i - 1][q - j - 1]; } }

cudaMemcpyToSymbol(N, B2, p * q * sizeof(int));

//申請(qǐng)數(shù)組并拷貝數(shù)據(jù)

cudaChannelFormatDesc channelDesc =cudaCreateChannelDesc

? ? ? ? ? ? ? ? ? ? ? ? ? ? ? ? ? ? ? ? ? ? ? ?(32, 0, 0, 0, cudaChannelFormatKindUnsigned);

cudaArray_t cuArray;

cudaMallocArray(&cuArray, &channelDesc, n, m);

const size_t spitch = n * sizeof(int);

cudaMemcpy2DToArray

? ? ? ? ? ? ? ? ? ? ? ? ? (cuArray, 0, 0, a, spitch, n * sizeof(int), m, cudaMemcpyHostToDevice);

//定義紋理相關(guān)屬性

struct cudaResourceDesc resDesc;

memset(&resDesc, 0, sizeof(resDesc));

resDesc.resType = cudaResourceTypeArray;

resDesc.res.array.array = cuArray;

struct cudaTextureDesc texDesc;

memset(&texDesc, 0, sizeof(texDesc));

texDesc.addressMode[0] = cudaAddressModeBorder;

texDesc.addressMode[1] = cudaAddressModeBorder;

texDesc.filterMode = cudaFilterModePoint;

texDesc.readMode = cudaReadModeElementType;

texDesc.normalizedCoords = 0;

//創(chuàng)建紋理

cudaTextureObject_t texObj = 0;

cudaCreateTextureObject(&texObj, &resDesc, &texDesc, NULL);

dim3 threads(8, 8);

dim3 blocks((m + threads.x - 1) / threads.x, (n + threads.y - 1) / threads.y);

conv2 << <blocks, threads >> > (d_s, texObj, bdy_x, bdy_y);

cudaDeviceSynchronize();

cudaMemcpy(a, d_s, m * n * sizeof(int), cudaMemcpyDeviceToHost);

cudaDestroyTextureObject(texObj);

cudaFreeArray(cuArray);

cudaFreeHost(a);

cudaFree(d_s);

}