Cmake C++ cuda 混合编程

自从Cmake3.9之后，Cmake就支持了C++ 和 Cuda的混合编译。就可以不用像以往的方法那样find_pacakge()来添加相关依赖，与此同时可以在project选项中采用
project(demo LANGUAGES CXX CUDA)
的方式进行支持，或者在后面的选项中使用
enable_language(CUDA)
的方法，开启对CUDA语言的支持。

示例项目结构

推荐平台：Linux
虽然Cmake跨平台效果不错，但是在Windows平台上容易出现找不到nvcc编译器的情况。所以推荐在Linux平台上编译这个项目

• 项目目录
  • bin
  • build
  • include（头文件存放位置）
    • foo.cuh （CUDA头文件）
  • src（源文件存放位置）
    • foo.cu （CUDA源文件）
    • main.cpp （主函数）
  • CMakeLists.txt（Cmake文件）

项目代码

CMakeLists.txt的写法

CMakeLists.txt

cmake_minimum_required(VERSION 3.9)#Cmake最低版本
#project(demo LANGUAGES CXX CUDA) #这个写法也是可以的
project(demo)
enable_language(CUDA)#激活CUDA语言支持，使用第一个写法时要进行注释
#设置语言标准
set(CMAKE_CXX_STANDARD 11)
set(CMAKE_CUDA_STANDARD 11)   

#添加项目自身的库和依赖
set(EXECUTABLE_OUTPUT_PATH ${PROJECT_SOURCE_DIR}/bin)#设置编译后的程序存放位置
aux_source_directory(src SRC_LIST)#添加src目录下的所有源文件
include_directories(include)#添加头文件路径

#添加OpenCV库（这个库比较常用，所以写个注释在这里）
#find_package(OpenCV 4 REQUIRED)
#include_directories(${OpenCV_INCLUDE_DIRS})
#message(STATUS "include:${OpenCV_INCLUDE_DIRS}")
#link_libraries(${OpenCV_LIBS})

#编译列表
add_executable(${PROJECT_NAME} ${SRC_LIST})

foo.cuh

该文件存放数组的相关

#ifndef FOO_CUH
#define FOO_CUH

#include <stdio.h>
//非模板函数才可以进行C连接
extern "C" float *matAdd(float *a,float *b,int length);

#endif

CUDA算法讲解

在不进行CUDA加速的时候，我们一般采用循环的方式进行相关的计算，如下

float *matAdd(float *a,float *b,int length)
{
    int size = length * sizeof(float);//计算空间大小
    float *sum =(float *)malloc(size) ;//开辟动态内存空间
    return sum;
	for (int i=0;i<length;i++){
		sum[i] = a[i]+b[i];
	}
	return sum;
}

最外层有一个循环，时间复杂度为O(n)，
为了能够缩减其效率，我们对其采用并行运算，将其最外层循环去掉，降低时间复杂度，因此我们可以将矩阵的每个数字存放在显卡线程上进行运算，因此我们可以定义线程上执行的核函数:

template<typename T>
__global__ void matAdd_cuda(T *a,T *b,T *sum)
{
	//blockIdx代表block的索引,blockDim代表block的大小，threadIdx代表thread线程的索引，因此对于一维的block和thread索引的计算方式如下
    int i = blockIdx.x*blockDim.x+ threadIdx.x;
    sum[i] = a[i] + b[i];
}

这样我们可以在显卡上执行并行运算

matAdd_cuda<float><<<block,thread>>>(aGPU,bGPU,sumGPU);

得到相应的输出
执行CUDA程序的步骤如下：
1.开辟显存空间
2.数据从内存拷贝到显存
3.执行核函数
4.数据从显存拷贝到内存
5.释放显存

相关详细内容如下

foo.cu

#include "foo.cuh"
//添加cuda库
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
//定义核函数 __global__为声明关键字
template<typename T>
__global__ void matAdd_cuda(T *a,T *b,T *sum)
{
	//blockIdx代表block的索引,blockDim代表block的大小，threadIdx代表thread线程的索引，因此对于一维的block和thread索引的计算方式如下
    int i = blockIdx.x*blockDim.x+ threadIdx.x;
    sum[i] = a[i] + b[i];
}
//核函数用模板不会报错，模板名字是具有链接的，但它们不能具有C链接，因此不能用在供调用的函数上
float *matAdd(float *a,float *b,int length)
{
    int device = 0;//设置使用第0块GPU进行运算
    cudaSetDevice(device);//设置运算显卡
    cudaDeviceProp devProp;
    cudaGetDeviceProperties(&devProp, device);//获取对应设备属性
    int threadMaxSize = devProp.maxThreadsPerBlock;//每个线程块的最大线程数
    int blockSize = (length+threadMaxSize-1)/threadMaxSize;//计算Block大小,block一维度是最大的，一般不会溢出
    dim3 thread(threadMaxSize);//设置thread
    dim3 block(blockSize);//设置block
    int size = length * sizeof(float);//计算空间大小
    float *sum =(float *)malloc(size) ;//开辟动态内存空间
    //开辟显存空间
    float *sumGPU,*aGPU,*bGPU;
    cudaMalloc((void**)&sumGPU,size);
    cudaMalloc((void**)&aGPU,size);
    cudaMalloc((void**)&bGPU,size);
    //内存->显存
    cudaMemcpy((void*)aGPU,(void*)a,size,cudaMemcpyHostToDevice);
    cudaMemcpy((void*)bGPU,(void*)b,size,cudaMemcpyHostToDevice);
    //运算
    matAdd_cuda<float><<<block,thread>>>(aGPU,bGPU,sumGPU);
    //cudaThreadSynchronize();
    //显存->内存
    cudaMemcpy(sum,sumGPU,size,cudaMemcpyDeviceToHost);
    //释放显存
    cudaFree(sumGPU);
    cudaFree(aGPU);
    cudaFree(bGPU);
    return sum;

}

main.cpp

用来测试程序

#include <stdio.h>
#include <iostream>
#include "foo.cuh"
//#include "opencv2/opencv.hpp"
int main()
{

    //cv::Mat image = cv::imread("./images/P0000.png");
    //unsigned char *array = image.data;
    //创建数组
    const int length = 10;
    float a[length],b[length];
    float *c = (float *)malloc(length *sizeof(float));
    for(int i=0;i<length;i++)
    {
        a[i] = 1;
        b[i] = 2;
    }
    //矩阵加法运算
    c = matAdd(a,b,length);
    //输出查看是否完成计算
    for(int i=0;i<length;i++)
    {
        std::cout<<a[i]<<" "<<b[i]<<" "<<c[i]<<std::endl;
    }
    return 0;
}

运行

进入build

cmake ..
make //windows平台可用MinGW编译器运行mingw32-make.exe
../bin/demo

参考资料

1.CUDA编程入门极简教程
2.CUDA编译（二）—用CMake混合编译C++与cuda
3.Linux下CMake简明教程