前言
上一节环境配置好了，我们开始吧！
一、A First Program
1. Hello, World!
我们先写一个C语言的 Hello, World! 作为对比
int main(void){printf("Hello, World!\n");return 0;
}
大家应该知道这个代码运行在CPU上吧，我们CPU和系统的内存称作 host，GPU及其内存称作 device。
#include <stdio.h>__global__ void kernel(void){}int main(void){kernel<<<1,1>>>();printf("Hello,World!\n");return 0;
}
这个程序：
① 有一个叫kernal且被 global 修饰的函数
② 调用这个kernal函数用了 <<<1,1>>>
C语言部分的代码就是用vs编译的，CUDA C 这部分，也就是带有 global 的代码，将在GPU上编译（在GPU上执行的函数称为CUDA核函数（Kernel Function)）。
（NVIDIA工具将C语言交给C编译器，核函数交给device：函数kernel()将被交给编译device代码的编译器，而main()函数将被交给host编译器）
2. Passing Parameters
#include "stdio.h"__global__ void add(int a, int b, int *c){*c = a+b;
}int main(){int c;int *dev_c;
cudaMalloc((void**)&dev_c, sizeof(int));add<<<1,1>>>(2,7,dev_c);cudaMemcpy(&c,dev_c,sizeof(int),cudaMemcpyDeviceToHost);printf("2+7 = %d\n",c);cudaFree(dev_c);return 0;
}
① 我们可以像C代码那样传递参数调用 global 函数
② 要先在device上申请内存，然后操作后返回给host
我们说一下 cudaMalloc：
和C的malloc很相似，不同的是，这个函数是要在 device – GPU 上申请内存。
第一个参数：
是一个指向指针的指针，因为我们需要新开辟内存的地址
第二个参数：
是内存大小
除了传参不同，这个函数返回值是CUDA中定义的一个错误代码。
CUDA C淡化了主机代码和设备代码之间的差异，这样的话，我们思考一个问题，我们可以在host端使用(解引用) dev_c 吗？
大家思考3秒钟…
当然不能，这是 CUDA 返回过来的，这可是GPU的内存哦。
你使用的话，编译器可是识别不出来的，要小心！
所以这里有几条限制：
① 可以将cudaMalloc()分配的指针传递给在device上执行的函数
② 可以在device代码中使用cudaMalloc()分配的指针进行内存读写操作
③ 可以将cudaMalloc()分配的指针传递给在host上执行的函数
④ 不能在主机代码中使用cudaMalloc()分配的指针进行内存读写操作
这③和④大家觉得是不是有歧义哦，是这样的：③的意思是在host可以传递参数 – 赋值操作是没问题的；④的意思是在host上解引用操作是不行的。大家可以理解吧。同样的道理，host的指针在device中也会出错哦。
cudaFree() 和 cudaMalloc() 相对应，一个释放，一个申请。
下面说一下 cudaMemcpy(dst, src, count, kind) :
- 前2个为指针，指向device 和 host 内存
- 第三个为拷贝的数据大小
- 第四个为拷贝的方向
- 返回值是错误代码
用kind标注这个复制的方向，这次程序里的是 cudaMemcpyDeviceToHost，也就是说源地址是device，目的地址是host。
反方向的是cudaMemcpyHostToDevice，还有一个很奇怪的cudaMemcpyDeviceToDevice，不知道这和 memcpy 有什么区别…
这些使用方法大家可以再去英伟达官网上瞅瞅：
https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__MEMORY.html#group__CUDART__MEMORY_1ga042655cbbf3408f01061652a075e094
2. Querying Devices
我们可以通过代码接口查看设备的情况：
#include "stdio.h"int main() {int count;cudaGetDeviceCount(&count);cudaDeviceProp prop;for (int i = 0; i < count; i++) {cudaGetDeviceProperties(&prop, i);printf(" --- General Information for device %d ---\n", i);printf("Name: %s\n", prop.name);printf("Compute capability: %d.%d\n", prop.major, prop.minor);printf("Clock rate: %d\n", prop.clockRate);printf("Device copy overlap: ");if (prop.deviceOverlap)printf("Enabled\n");elseprintf("Disabled\n");printf("Kernel execution timeout : ");if (prop.kernelExecTimeoutEnabled)printf("Enabled\n");elseprintf("Disabled\n");printf(" --- Memory Information for device %d ---\n", i);printf("Total global mem: %ld\n", prop.totalGlobalMem);printf("Total constant Mem: %ld\n", prop.totalConstMem);printf("Max mem pitch: %ld\n", prop.memPitch);printf("Texture Alignment: %ld\n", prop.textureAlignment);printf(" --- MP Information for device %d ---\n", i);printf("Multiprocessor count: %d\n",prop.multiProcessorCount);printf("Shared mem per mp: %ld\n", prop.sharedMemPerBlock);printf("Registers per mp: %d\n", prop.regsPerBlock);printf("Threads in warp: %d\n", prop.warpSize);printf("Max threads per block: %d\n",prop.maxThreadsPerBlock);printf("Max thread dimensions: (%d, %d, %d)\n",prop.maxThreadsDim[0], prop.maxThreadsDim[1],prop.maxThreadsDim[2]);printf("Max grid dimensions: (%d, %d, %d)\n",prop.maxGridSize[0], prop.maxGridSize[1],prop.maxGridSize[2]);printf("\n");}
}
大家见笑了，我这电脑也就是凑活能用😅
大家去这里搜这俩函数，这个 cudaDeviceProp 内容实在太多，就写了这些，大家去这里看详情：
https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__DEVICE.html#group__CUDART__DEVICE
3.Using Device Properties
假设我们需要找特定能力的GPU设备，要怎么办呢？像上面一样，把每个都打印出来一条一条挨着看吗？
int main( void ) {cudaDeviceProp prop;int dev;cudaGetDevice( &dev );printf( "ID of current CUDA device: %d\n", dev );memset( &prop, 0, sizeof( cudaDeviceProp ) );prop.major = 1;prop.minor = 3;cudaChooseDevice( &dev, &prop );printf( "ID of CUDA device closest to revision 1.3: %d\n", dev );cudaSetDevice( dev );
}
我们可以用这个 cudaDeviceProp 设定出我们的需求设备，然后 cudaChooseDevice 去找找看有没有（函数返回码），有的话就去设定。
总结
我们已经在GPU上运行出来了经典的 Hello world ！，大家已经入门了，继续加油哦~