CUDA的安装

目录
• • • • 1.cuda安装及配置 2.cuda基本介绍 3.cuda编程实例 4.总结
用CUDA求随机数的平方和
首先，用以下函数产生一系列随机数，随机数的数 量我们设定为1048576： void GenerateNumbers(int *number, int size) { for(int i = 0; i < size; i++) { number[i] = rand() % 10; } }
利用cuda进行计算之前，在主函数中我们加入如下的代 码。 GenerateNumbers(data, DATA_SIZE); int* gpudata, *result; cudaMalloc((void**) &gpudata, sizeof(int) * DATA_SIZE); cudaMalloc((void**) &result, sizeof(int)); cudaMalloc((void**) &time, sizeof(clock_t)); cudaMemcpy(gpudata, data, sizeof(int) * DATA_SIZE, cudaMemcpyHostToDevice);
int final_sum = 0; for(int i = 0; i < THREAD_NUM; i++) { final_sum += sum[i]; } printf("sum: %d time: %d\n", final_sum, time_used); final_sum = 0; for(int i = 0; i < DATA_SIZE; i++) { sum += data[i] * data[i]; } printf("sum (CPU): %d\n", final_sum);
改进1：把随机的1048576个数分成256组，每组有4096个数，然 后利用256个线程，每个线程处理一组数据。 __global__ static void sumOfSquares(int *num, int* result, clock_t* time) { const int tid = threadIdx.x;（线程数为256） const int size = DATA_SIZE / THREAD_NUM; int sum = 0; int i; clock_t start; if(tid == 0) start = clock(); for(i = tid * size; i < (tid + 1) * size; i++) { sum += num[i] * num[i]; } result[tid] = sum; if(tid == 0) *time = clock() - start; } 主函数里面改为：
最后，在主函数里面把结果复制回内存上
• sumOfSquares<<<1, 1, 0>>>(gpudata, result,time); • int sum ; clock_t time_used; • cudaMemcpy(&sum, result, sizeof(int), cudaMemcpyDeviceToHost); • cudaMemcpy(&time_used, time, sizeof(clock_t), cudaMemcpyDeviceToHost); • cudaFree(gpudata); • cudaFree(result); • cudaFree(time); • printf("sum: %d time: %d\n", sum, time_used); • sum = 0; • for(int i = 0; i < DATA_SIZE; i++) { sum += data[i] * data[i]; } printf("sum (CPU): %d\n", sum); • 程序执行所花费的时间是756384988。
• 在 CUDA中，要执行一个函式，使用以下的语法： 函式名称<<<block 数目, thread 数目, shared memory 大小>>>(参数...)
__syncthreads(); 同步块中的所有线程。一旦所有线程均达到 此同步点，执行将正常恢复。 clock_t clock(); 在设备代码中执行时，返回随每一次时钟周 期而递增的每个多处理器计数器的值。
• • • • 1.cuda安装及配置 2.cuda基本介绍 3.cuda编程实例 4.总结
cuda的运行模型
cuda的内存模型
2.基本的一些函数 • cudaMalloc((void**)&data,size) 用来取得显卡内存。 cudaFree()程序执行完后，释放取得的内存。 • cudaMemcpy(&data1, &data2,size, ) 第四个参数指示内存的复制方向 cudaMemcpyHostToDevice cudaMemcpyDeviceToHost
主函数修改为： int* gpudata, *result; clock_t* time; cudaMalloc((void**) &gpudata, sizeof(int) * DATA_SIZE); cudaMalloc((void**) &result, sizeof(int) * THREAD_NUM * BLOCK_NUM); cudaMalloc((void**) &time, sizeof(clock_t) * BLOCK_NUM * 2); cudaMemcpy(gpudata, data, sizeof(int) * DATA_SIZE,cudaMemcpyHostToDevice); sumOfSquares<<<B LOCK_NUM, THREAD_NUM, 0>>>(gpudata, result, time); int sum[THREAD_NUM * BLOCK_NUM]; clock_t time_used[BLOCK_NUM * 2]; cudaMemcpy(&sum, result, sizeof(int) * THREAD_NUM * BLOCK_NUM,cudaMemcpyDeviceToHost); cudaMemcpy(&tim e_used, time, sizeof(clock_t) * BLOCK_NUM * 2,cudaMemcpyDeviceToHost); cudaFree(gpudata); cudaFree (result); cudaFree(time);
在HKEY_LOCAL_MACHINE\SOFTWARE Microsoft\Visua lStudio\8.0\Languages\File Extensions\下面添加子键 .cu ， 然后copy .cpp的键值到.cu。这样才能表示cu也是VS下的VC的 工程文件。
在HKEY_CURRENT_USER\Software\Whole Tomato\Visual Assist X\VANet8 在ExtSource键添加键值.cu。
通过一次改进后得到的时间为8731744，比上一个程序快了87倍， 速度相当可观。
改进2：利用256线程对所有数据进行处理，让 thread 0 读取第一 个数字，thread 1 读取第二个数字…依此类推 __global__ static void sumOfSquares(int *num, int* result,clock_t* time) { const int tid = threadIdx.x（线程数是256）; int sum = 0; int i; clock_t start; if(tid == 0) start = clock(); for(i = tid; i < DATA_SIZE; i += THREAD_NUM) { sum += num[i] * num[i]; } result[tid] = sum; if(tid == 0) *time = clock() - start; }
int* gpudata, *result; clock_t* time; cudaMalloc((void**) &gpudata, sizeof(int) * DATA_SIZE); cudaMalloc((void**) &result, sizeof(int) *THREAD_NUM); cudaMalloc((void**) &time, sizeof(clock_t)); cudaMemcpy(gpudata, data, sizeof(int) * DATA_SIZE, cudaMemcpyHostToDevice); sumOfSquares<<<1, THREAD_NUM, 0>>>(gpudata, result, time); int sum[THREAD_NUM]; clock_t time_used; cudaMemcpy(&sum, result, sizeof(int) * THREAD_NUM, cudaMemcpyDeviceToHost); cudaMemcpy(&time_used, time, sizeof(clock_t), cudaMemcpyDeviceToHost); cudaFree(gpudata); cudaFree(result); cudaFree(time); 在CPU端把各组数据的平方和加总：
在显卡上面写代码，在函式前面加上 __global__ 表示这 个函式是要在显示芯片上执行的，代码如下：
__global__ static void sumOfSquares (int *num, int* result,clock_t* time) { int sum = 0；int i； clock_t start = clock(); for(i = 0; i < DATA_SIZE; i++) { sum += num[i] * num[i]; } *result = sum; *time = clock() - start; }
Thead3
Thread255
…
1
2
3
4
256
257
258
512
改进3：建立32个blocks，每个 block有256个线程。这样总共有 32*256 = 8192 个 threads。 __global__ static void sumOfSquares(int *num, int* result, clock_t* time) { const int tid = threadIdx.x;（256个线程） const int bid = blockIdx.x; （32个块） int sum = 0; int i; if(tid == 0) time[bid] = clock(); for(i = bid * THREAD_NUM + tid; i < DATA_SIZE; i += BLOCK_NUM * THREAD_NUM) { sum += num[i] * num[i]; } result[bid * THREAD_NUM + tid] = sum; if(tid == 0) time[bid + BLOCK_NUM] = clock(); }
程序执行以后，记录的时间为3086794，比上一次又快了接近3倍，把线程数增 加到512，记录的时间为1783682，效率又提高了。
上面两种ຫໍສະໝຸດ Baidu同的内存读取方式
Thread0 Thread1 Thread…
1-4096
4097-…
……
Thread 取内存中 随机数的顺序
Thread0
Thread1
Thread2
CUDA深入浅出学习
于文茂
目录
• • • • 1.cuda安装及配置 2.cuda基本介绍 3.cuda编程实例 4.总结
cuda安装及配置
*1.分别安装cuda驱动、Toolkit、SDK三个文件。 *2.Vs2005选择工具/选项/文本编辑器/文件扩展名，
然后在扩展名里添加cu，编辑器选择MVC++（默 认是MVBasic） *3.下载安装CUDA VS2005Wizard * 4 .修改regedit，Visual Assist 添加支持*.cu文件
*5.打开Visual Assist属性，在projects 的C/C++ Directories custom下面添加 CUDA的头文件目录，这样才能在Visual Assist 生成规则的时候找到CUDA自身 的特殊定义才能生成Visual Assist的关键字，如__global__.
目录