在《nVidia CUDA API(上)》的部分,已经大致说明了 extension 的部分;这边要来讲的则是 runtime library 的部分。
CUDA 的 runtime library 分成下面三部分:
-
common component
提供 CUDA 的 vector 等型别,以及可以同时在 host 及 device 上可以执行的函式。主要包括:-
内建的 vector 型别
-
dim3 型别(用於指定 kernel 参数的型别)
-
texture 型别
-
数学函式
-
-
device component
在 device 上执行的部分,提供 device 上特殊的函式。主要有:-
数学函式。这边的数学函式是 device 的特殊版本,精确度较低,但是速度较快。
-
同步函式
-
Atomic 函式
-
型别转换函式(Conversion/Casting)
-
texture 函式
-
-
host component
在 host 上执行的部分,负责控制 device。主要是处理:-
Device 管理
-
Context 管理(所谓 Context 大致相当於 CPU?)
-
记忆体管理
-
Code module 管理(Code module 似乎相当於 dynamic library)
-
Execution control
-
Texture reference 管理
-
Interoperability with OpenGL and Direct3D.
-
不过,这边应该就不会列举了~个别的说明,请参考《CUDA Programming Guide 1.0》这份文件。而这边,就只大概介绍一些 Heresy 觉得比较会用到的部分了。
记忆体管理
要说最重要、一定会用到的部分,应该就是记忆体管理的部分了!因为在 CUDA 中,要在 device 的程式中存取的资料,一定要先存在 device 上,所以在记忆体管理的部分, CUDA 提供了很类似 C 的一些函式: cudaMalloc()、 cudaFree()。原则上, cudaMalloc 大致上就是 C 语言中的 malloc,也就像是 C++ 的 new;而 cudaFree 则就相当於 C 的 free,以及 C++ 的 delete。
而要使用的方法,则是要先宣告 pointer,在 allocate 记忆体给他;下面就是一个简单的范例:
float *dev_Array;
cudaMalloc( (void**)&dev_Array, ArraySize * sizeof( float ) );
这样,就可以在 device 上配置一块大小是 ArraySize 的 float 阵列了~
如果要把已经存在一般程式中的资料复制到 device 上呢?这时候就要用对应到 C 语言中 memcpy 的函式- cudaMemcpy 了~以 CUDA SDK 范例里的写法,会是:
// define, allocate on host
float *host_Array = new float[100];
for( int i = 0; i < 100; ++ i )
host_Array = i;
// define, allocate on device
float *dev_Array;
cudaMalloc( (void**)&dev_Array, 100 * sizeof( float ) );
// copy from host to device
cudaMemcpy>( dev_Array, host_Array, 100 * sizeof( float ),
cudaMemcpyHostToDevice );
cudaMemcpy 这个函式的四个参数依序为:目标、来源、大小、方法。最后一项方法的型别是 CUDA 的一个列举型别 cudaMemcpyKind,他有四种值:cudaMemcpyHostToHost, cudaMemcpyHostToDevice, cudaMemcpyDeviceToHost, cudaMemcpyDeviceToDevice;相当直觉的,就是代表由 host/device 复制到 device/host。所以,如果是要反过来把 device 上的资料复制回 host 的话,就只要把最后一项参数改成 cudaMemcpyDeviceToHost 就可以了!
而要释放掉 allocate 出来的记忆体空间的话,就直接使用 cudaFree 就可以了。
cudaFree( dev_Array );
实际上 CUDA 的记忆体管理除了这边提到的基本类型,还有 cudaMallocPitch(), cudaMemset(), cudaMemcpy2D(), cudaMemcpyToArray(), cudaMemcpyToSymbol() 等其他的记忆体管理函式,在这边就不一一提出来了。
装置管理
一般情况下,可能用不到这一部分的 API 吧~但是如果有两个以上的 CUDA Device 的话,就须要透过这些函式来选择要用的装置了~此外,也可以先透过这些函式,来确认机器上的一些基本规格。
这部分的函式不多,有下面几个:
-
cudaGetDeviceCount( int* count )
取得目前机器上,支援 CUDA 的 device 数目,会把结果存在 count 中。函式则会 return 一个 cudaError_t,用来表示 CUDA 的错误。
-
cudaGetDeviceProperties( cudaDeviceProp* prop, int dev )
取得目前机器上第 dev 个 device 的属性;prop 的结构会是:
struct cudaDeviceProp { char name[256]; //装置名称 size_t totalGlobalMem; //Device 的 global memory 总量 size_t sharedMemPerBlock; //Device 的 shared memory 总量 int regsPerBlock; //每一个 block 的 registers 数量 int warpSize; //warp size size_t memPitch; //最大的 memory pitch int maxThreadsPerBlock; //Block 的最大 thread 数 int maxThreadsDim[3]; //Block 中 thread 最大的维度 int maxGridSize[3]; //Grid 中 block 的最大维度 size_t totalConstMem; //Device 的 constant memory 总量 int major; //主要版本编号 int minor; //次要版本编号 int clockRate; //时脉 size_t textureAlignment; //alignment requirement mentioned };
-
cudaChooseDevice( int* dev, const cudaDeviceProp* prop)
在目前的 device 中,找一个最符合给定的 prop 的,并把他的索引直存在 dev 里。
-
cudaSetDevice(int dev)
指定现在要使用第 dev 个 device。
-
cudaGetDevice(int* dev)
取得目前正在使用的 device 的索引编号,储存在 dev 中。
快速数学计算
从上面的列表可以发现,在 common component 和 device component 中,都有数学函式的功能。其中,common component 中的应该算是一般的数学计算函式,而 device component 所提供的,则是精确度较低,但是速度较快的版本!
也就是说 CUDA 在 device component 的部分,提供了部分数学计算的函式,是以 GPU 来做计算的快速版本。这些函式在 CUDA 中统一以「__」来做开头,例如 __sinf(x)、__fadd_rz(x,y)、__logf(x) 等等。这些函式在《CUDA Programming Guide 1.0》的 Appendix B 中,可以找到比较完整的列表可以参考。
而在 nvcc 的编译参数中,也还提供了一个「-use_fast_math」的参数,可以强制所有可以的数学计算函数,都使用叫快速的版本来计算。这在精确度比较不重要的情况下,应该可以对效能有些增益。
上面讲的部分,都是 runtime API 的部分;大部分的函式,都是以 cuda 来做为开头。而实际上,CUDA 还有提供 Driver API;这一部分的函式会是以 cu 来做为开头。而许多函式,都可以和 runtime API 来做一一对应。而以文件里的说法,Driver API 应该是比 runtime API 来的低阶,但是实际的差异,Heresy 也不是很清楚就是了。
本文主要参考nVidia 的《CUDA Programming Guide 1.0》(PDF 文件)的第四章。