CUDA编程模型是一个异构模型,需要CPU和GPU协同工作.
host和device
host和device是两个重要的概念
- host指代CPU及其内存
- device指代GPU及其内存
__global__: host调用,device上执行
__device__:device调用,device执行
__host__:host调用, host执行
典型编程流程
- 分配host内存,并进行数据初始化
- 分配device内存,并从host将数据拷贝到device上
- 调用CUDA的核函数在device上完成指定的运算
- 将device上的运算结果拷贝到host上
- 释放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);