前言
上一节环境配置好了,我们开始吧!
一、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 !,大家已经入门了,继续加油哦~