前言
上一节环境配置好了，我们开始吧！
一、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");
else
printf("Disabled\n");
printf("Kernel execution timeout : ");
if (prop.kernelExecTimeoutEnabled)
printf("Enabled\n");
else
printf("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 ！，大家已经入门了，继续加油哦~
