0. Introduction

CUDA是什么

英伟达为自家显卡开发的工具，方便程序员进行与显卡相关的编程。平时我们的程序如果没有使用类似CUDA的工具，都是在CPU上运行的，想让程序在GPU上运行就需要专门编程来实现。

准备工作

搭建CUDA编程环境。在Windows环境中，安装CUDA后可以用VS创建CUDA项目，直接用nvcc编译器进行编译会有无法使用cl.exe链接等问题。
本系列文章使用Ubuntu 22.04 LTS。

按照正常流程安装配置完，就ok了。

# include <stdio.h>
__global__ void hello_from_GPU()
{
    printf("hello from GPU \n");
}
int main(void){
    hello_from_GPU<<<4, 4>>>();
    cudaDeviceSynchronize();
    return 0;
}

copy一段简单代码，使用nvcc进行编译即可。这里的<<<X, Y>>>稍后再解释(这里不是正常的C/C++语法)。
__global__用于定义核函数。同样，__global__也在是nvcc中定义的。
cudaDeviceSynchronize()类似多线程编程中的thread_join，或者多进程编程中的waitpid()，用于等待GPU上任务完成实现同步。

函数类型

__host__，__device__和__global__是三种特殊关键字，用于定义不同类型的函数。
在CUDA编程中，主机(host)指CPU端；设备(device)指GPU端。

__host__

定义一个在CPU端的函数。没有这三种关键字时默认是__host__类型的函数。

__device__

定义一个在GPU端的函数，在GPU端执行和调用。

__global__

定义核函数，在CPU端使用func<<<X, Y>>>()的形式进行调用。不能有返回值，可以有参数，host可以调用global，global可以调用device，device可以调用device。
另外__host__可以和__device__结合使用，此时既可以在CPU端被调用，也可以在GPU端被调用。此时调用情况比较复杂。以下代码虽然能运行，但是会出现很多警告。

#include <iostream>
using namespace std;
__host__ __device__ void func1() {
    // cout<<"1\n";
    string s = "abc";
    printf("%c\n", s[0]);
}
__host__ void func2(){
    cout<<"2\n";
}
__device__ void func3(){
    // cout<<"3\n";
}
__global__ void kernal() {
    func3();
}

int main(){
    func1();
    func2();
    kernal<<<1, 1>>>();
    cudaDeviceSynchronize();
    return 0;
}

警告的主要内容是warning #20011-D: calling a __host__ function("std::allocator<char> ::~allocator()") from a __host__ __device__ function("func1") is not allowed，也就是说string类属于host，在__host__ __device__标识的函数中调用不合法。
注意如果func1()和func3()内使用cout，就会报错，而非只是警告。
那么__host__ __device__标识的函数和核函数的调用情况呢？

#include <iostream>
using namespace std;
__host__ __device__ void func1() {
    printf("func1() called.\n");
}
__global__ void kernal() {
    func1();
}
__host__ __device__ void func2() {
    kernal<<<1, 1>>>();
}
int main(){
    func2();
    cudaDeviceSynchronize();
    return 0;
}

此时编译参数要加上-rdc=true，然后可以正常输出func1() called.。说明__host__ __device__标识的函数可以和核函数互相调用。

参考资料

教程汇总帖：https://zhuanlan.zhihu.com/p/346910129
安装CUDA：https://blog.csdn.net/qq_74083422/article/details/134745040