目录
一、CUDA极简入门教程
本部分只是CUDA 的一个超级简单且不完整的内容，关于CUDA配置和编程，请参考官方文档或其他教程。
1、Kernel
Kernel是在GPU上执行的函数，访问的数据都应该在显存中；函数没有返回值，需用void作为返回类型；语法和C++相同，也能使用C++的一些标准库函数(因为这些库函数有GPU实现，不过函数名字和参数相同而已)。kernel是函数的名字，可以随便改。
__global__ void kernel(参数1,参数2,...){
int i = threadIdx.x + blockIdx.x * blockDim.x; //列坐标
int j = threadIdx.y + blockIdx.y * blockDim.y; //行坐标
}__global__是Kernel的一个标识符，与之相对的还有__host__和__device__；由__host__标识的函数和普通函数无异，在CPU上执行；__device__标识的函数只能有__gloabal__标识的函数调用或者被其他用__device__标识的函数调用。
2、线程组织模型
GPU有很多个流处理器，每个流处理器相互独立，可以执行不同的代码；每个流处理器里面还有很多小核心，这些核心在同一时刻执行相同的代码，不过可以通过索引去访问不同的数据。在CUDA的线程模型里面，这些小核心对应的概念叫做Thread，每个Thread都可以计算出一个全局唯一从0开始的索引(索引可以是一维的，可以是二维的，甚至可以时是三维的)。
下面这图是在官方文档中Copy过来的图，图是二维线程模型的一个例子。
Grid由许多个Block组成，一个Block由许多Thread组成。
- 一维：Block的索引为(b_x)，Thread的索引是(t_x)
- 二维：Block的索引为(b_x, b_y)，Thread的索引是(t_x, t_y)
- 三维：Block的索引为(b_x, b_y, b_z)，Thread的索引是(t_x, t_y, t_z)
b_x, b_y, b_z分别是Block在Grid里的x、 y、 z坐标，分别对应blockIdx.x、blockIdx.y、blockIdx.z。t_x, t_y, t_z分别是Thread在Block里的x、 y、z坐标，分别对应.x、threadIdx.y、threadIdx.z。每个Block在x、 y、z三个方向的大小是blockDim.x、blockDim.y、blockDim.z；每个Grid在x、 y、z三个方向的大小是gridDim.x、gridDim.y、gridDim.z。
blockIdx、threadIdx、blockDim、gridDim是kernel程序中的内置变量。
故对于一维线程模型(Block是一维的)，Thread的全局索引(t_x)就可以使用如下代码来计算：
int t_x = threadIdx.x + blockIdx.x*blockDim.x;对于二维线程模型(Block是二维的)，Thread的全局索引(t_x, t_y)就可以使用如下代码来计算：
int t_x = threadIdx.x + blockIdx.x * blockDim.x; // 列坐标
int t_y = threadIdx.y + blockIdx.y * blockDim.y; // 行坐标对于三维线程模型(Block是三维的)，Thread的全局索引(t_x, t_y, t_z)就可以使用如下代码来计算：
int t_x = threadIdx.x + blockIdx.x * blockDim.x;
int t_y = threadIdx.y + blockIdx.y * blockDim.y;
int t_z = threadIdx.z + blockIdx.z * blockDim.z;3、第一个CUDA程序，一维线程模型示例
Kernel函数的调用形式为：函数名<<<Block数量, Block的大小>>>(参数1, 参数2, ...)
// 命名为main.cu
#include <iostream>
#include <string>
#include <stdexcept>
#include <cuda.h>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
using namespace std;
#define SIZE 100
// 检查cuda函数的返回值,出错的时候抛出异常
#define CE(status,error_msg) \
if (status != cudaSuccess) \
{\
char err_buffer[2048];\
sprintf(err_buffer,"第%d行： %s，错误详细信息>> %s\n",__LINE__,error_msg,cudaGetErrorString(status));\
throw runtime_error(err_buffer);\
}
/**
* Keernel。计算c=a+b
* @param a 数组a
* @param b 数组b
* @param c 数组c
*/
__global__ void add_kernel(int *a, int *b, int *c){
int id = threadIdx.x + blockIdx.x*blockDim.x; // 获取当前thread的索引
if(id<SIZE) //! 要判断是否越界！！！
c[id] = a[id]+b[id];
}
void add_caller(int *a, int *b, int *c, size_t size) {
int *dev_a = nullptr;
int *dev_b = nullptr;
int *dev_c = nullptr;
CE(cudaSetDevice(0),"初始化设备失败");
// 在GPU上分配显存
CE(cudaMalloc(&dev_a,size* sizeof(int)),"分配内存失败");
CE(cudaMalloc(&dev_b,size* sizeof(int)),"分配内存失败");
CE(cudaMalloc(&dev_c,size* sizeof(int)),"分配内存失败");
// 从内存把数据复制到显存
CE(cudaMemcpy(dev_a,a, size * sizeof(int),cudaMemcpyHostToDevice),"复制数据失败");
CE(cudaMemcpy(dev_b,b, size * sizeof(int),cudaMemcpyHostToDevice),"复制数据失败");
// 调用Kernel
size_t numBlock = 2; // 使用两个block
size_t blockSize = size/2; // block的大小是size/2
add_kernel<<<numBlock,blockSize>>>(dev_a,dev_b,dev_c);
// 等待Kernel执行完
CE(cudaThreadSynchronize(),"同步失败");
// 从显存中把数据复制回内存
CE(cudaMemcpy(c,dev_c, size * sizeof(int),cudaMemcpyDeviceToHost),"复制数据失败");
// 释放显存
CE(cudaFree(dev_a),"释放内存失败");
CE(cudaFree(dev_b),"释放内存失败");
CE(cudaFree(dev_c),"释放内存失败");
}
int main() {
std::cout << "Hello, World!" << std::endl;
const size_t size = SIZE;
int a[size],b[size],c[size];
for(int i=0;i<size;i++){
a[i] = i;
b[i] = i;
}
add_caller(a,b,c,size);
for (int i : c) {
cout<< i <<endl;
}
return 0;
}CMakeLists.txt
cmake_minimum_required(VERSION 3.0)
project(OCSample)
set(CUDA_USE_STATIC_CUDA_RUNTIME ON) #这一句解决 cannot find -lopencv_dep_cudart
set(CMAKE_CXX_STANDARD 11)
set(CMAKE_RUNTIME_OUTPUT_DIRECTORY ${CMAKE_SOURCE_DIR})
find_package(CUDA REQUIRED)
message(STATUS "CUDA版本: ${CUDA_VERSION}")
message(STATUS " 头文件目录：${CUDA_INCLUDE_DIRS}")
message(STATUS " 库文件列表：${CUDA_LIBRARIES}")
set(CUDA_NVCC_FLAGS -G;-g;-std=c++11) # nvcc flags
include_directories(${CUDA_INCLUDE_DIRS})
CUDA_ADD_EXECUTABLE(main main.cu)
target_link_libraries(main ${CUDA_LIBRARIES})二、访问GpuMat的每个元素
要访问GpuMat的每一个元素，实现自定义的算法，就得自己重新实现一个Kernel，然后把GpuMat作为参数传进去。但实际上，为了提高程序性能，一般不直接使用GpuMat作为参数，而是使用它的精简版PtrStepSz或者PtrStep代替。
- Block数量只能多不能少，否则有的像素访问不到。
- 观察全局Thread索引是怎么算的。
- 在Kernel里面一定要判断是否越界。当然，rows和cols分别是threadsPerBlock.x和threadsPerBlock.y的整倍数时，不需要判断。
- 访问src的一个元素的方法是src(行坐标, 列坐标)。
// main.cu
#include "common.h"
//---------------------CUDA头文件----------------
#include <cuda.h>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
//---------------------CUDA头文件----------------
/**
* CUDA kernel,在GPU上执行的函数。
* 上千个线程都是执行这个函数，每个Thread根据全局id作为坐标来访问像素
* @param src 类型是PtrStepSz<uchar3>，相当于是GpuMat的精简版
*/
__global__ void kernel(GPU::PtrStepSz<uchar3> src){
int i = threadIdx.x + blockIdx.x * blockDim.x; // thread在x方向的全局索引，也就是列坐标
int j = threadIdx.y + blockIdx.y * blockDim.y; // thread在y方向的全局索引，也就是行坐标
if (i ==0 && j==0)
printf("Grid size: (%d, %d)\n", gridDim.x, gridDim.y); //可用printf来debug
if(j<src.rows && i<src.cols) //! 要判断是否越界！！！
src(j,i) = make_uchar3(j%255,i%255,255); // 访问一个像素的方式是src(行坐标,列坐标)
__syncthreads(); // 同步所有线程
}
int main() {
const int rows = 16*10;
const int cols = 32*10;
if(GPU::getCudaEnabledDeviceCount()==0){
cerr<<"此OpenCV编译的时候没有启用CUDA模块"<<endl;
return -1;
}
// 定义一个纯黑GpuMat
GPU::GpuMat gpuMat(rows,cols,CV_8UC3);
// 每个block共有8*16个线程
dim3 threadsPerBlock(8, 16); // blockDim.x = 8, blockDim.y = 16
printf("Block size: (%d, %d)\n", blockDim.x, blockDim.y);
// 计算竖直需要多少个block
uint block_num_vertical = (rows+threadsPerBlock.y-1)/threadsPerBlock.y;
// 计算水平需要多少个block
uint block_num_horizontal = (cols+threadsPerBlock.x-1)/threadsPerBlock.x;
dim3 numBlocks(block_num_horizontal, block_num_vertical); // gridDim.x=numBlocks.x, gridDim.y=numBlocks.y
// 调用kernel，gpuMat的类型会自动转换为PtrStepSz<uchar3>
kernel<<<numBlocks,threadsPerBlock>>>(gpuMat);
// 从显存把数据下载到内存
Mat local;
gpuMat.download(local);
// 显示
imshow("s",local);
imwrite("s.jpg",local);
waitKey(0);
return 0;
}
kernel的参数是PtrStepSz<uchar3>，类型uchar3是个结构体，有三个分量x,y,z。为什么使用uchar3呢，因为gpuMat的类型是CV_8UC3，就是每个元素有三个分量都是uchar这种类型的，当使用uchar3作为PtrStepSz的元素类型时，src(j,i)刚好返回第(j,i)个像素的引用。再举一个例子，如果gpuMat的类型是CV_32FC2(每个元素有两个float)，那么，kenel的参数类型就应该为PtrStepSz<float2>或者PtrStep<float2>。
我使用cmake编译，CMakeLists.txt如下（common.cpp是个空文件，也可以写你自己的代码）：
cmake_minimum_required(VERSION 3.0)
project(OCSample)
set(CUDA_USE_STATIC_CUDA_RUNTIME ON) #这一句解决 cannot find -lopencv_dep_cudart
set(CMAKE_CXX_STANDARD 11)
set(CMAKE_RUNTIME_OUTPUT_DIRECTORY ${CMAKE_SOURCE_DIR})
find_package(CUDA REQUIRED)
message(STATUS "CUDA版本: ${CUDA_VERSION}")
message(STATUS " 头文件目录：${CUDA_INCLUDE_DIRS}")
message(STATUS " 库文件列表：${CUDA_LIBRARIES}")
set(CUDA_NVCC_FLAGS -G;-g;-std=c++11) # nvcc flags
include_directories(${CUDA_INCLUDE_DIRS})
set(OpenCV_DIR "/usr/local/opencv343-cuda90/share/OpenCV") # 指定OpenCV安装路径来区分不同的OpenCV版本
find_package(OpenCV REQUIRED)
set(OpenCV_LIB_DIR ${OpenCV_INSTALL_PATH}/lib)
message(STATUS "OpenCV版本: ${OpenCV_VERSION}")
message(STATUS " 头文件目录：${OpenCV_INCLUDE_DIRS}")
message(STATUS " 库文件目录：${OpenCV_LIB_DIR}")
message(STATUS " 库文件列表：${OpenCV_LIBS}")
include_directories(${OpenCV_INCLUDE_DIRS})
link_directories(${OpenCV_LIB_DIR})
CUDA_ADD_EXECUTABLE(main main.cu common.h common.cpp)
target_link_libraries(main ${OpenCV_LIBS} ${CUDA_LIBRARIES})编译运行上述代码，最后可得到下图
注：实际上这种一个Thread处理一个像素，对显卡的性能浪费非常严重，要想更高效的使用自定义算法，请参考CUDA官方的文档。
扫一扫
5696
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
