CUDA編程學習筆記1

CUDA編程模型是一個異構模型,需要CPU和GPU協同工作.

---

host和device

host和device是兩個重要的概念

• host指代CPU及其內存
• device指代GPU及其內存

__global__: host調用,device上執行

__device__:device調用,device執行

__host__:host調用, host執行

---

典型編程流程

1. 分配host內存,並進行數據初始化
2. 分配device內存,並從host將數據拷貝到device上
3. 調用CUDA的核函數在device上完成指定的運算
4. 將device上的運算結果拷貝到host上
5. 釋放device和host上的內存

---

核函數

核函數(kernel)是在device上線程中並行的函數.

• 初始化:核函數用__global__符號聲明
• 每一個線程有唯一的縣稱號thread ID,這個用內置變量threadIdx
• 在調用時候用<<<grid,block>>>來指定kernel要執行的線程數量.
  其中,一個kernel所啟用的所有的線程稱為grid,同一個grid上的線程共享相同的全局內存空間,grid又可以分割為很多的block,block里包含很多線程.

dim3 grid(3, 2);
dim3 block(5, 3);
kernel_fun<<<grid, block>>>(params...);

---

內存模型

• 每個線程有自己的local memory
• 每個線程塊(block)有shared memory.可以block中所有的thread共享,其生命周期與block一致.
• 所有的thread都可以訪問全局內存global memory.還可以訪問一些只讀模塊,constant memory 和 texture memory.

---

GPU硬件實現的基本認識

一個kernel會啟動很多線程,這些線程邏輯上是並行的,但是在物理上卻不一定.這個和CPU的多線程有類似支出,多線程如果沒有多核支持,在物理層也是無法實現的.

但是好在GPU存在很多CUDA核心,充分利用CUDA核心可以充分發揮GPU的並行計算能力.

GPU硬件的一個核心組件是SM(streaming multiprocessor),流式多處理器.

SM的核心組件包括CUDA核心,共享內存,寄存器等.

一個線程塊只能在一個SM上被調度。SM一般可以調度多個線程塊，這要看SM本身的能力。

那么有可能一個kernel的各個線程塊被分配多個SM，所以grid只是邏輯層，而SM才是執行的物理層。

由於SM的基本執行單元是包含32個線程的線程束，所以block大小一般要設置為32的倍數。

在進行CUDA編程前，可以先檢查一下自己的GPU的硬件配置，這樣才可以有的放矢，可以通過下面的程序獲得GPU的配置屬性

  int dev = 0;
    cudaDeviceProp devProp;
    CHECK(cudaGetDeviceProperties(&devProp, dev));
    std::cout << "使用GPU device " << dev << ": " << devProp.name << std::endl;
    std::cout << "SM的數量：" << devProp.multiProcessorCount << std::endl;
    std::cout << "每個線程塊的共享內存大小：" << devProp.sharedMemPerBlock / 1024.0 << " KB" << std::endl;
    std::cout << "每個線程塊的最大線程數：" << devProp.maxThreadsPerBlock << std::endl;
    std::cout << "每個EM的最大線程數：" << devProp.maxThreadsPerMultiProcessor << std::endl;
    std::cout << "每個EM的最大線程束數：" << devProp.maxThreadsPerMultiProcessor / 32 << std::endl;

    // 輸出如下
    使用GPU device 0: GeForce GT 730
    SM的數量：2
    每個線程塊的共享內存大小：48 KB
    每個線程塊的最大線程數：1024
    每個EM的最大線程數：2048
    每個EM的最大線程束數：64

---

加法實例

cudaError_t cudaMalloc(void** devPtr, size_t size);
cudaError_t cudaMemcpy(void* dist, const void* src, size_t count, cudaMemcpyKind kind);

其中cudaMemcpyKind是一個enum

enum cudaMemcpyKind {
    cudaMemcpyHostToHost,
    cudaMemcpyHostToDevice,
    cudaMemcpyDeviceToHost,
    cudaMemcpyDeviceToDevice
};

// -- grid 和 block 都是1-dim, 先定義kernel
__global__ void add(float* x, float* y, float* z, int n) {
    int index = threadIdx.x + blockIdx.x * blockDim*x;
    int stride = blockDim.x * gridDim.x; // -- 整個grid的總線程數
    for (int i = index; i < n; i += stride) {
        z[i] = x[i] + y[i];
    }
}

int main() {
    int N = 1 << 20;
    int nBytes = N * sizeof(float);
    // 申請host內存
    float *x, *y, *z;
    x = (float*)malloc(nBytes);
    y = (float*)malloc(nBytes);
    z = (float*)malloc(nBytes);
    
    // -- init data
    for (int i = 0; i < N; ++i) {
        x[i] = 10.0;
        y[i] = 20.0;
    }
    
    // --申請device內存
    float *d_x, *d_y, *d_z;
    cudaMalloc((void**)&d_z, nBytes);
    cudaMalloc((void**)&d_y, nBytes);
    cudaMalloc((void**)&d_z, nBytes);
    // -- host copy to device
    cudaMemcpy((void*)d_x, (void*)x, nBytes, cudaMemcpyHostToDevice);
    cudaMemcpy((void*)d_y, (void*)y, nBytes, cudaMemcpyHostToDevice);
    // -- 定義kernel的執行配置
    dim3 blockSize(256);
    dim3 gridSize((N + blockSize.x - 1) / blockSize.x);
    // -- 執行kernel
    add <<<gridSize, blockSize>>>(d_x, d_y, d_z, N);
    // -- 
}

---

#include <iostream>
#include <time.h>
#include "opencv2/highgui.hpp"
#include "opencv2/opencv.hpp"
using namespace cv;
using namesapce std;

__global__ void rgb2grayincuda(uchar3* const d_in, unsigned char* const d_out, uint imgheight, uint imgwidth) {
    const unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
    const unsigned int idy = blocKIdx.y * blockDim.y + threadIdx.y;
    
    if (idx < imgwidth && idy < imgheight) {
        uchar3 rgb = d_in[idy * imgwidth + id];
        d_out[idy * imgwidth + idx] = 0.229f * rgb.x + 0.587f * rgb.y + 0.114f * rgb.z;
    }
}

int main(void) {
    Mat srcImage = imread("./test.jpg");
    
    const uint imgheight = srcImage.rows;
    const uint imgwidth = srcImage.cols;
    
    Mat grayImage(imgheight, imgwidth, CV_8UC1, Scalar(0));
    
    uchar3 *d_in;
    unsighed char * d_out;
    cudaMalloc((void**)&d_in, imgheight * imgwidth * sizeof(uchar3));
    cudaMalloc((void**)&d_out, imgheight * imgwidht * sizeof(unsigned char));
    
    cudaMemcpy(d_in, srcImage.data, imgheight * imgwidth * sizeof(uchar3), cudaMemcpyHostToDevice);
    
    dim3 threadsPerBlock(32, 32);
    dim3 blocksPerGrid((imgwidth + threadsPerBlock.x - 1) / (threadPerBlock.x,, (imgheight + threadPerBlock.y - 1) / threadsPerBlock.y);
    
    rgb2grayincuda <<<blocksPerGrid, threadsPerBlock>>>(d_in, d_out, imgheight, imgwidth);
                       
    cudaDeviceSynchronize();
                       
}

CMakeLists.txt

cmake_minumum_requred(VERSION 2.8)
project(testcuda)
find_package(CUDA REQUIRED)
find_package(OpenCV REQUIRED)
cuda_add_executable(testcuda main.cu)
target_link_libraries(testcuda ${OpenCV_LIBS})

---

設備內存

CUDA運行庫提供了函數以分配/釋放設備端的內存，以及與主機端內存傳輸數據。

這里的設備內存，指的是全局內存+常量內存+紋理內存。

線性內存是我們常用的內存方式，在GPU上用40位的地址線尋址。線性內存可以用cudaMalloc()分配，用cudaFree()釋放，用cudaMemcpy()復制數據，用cudaMemset()賦值。

對於2D或3D數組，可以使用cudaMallocPitch()和cudaMalloc3D()來分配內存。這兩個函數會自動padding，以滿足內存對齊的要求，提高內存讀寫效率。內存對齊的問題，會在第五章里詳細闡述。

另外，如果要在設備內存中定義全局變量，則需要使用使用__constant__或__device__來修飾，並使用cudaMemcpyToSymbol()和cudaMemcpyFromSymbol()來讀寫。如下例:

__constant__ float constData[256];
float data[256];
cudaMemcpyToSymbol(constData, data, sizeof(data));
cudaMemcpyFromSymbol(data, constData, sizeof(data));

__device__ float devData;
float value = 3.14f;
cudaMemcpyToSymbol(devData, &value, sizeof(float));

__device__ float* devPointer;
float* ptr;
cudaMalloc(&ptr, 256 * sizeof(float));
cudaMemcpyToSymbol(devPoint, &ptr, sizeof(ptr));

實際上，當使用__constant__關鍵字時，是申請了一塊常量內存；而使用__device__時，是普通的全局內存。因此__device__申請的內存需要申請，而__constant__不用。不管是全局內存，還是常量內存，需要用帶有Symbol的函數拷貝。

---

Texture

enum cudaTextureAddressMode {
    cudaAddressModeWrap, // -- warpping address mode
    cudaAddressModeClamp, // -- 將超出坐標截斷為最大值或最小值,即返回圖像邊緣像素值
    cudaAddressModeMirror, // -- 將圖像看成周期函數訪問
    cudaAddressModeBorder // -- 如果超出邊緣就返回0
};

enum cudaTextureFilterMode {
    cudaFilterModePoint,	// -- point filter mode 最近領插值
    cudaFilterModeLinear	// -- linear filter mode	雙線性插值 必須配合float使用
};

enum cudaTextureReadMode {
    cudaReadModeElementType,		// -- read texture as specifed element type
    cudaReadModeNormalizedFloat		// -- read texture as normalized float
}

紋理的聲明

texture<Datatype, Type, ReadMode> texRef;
// Datatype, 數據類型, uchar, float, double
// Type, 紋理維度, Type = 2(二維)
// ReadMode, 訪問模式, 
enum cudaTextureFilterMode filterMode;

---

關於cudaMalloc

cudaError_t cudaMalloc(void ** devPtr, size_t size);
cudaError_t cudaMalloc3D(struct cudaPitchedPtr* pitchedDevPtr, struct cudaExtext extext);
cudaError_t cudaMallocArray(struct cudaArray** array, const struct cudaChannelFormatDesc* desc, size_t width, size_t height, unsigned int flags = 0);