用GPU输出Hello World修饰符__global__告诉编译器这个函数将会从CPU调用，然后在GPU上执行。使用

// hello.cu

#include <stdio.h>

__global__ void helloFormGPU_kernel(void){
    int idx = blockIdx.x * blockDim.x + threadIdx.x; 
    printf("Hello World from GPU!, %d\n", idx);
}
int main(void){
    printf("Hello World from CPU!\n");

    helloFormGPU_kernel <<<1, 10>>>();
    cudaDeviceReset();
    return 0;
}

修饰符__global__告诉编译器这个函数将会从CPU调用，然后在GPU上执行。使用下面的代码启动内核函数。

helloFormGPU_kernel <<<1, 10>>>();

三重尖括号意味着从主线程到设备端代码的调用。一个内核函数通过一组线程来执行，所有线程执行相同的代码。三重尖括号里面的参数是执行配置，用来说明使用多少线程来执行内核函数，本例中有10个GPU线程被调用。

nvcc编译：

nvcc hello.cu -o hello

运行可执行文件：

./hello

输出：

Hello World from CPU!
Hello World from GPU!, 0
Hello World from GPU!, 1
Hello World from GPU!, 2
Hello World from GPU!, 3
Hello World from GPU!, 4
Hello World from GPU!, 5
Hello World from GPU!, 6
Hello World from GPU!, 7
Hello World from GPU!, 8
Hello World from GPU!, 9

使用-arch sm_80进行编译:

nvcc -arch sm_80 hello.cu -o hello

使用cuda-gdb调试，编译时需加-g -G，-g 表示将CPU代码(host)编译成可调式版本，-G表示将GPU代码(kernel)编译成可调式版本：

nvcc -g -G -arch sm_80 hello.cu -o hello

cuda-gdb调试：

cuda-gdb
set cuda software_preemption on
//运行可执行程序
file hello

常用参数： | set args arg1 arg2 | | 设置命令行参数 | set args1 1 2 | | ------------------ | --- | --- | --- | | breakpoint | b | 设置断掉，其参数可以是方法名或者是行号 | b main；b数字 | | run | r | 在调试器中执行程序 | | | start | | 开始执行程序，并在main的第一行停住 | | | next | n | 单步执行到下一行 | | | continue | c | 继续执行已暂停的程序至下一个断点或程序结尾 | | | step | s | 单步执行，会进入函数内部执行 | | | print | p | 打印参数信息，查看变量 | p var1 | | thread | | 列出当前主机线程 | | | cuda | | 列出当前活跃的kernel/grid/block/thread内容，并允许将焦点移至此处| cuda thread(1,1,1);cuda kernel 1 block(1,2,1)| | cuda | | 列出当前活跃的gpu线程（若有的话） | | | cuda kernel | | 列出当前活跃的GPU kernel，并允许将焦点移到指定的GPU线程| | | info | | 查看参数所包含的具体信息 | info devices;info kernel;info threads| | backtrace | bt | 显示当前函数调用栈的内容 | |
NVIDIA Profiling

nsys profile -t nvtx,cuda --stats=true --force-overwrite true --wait=all -o my_report ./hello

nsys stats my_report.sqlite

cuda-memcheck工具：

cuda$ cuda-memcheck  ./hello

nsys profile ./hello