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);