c++和cuda混合编程总结

---

前言

C++ 与 CUDA 混合编程是指在同一个程序中同时使用 C++ 和 CUDA 编程，以实现在 GPU 上并行处理任务。这种混合编程方式允许你充分利用 GPU 的并行计算能力来加速特定的任务，同时还能在同一程序中使用 C++ 进行控制逻辑和数据处理。

编译命令

gcc/g++

命令	解释
-w	关闭所有警告信息
-g	编译生成的可执行文件中包含调试信息。这些信息可以被调试器使用
-m64	将源代码编译成64位机器指令，以便在64位操作系统上运行。否则默认生成32位代码。
-fPIC	表示生成位置独立代码(Position Independent Code)，也就是生成可重定位目标文件。
-rpath	指定运行时库搜索路径
-Wl	表示将后面的选项传递给连接器ld(例如 -Wl,-rpath=/usr/local/lib)

---

nvcc

命令	解释
-Xcompiler=fPIC	nvcc设置fPIC的方式
-rdc=true	启用CUDA的远程设备代码(RDC)特性，(打包静态库时要设置为false)
-ccbin=g++	指定使用的C++编译器

---

ar

参数	解释
r	替换或添加文件到档案中。
c	创建新的档案。
s	在建立档案时，将目标文本中所有符号保存到符号表中。
v	显示详细的执行过程。
q	将新增文件添加到档案中。
t	列出目前档案中所有的成员文件名。
d	从档案中删除指定的文件。

打包静态库命令ar crv lib.a file1.o file2.o，其中，c表示创建新的档案，r表示替换或添加文件到档案中，v表示显示详细的执行过程。

---

打包库

动态库

命令

# nvcc -Xcompiler=-fPIC -shared -o libkernel.so kernel.cu
nvcc -c -Xcompiler=-fPIC -o kernel.cu.o kernel.cu
nvcc -shared -o libkernel.so kernel.cu.o
g++ -c main.cpp -o main.o -I/usr/local/cuda-11.3/targets/x86_64-linux/include/
g++ main.o -o main -L. -L/usr/local/cuda-11.3/targets/x86_64-linux/lib -lcuda -lcudart -lkernel

cmake

cmake_minimum_required(VERSION 3.10)
# set(CMAKE_VERBOSE_MAKEFILE on)
project(kernel CUDA CXX)
set(CMAKE_LIBRARY_OUTPUT_DIRECTORY ${CMAKE_ROOT_DIR}/lib)
set(CMAKE_ARCHIVE_OUTPUT_DIRECTORY ${CMAKE_ROOT_DIR}/lib)
set(CMAKE_CXX_COMPILER g++)
set(CMAKE_CXX_STANDARD 14)
set(CMAKE_CUDA_STANDARD 14)
set(CMAKE_CUDA_ARCHITECTURES xx)

find_package(CUDA REQUIRED)
include_directories(/usr/local/cuda-11.3/targets/x86_64-linux/include)
add_library(${PROJECT_NAME} SHARED vector_add.cu)

set_target_properties(${PROJECT_NAME} PROPERTIES
       CUDA_SEPARABLE_COMPILATION ON)
target_link_libraries(${PROJECT_NAME} cuda cudart stdc++)

cmake_minimum_required(VERSION 3.10)
project(main CXX)
set(CMAKE_CXX_STANDARD 14)

include_directories(xxx)
link_directories(xxx)
include_directories(/usr/local/cuda-11.3/targets/x86_64-linux/include)
link_directories(/usr/local/cuda-11.3/targets/x86_64-linux/lib)

add_executable(${PROJECT_NAME} main.cpp)
target_link_libraries(${PROJECT_NAME} cuda cudart stdc++ m kernel)

静态库

命令

nvcc -c kernel.cu -o kernel.cu.o
ar rc libvectoradd.a kernel.cu.o
g++ -c main.cpp -o main.o -I/usr/local/cuda-11.3/targets/x86_64-linux/include/
g++ main.o -o main -L. -L/usr/local/cuda-11.3/targets/x86_64-linux/lib -lcuda -lcudart -lkernel

cmake

cmake_minimum_required(VERSION 3.10)
# set(CMAKE_VERBOSE_MAKEFILE on)
project(kernel CUDA CXX)
set(CMAKE_LIBRARY_OUTPUT_DIRECTORY ${CMAKE_ROOT_DIR}/lib)
set(CMAKE_ARCHIVE_OUTPUT_DIRECTORY ${CMAKE_ROOT_DIR}/lib)
set(CMAKE_CXX_COMPILER g++)
set(CMAKE_CXX_STANDARD 14)
set(CMAKE_CUDA_STANDARD 14)
set(CMAKE_CUDA_ARCHITECTURES xx)

find_package(CUDA REQUIRED)
include_directories(/usr/local/cuda-11.3/targets/x86_64-linux/include)
add_library(${PROJECT_NAME} STATIC vector_add.cu)
# 下面要注释掉
# set_target_properties(${PROJECT_NAME} PROPERTIES
#       CUDA_SEPARABLE_COMPILATION ON)
target_link_libraries(${PROJECT_NAME} cuda cudart stdc++)

cmake_minimum_required(VERSION 3.10)
project(main CXX)
set(CMAKE_CXX_STANDARD 14)

include_directories(xxx)
link_directories(xxx)

include_directories(/usr/local/cuda-11.3/targets/x86_64-linux/include)
link_directories(/usr/local/cuda-11.3/targets/x86_64-linux/lib)
add_executable(${PROJECT_NAME} main.cpp)
target_link_libraries(${PROJECT_NAME} cuda cudart stdc++ m libkernel.a)

案例

一个简单的demo

c++代码

#include <cuda_runtime.h>
#include <stdio.h>
#define checkRuntime(op)  __check_cuda_runtime((op), #op, __FILE__, __LINE__)
bool __check_cuda_runtime(cudaError_t code, const char* op, const char* file, int line){
    if(code != cudaSuccess){    
        const char* err_name = cudaGetErrorName(code);    
        const char* err_message = cudaGetErrorString(code);  
        printf("runtime error %s:%d  %s failed. \n  code = %s, message = %s\n", file, line, op, err_name, err_message);   
        return false;
    }
    return true;
}
void test_print(const float* pdata, int ndata);
int main(){
    float* parray_host = nullptr;
    float* parray_device = nullptr;
    int narray = 10;
    int array_bytes = sizeof(float) * narray;

    parray_host = new float[narray];
    checkRuntime(cudaMalloc(&parray_device, array_bytes));

    for(int i = 0; i < narray; ++i)
        parray_host[i] = i;
    
    checkRuntime(cudaMemcpy(parray_device, parray_host, array_bytes, cudaMemcpyHostToDevice));
    test_print(parray_device, narray);
    checkRuntime(cudaDeviceSynchronize());

    checkRuntime(cudaFree(parray_device));
    delete[] parray_host;
    return 0;
}

cu代码

#include <stdio.h>
#include <cuda_runtime.h>

__global__ void test_print_kernel(const float* pdata, int ndata){

    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    /*    dims                 indexs
        gridDim.z            blockIdx.z
        gridDim.y            blockIdx.y
        gridDim.x            blockIdx.x
        blockDim.z           threadIdx.z
        blockDim.y           threadIdx.y
        blockDim.x           threadIdx.x

        Pseudo code:
        position = 0
        for i in 6:
            position *= dims[i]
            position += indexs[i]
    */
    printf("Element[%d] = %f, threadIdx.x=%d, blockIdx.x=%d, blockDim.x=%d\n", idx, pdata[idx], threadIdx.x, blockIdx.x, blockDim.x);
}

void test_print(const float* pdata, int ndata){

    // <<<gridDim, blockDim, bytes_of_shared_memory, stream>>>
    test_print_kernel<<<1, ndata, 0, nullptr>>>(pdata, ndata);

    // 在核函数执行结束后，通过cudaPeekAtLastError获取得到的代码，来知道是否出现错误
    // cudaPeekAtLastError和cudaGetLastError都可以获取得到错误代码
    // cudaGetLastError是获取错误代码并清除掉，也就是再一次执行cudaGetLastError获取的会是success
    // 而cudaPeekAtLastError是获取当前错误，但是再一次执行 cudaPeekAtLastError 或者 cudaGetLastError 拿到的还是那个错
    // cuda的错误会传递，如果这里出错了，不移除。那么后续的任意api的返回值都会是这个错误，都会失败
    cudaError_t code = cudaPeekAtLastError();
    if(code != cudaSuccess){    
        const char* err_name    = cudaGetErrorName(code);    
        const char* err_message = cudaGetErrorString(code);  
        printf("kernel error %s:%d  test_print_kernel failed. \n  code = %s, message = %s\n", __FILE__, __LINE__, err_name, err_message);   
    }
}

编译命令

# 编译生成kernel.o文件
nvcc -c kernel.cu
# 编译生成main.o文件
g++ -c main.cpp  -I/usr/local/cuda-11.3/targets/x86_64-linux/include/
# 链接
g++ kernel.o main.o -L/usr/local/cuda-11.3/lib64 -lcudart -lstdc++