关于Ubuntu12.04 下 CUDA5.5 的安装请参看如下链接
1、使用 Runtime API 创建第一个 CUDA 程序
CUDA 初始化函数
由于是使用 Runtime API, 所以在文件开头要加入 cuda_runtime.h 头文件。
初始化函数包括一下几个步骤：
- 获取 CUDA 设备数
- 获取 CUDA 设备属性
- 设置 CUDA 设备
/* *******************************************************************
##### File Name: first_cuda.cu
##### File Func: initial CUDA device and print device prop
##### Author: Cai*
##### E-mail: cai*220@gmail.com
##### Create Time: 2014-4-21
* ********************************************************************/
#include <stdio.h>
#include <cuda_runtime.h>
void printDeviceProp(const cudaDeviceProp &prop)
{
printf("Device Name : %s.\n", prop.name);
printf("totalGlobalMem : %d.\n", prop.totalGlobalMem);
printf("sharedMemPerBlock : %d.\n", prop.sharedMemPerBlock);
printf("regsPerBlock : %d.\n", prop.regsPerBlock);
printf("warpSize : %d.\n", prop.warpSize);
printf("memPitch : %d.\n", prop.memPitch);
printf("maxThreadsPerBlock : %d.\n", prop.maxThreadsPerBlock);
printf("maxThreadsDim[0 - 2] : %d %d %d.\n", prop.maxThreadsDim[0], prop.maxThreadsDim[1], prop.maxThreadsDim[2]);
printf("maxGridSize[0 - 2] : %d %d %d.\n", prop.maxGridSize[0], prop.maxGridSize[1], prop.maxGridSize[2]);
printf("totalConstMem : %d.\n", prop.totalConstMem);
printf("major.minor : %d.%d.\n", prop.major, prop.minor);
printf("clockRate : %d.\n", prop.clockRate);
printf("textureAlignment : %d.\n", prop.textureAlignment);
printf("deviceOverlap : %d.\n", prop.deviceOverlap);
printf("multiProcessorCount : %d.\n", prop.multiProcessorCount);
}
bool InitCUDA()
{
//used to count the device numbers
int count;
// get the cuda device count
cudaGetDeviceCount(&count);
if (count == 0) {
fprintf(stderr, "There is no device.\n");
return false;
}
// find the device >= 1.X
int i;
for (i = 0; i < count; ++i) {
cudaDeviceProp prop;
if (cudaGetDeviceProperties(&prop, i) == cudaSuccess) {
if (prop.major >= 1) {
printDeviceProp(prop);
break;
}
}
}
// if can‘t find the device
if (i == count) {
fprintf(stderr, "There is no device supporting CUDA 1.x.\n");
return false;
}
// set cuda device
cudaSetDevice(i);
return true;
}
int main(int argc, char const *argv[])
{
if (InitCUDA()) {
printf("CUDA initialized.\n");
}
return 0;
}
Runtime API 函数解析
cudaGetDeviceCount —— 返回具有计算能力的设备的数量函数原型： cudaError_t cudaGetDeviceCount( int* count )
函数说明： 以*count形式返回可用于执行的计算能力大于等于1.0的设备数量。如果不存在此类设备，将返回1
返回值： cudaSuccess，注意，如果之前是异步启动，该函数可能返回错误码。
cudaGetDeviceProperties —— 返回关于计算设备的信息
函数原型： cudaError_t cudaGetDeviceProperties( struct cudaDeviceProp* prop,int dev )
函数说明： 以*prop形式返回设备dev的属性。
返回值： cudaSuccess、cudaErrorInvalidDevice，注意，如果之前是异步启动，该函数可能返回错误码。
cudaDeviceProp 结构定义如下:
struct cudaDeviceProp {
char name [256];
size_t totalGlobalMem;
size_t sharedMemPerBlock;
int regsPerBlock;
int warpSize;
size_t memPitch;
int maxThreadsPerBlock;
int maxThreadsDim [3];
int maxGridSize [3];
size_t totalConstMem;
int major;
int minor;
int clockRate;
size_t textureAlignment;
int deviceOverlap;
int multiProcessorCount;
}其中:name
用于标识设备的ASCII字符串;
totalGlobalMem
设备上可用的全局存储器的总量,以字节为单位;
sharedMemPerBlock
线程块可以使用的共享存储器的最大值,以字节为单位;多处理器上的所有线程块可以同时共享这些存储器;
regsPerBlock
线程块可以使用的32位寄存器的最大值;多处理器上的所有线程块可以同时共享这些寄存器;
warpSize
按线程计算的warp块大小;
memPitch
允许通过cudaMallocPitch()为包含存储器区域的存储器复制函数分配的最大间距(pitch),以字节为单位;
maxThreadsPerBlock
每个块中的最大线程数
maxThreadsDim[3]
块各个维度的最大值:
maxGridSize[3]
网格各个维度的最大值;
totalConstMem
设备上可用的不变存储器总量,以字节为单位;
major,minor
定义设备计算能力的主要修订号和次要修订号;
clockRate
以千赫为单位的时钟频率;
textureAlignment
对齐要求;与textureAlignment字节对齐的纹理基址无需对纹理取样应用偏移;
deviceOverlap
如果设备可在主机和设备之间并发复制存储器,同时又能执行内核,则此值为 1;否则此值为 0;
multiProcessorCount
设备上多处理器的数量。
cudaSetDevice —— 设置设备以供GPU执行使用
函数原型： cudaError_t cudaSetDevice(int dev)
函数说明： 将dev记录为活动主线程将执行设备码的设备。
返回值： cudaSuccess、cudaErrorInvalidDevice，注意，如果之前是异步启动，该函数可能返回错误码。
nvcc 编译代码nvcc 是 CUDA 的编译工具，它可以将 .cu 文件解析出在 GPU 和 host 上执行的部分，也就是说，它会帮忙把 GPU 上执行和主机上执行的代码区分开来，不许要我们手动去做了。在 GPU 执行的部分会通过 NVIDIA 提供的 编译器编译成中介码，主机执行的部分则调用 gcc 编译。
nvcc -o first_cuda first_cuda.cu这样就可以生成可执行文件 first_cuda 了。
运行结果如下所示：
2、利用 CUDA 实现 GPU 计算
例子：计算一个数组中数字的平方和。
首先、修改程序头文件
#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>
其次、在文件头部进行宏定义和全局变量定义
再次、添加一个产生随机数组的函数
DATA_SIZE 表示将要处理 1M 个 int 类型的数据， 也就是 4M Byte 的数据。// ======== define area ========
#define DATA_SIZE 1048576 // 1M
// ======== global area ========
int data[DATA_SIZE];
// ======== used to generate rand datas ========
void generateData(int *data, int size)
{
for (int i = 0; i < size; ++i) {
data[i] = rand() % 10;
}
}
然后、拷贝 HOST 内存到 GPU 内存中，在 main 函数中添加
// generate rand datas generateData(data, DATA_SIZE); // malloc space for datas in GPU int *gpuData, *sum; clock_t *time; cudaMalloc((void**) &gpuData, sizeof(int) * DATA_SIZE); cudaMalloc((void**) &sum, sizeof(int)); cudaMalloc((void**) &time, sizeof(clock_t)); cudaMemcpy(gpuData, data, sizeof(int) * DATA_SIZE, cudaMemcpyHostToDevice);这段程序首先生成一个随机的数组 data ，然后通过 cudaMalloc 函数在 GPU 中创建一块内存空间。利用 cudaMemcpy 函数把 HOST 内存中的数据 data 拷贝到 GPU 内存 gpuData 中，其最后一个参数表示数据拷贝的方向性，另外还有 cudaMemcpyDeviceToHost 表示将 GPU 内存拷贝到 HOST 内存。
同时 time 参数是用来返回 GPU 运行的时钟周期数。
最后、就是写在 GPU 中运行的函数。
在 CUDA 中，通过在函数之前显示的加上 __global__ 表示函数是在 GPU 中执行的。函数如下：
__global__ static void squaresSum(int *data, int *sum, clock_t *time) { int sum_t = 0; clock_t start = clock(); for (int i = 0; i < DATA_SIZE; ++i) { sum_t += data[i] * data[i]; } *sum = sum_t; *time = clock() - start; }注意：上述在 GPU 中执行的函数不能有返回值。接下来让主函数调用这个函数，CUDA 中需要通过一下方式调用。
函数名称<<<block num, thread num, shared memory size>>>(param,...);
在 main 函数需要增加如下代码：
// calculate the squares‘s sum squaresSum<<<1, 1, 0>>>(gpuData, sum, time); // copy the result from GPU to HOST int result; clock_t time_used; cudaMemcpy(&result, sum, sizeof(int), cudaMemcpyDeviceToHost); cudaMemcpy(&time_used, time, sizeof(clock_t), cudaMemcpyDeviceToHost); // free GPU spaces cudaFree(gpuData); cudaFree(sum); cudaFree(time); // print result printf("(GPU) sum:%d time:%ld\n", result, time_used);这个程序只使用一个 thread，一个 block，没有用共享内存。同时提供一个 CPU 程序验证运算结果的正确性。
// CPU calculate result = 0; clock_t start = clock(); for (int i = 0; i < DATA_SIZE; ++i) { result += data[i] * data[i]; } time_used = clock() - start; printf("(CPU) sum:%d time:%ld\n", result, time_used);
一下给出整个文件的代码：
/* *******************************************************************
##### File Name: squareSum.cu
##### File Func: calculate the sum of inputs‘s square
##### Author: Cai*
##### E-mail: cai*220@gmail.com
##### Create Time: 2014-4-21
* ********************************************************************/
#include <stdio.h>
#include <stdlib.h>
#include <cuda_runtime.h>
// ======== define area ========
#define DATA_SIZE 1048576 // 1M
// ======== global area ========
int data[DATA_SIZE];
void printDeviceProp(const cudaDeviceProp &prop);
bool InitCUDA();
void generateData(int *data, int size);
__global__ static void squaresSum(int *data, int *sum, clock_t *time);
int main(int argc, char const *argv[])
{
// init CUDA device
if (!InitCUDA()) {
return 0;
}
printf("CUDA initialized.\n");
// generate rand datas
generateData(data, DATA_SIZE);
// malloc space for datas in GPU
int *gpuData, *sum;
clock_t *time;
cudaMalloc((void**) &gpuData, sizeof(int) * DATA_SIZE);
cudaMalloc((void**) &sum, sizeof(int));
cudaMalloc((void**) &time, sizeof(clock_t));
cudaMemcpy(gpuData, data, sizeof(int) * DATA_SIZE, cudaMemcpyHostToDevice);
// calculate the squares‘s sum
squaresSum<<<1, 1, 0>>>(gpuData, sum, time);
// copy the result from GPU to HOST
int result;
clock_t time_used;
cudaMemcpy(&result, sum, sizeof(int), cudaMemcpyDeviceToHost);
cudaMemcpy(&time_used, time, sizeof(clock_t), cudaMemcpyDeviceToHost);
// free GPU spaces
cudaFree(gpuData);
cudaFree(sum);
cudaFree(time);
// print result
printf("(GPU) sum:%d time:%ld\n", result, time_used);
// CPU calculate
result = 0;
clock_t start = clock();
for (int i = 0; i < DATA_SIZE; ++i) {
result += data[i] * data[i];
}
time_used = clock() - start;
printf("(CPU) sum:%d time:%ld\n", result, time_used);
return 0;
}
__global__ static void squaresSum(int *data, int *sum, clock_t *time)
{
int sum_t = 0;
clock_t start = clock();
for (int i = 0; i < DATA_SIZE; ++i) {
sum_t += data[i] * data[i];
}
*sum = sum_t;
*time = clock() - start;
}
// ======== used to generate rand datas ========
void generateData(int *data, int size)
{
for (int i = 0; i < size; ++i) {
data[i] = rand() % 10;
}
}
其中：
void printDeviceProp(const cudaDeviceProp &prop);
bool InitCUDA();
可以参考 initCUDA.cu 中的实现，完全一致。
nvcc 编译代码
nvcc -o squareSum squareSum.cu运行结果：
CUDA initialized.
(GPU) sum:29909398 time:787124792
(CPU) sum:29909398 time:10000
从执行的结果可以看出， GPU 中运行的程序居然要比 CPU 中的消耗的时钟周期还要多得多。这是有原因的。
因为程序之中并没有使用 CUDA 并行执行的优势。
这里分析一下 GPU 运行的性能。
此 GPU 消耗的时钟周期： 787124792 cycles
GeForce G 103M 的 clockRate： 1.6 GHz
所以可以计算出 GPU 上运行时间是： 时钟周期 / clockRate = 0.49195 s
1 M 个 int 型数据有 4M Byte 的数据量，实际使用的 GPU 内存带宽是：数据量 / 运行时间 = 8.13 MB/s
可见这个程序没有很好的发挥 GPU 的性能，使用的内存带宽很小。
没有有效利用 GPU 性能的原因？？？
在 CUDA 中，一般的数据复制到的显卡内存的部份，称为 global memory。这些内存是没有 cache 的，而且，存取 global memory 所需要的时间（即 latency）是非常长的，通常是数百个 cycles。
由于我们的程序只有一个 thread，所以每次它读取 global memory 的内容，就要等到实际读取到数据、累加到 sum 之后，才能进行下一步。这就是为什么它的表现会这么的差。实际上 GPU 一直在等待上一个数据运行的结束，然后再拷贝一个内存数据，所以使用的时钟周期自然就长了。
由于 global memory 没有 cache，所以要避开巨大的 latency 的方法，就是要利用大量的 threads。假设现在有大量的 threads 在同时执行，那么当一个 thread 读取内存，开始等待结果的时候，GPU 就可以立刻切换到下一个 thread，并读取下一个内存位置。因此，理想上当 thread 的数目够多的时候，就可以完全把 global memory 的巨大 latency 隐藏起来了。
如何利用 GPU 进行并行化计算，并且增加内存带宽，减少时钟周期的方法会在下一篇文章中讲解！！！