借一栗子讲解基于C的CUDA并行计算

Love The Way You Lie 2022-11-14 11:49 101阅读 0赞

一、C语言接口回顾

代码1

  1. // ConsoleApplication6.cpp : 定义控制台应用程序的入口点。
  2. //
  4. #include "stdafx.h"
  5. #include <iostream>
  7. using namespace std;
  9. //自定义数据类型 数据对齐
  10. typedef struct student
  11. {
  13.     char label[1]; //1个字节
  14.     int number;    //4个字节
  15.     float score;   //4个字节
  16. }stu;
  18. int _tmain(int argc, _TCHAR* argv[])
  19. {
  21.     //在cpu-memory模式下面 数据是需要对齐的 因为这样cpu和memory的内存交互速度是最快的 
  22.     //在硬件设计的时候 就已经把这种属性 注入到硬件的平台 
  23.     //但是在gpu平台上面 对于数据对齐的处理要求会更严格
  24.     //因为并行计算 需要数据的提供延时更小 所以在不是4字节的类型里面 会自动进行填充 
  25.     //也就是说在gpu平台上面 这个student结构体12字节
  26.     std::cout << "sizeof(stu):" << sizeof(stu) << endl;
  27.     // printf(sizeof(stu));
  28.     //void*类型作为一个万能类型 专门给强制类型转换预留的接口 内核驱动里面和回掉函数里面 会有大量应用
  29.     // malloc free memset memcpy
  30.     char *p = (char *)malloc(100 * sizeof(char));
  32.     //申请的空间里面的值全部赋值成0 
  33.     memset(p, 0, 100);
  35.     char p_str[50] = "I love C++ program!!!";
  36.     //char *pstr = "I love C program!!!";
  38.     //把内存里面的值进行拷贝
  39.     memcpy(p, p_str, strlen(p_str));
  41.     //格式化输出
  42.     printf("%s", p);
  44.     //申请的内存空间需要释放 告诉编译器现在这段空间我不需要了 编译器进而通知系统 这段空间我现在不需要了 最后系统把这段空间重新利用
  45.     free(p);
  47.     system("pause");
  48.     return 0;
  49. }

总结：

  1. c/c++写一段代码流程：
    1.申请内存 malloc / new
    2.内存初始化 memset
    3.处理数据(此处是复制数据，memcpy)
    4.输出 （printf）
    5.释放内存 free/ delete
  2. 自定义数据，数据类型对齐
    在cpu-memory模式下面，数据是需要对齐的，因为这样cpu和memory的内存交互速度是最快的。
    在硬件设计的时候，就已经把这种属性，注入到硬件的平台。
    但是在gpu平台上面，对于数据对齐的处理要求会更严格。
    因为并行计算，需要数据的提供延时更小，所以在不是4字节的类型里面，会自动进行填充 。
    void*类型作为一个万能类型，专门给强制类型转换预留的接口，内核驱动里面和回掉函数里面会有大量应用。

二、 cuda编程基础概念

基础概念
主机：cpu和内存
设备：gpu和显存
API：
warp： thread
访问速度不同
变量类型限定符：device shared constant
函数类型限定符号 global 在cpu上定义，在gpu上显存上执行
device gpu上执行
host 主机上执行 cpu

thread：线程 cpu架构基本一致 同步
多少个流就可以创建多少个线程
block： 多个线程
grid: 多个block
SIMT：单指令多线程
内置变量：threadid blockid blockdim gridid griddim
显存带宽犹如高速公路的车道，优化的方法是尽可能接近理论的带宽。

代码2

  1. /***************************************************************************
  2. * first_cuda.cu
  3. *1.将数据从主机内存数据复制到设备显存
  4. *2.写好核函数
  5. *3.CUDA编译器执行核函数 在GPU上完成计算操作
  6. *4.把显存数据复制到主机内存
  7. *5.释放显存空间
  8. /***************************************************************************/
  10. #include <stdio.h>
  11. #include <stdlib.h>
  13. //CUDA RunTime API
  14. #include <cuda_runtime.h>
  15. #define DATA_SIZE 1048576
  16. int data[DATA_SIZE];
  18. //产生大量0-9之间的随机数
  19. // 指针不带内存大小，所以需要传入一个size，指明大小
  20. void GenerateNumbers(int *number, int size)
  21. {
  23.     for (int i = 0; i < size; i++) {
  25.         number[i] = rand() % 10;
  26.     }
  27. }
  29. //CUDA 初始化
  30. bool InitCUDA()
  31. {
  33.     int count;
  34.     //取得支持Cuda的装置的数目
  35.     cudaGetDeviceCount(&count);
  36.     if (count == 0) {
  38.         fprintf(stderr, "There is no device.\n");
  39.         return false;
  40.     }
  41.     int i;
  42.     for (i = 0; i < count; i++) {
  44.         cudaDeviceProp prop;
  45.         if (cudaGetDeviceProperties(&prop, i) == cudaSuccess) {
  47.             if (prop.major >= 1) {
  49.                 break;
  50.             }
  51.         }
  52.     }
  53.     if (i == count) {
  55.         fprintf(stderr, "There is no device supporting CUDA 1.x.\n");
  56.         return false;
  57.     }
  58.     cudaSetDevice(i);
  59.     return true;
  60. }
  62. // __global__ 函数 (GPU上执行) 计算立方和 
  63. //核函数是不可以有返回值类型的
  64. __global__ static void sumOfSquares(int *num, int* result)
  65. {
  67.     int sum = 0;
  68.     int i;
  69.     for (i = 0; i < DATA_SIZE; i++) {
  71.         sum += num[i] * num[i] * num[i];
  72.     }
  73.     *result = sum;
  74. }
  76. int main()
  77. {
  79.     //CUDA 初始化
  80.     if (!InitCUDA()) {
  82.         return 0;
  83.     }
  84.     //生成随机数
  85.     GenerateNumbers(data, DATA_SIZE);
  86.     /*把数据复制到显卡内存中*/
  87.     int* gpudata, *result;
  88.     //cudaMalloc 取得一块显卡内存 ( 其中result用来存储计算结果 )
  89.     cudaMalloc((void**)&gpudata, sizeof(int)* DATA_SIZE);
  90.     cudaMalloc((void**)&result, sizeof(int));
  91.     //cudaMemcpy 将产生的随机数复制到显卡内存中 
  92.     //cudaMemcpyHostToDevice - 从内存复制到显卡内存
  93.     //cudaMemcpyDeviceToHost - 从显卡内存复制到内存
  94.     cudaMemcpy(gpudata, data, sizeof(int)* DATA_SIZE, cudaMemcpyHostToDevice);
  95.     // 在CUDA 中执行函数 语法：函数名称<<<block 数目, thread 数目, shared memory 大小>>>(参数...);
  96.     sumOfSquares << <1, 1, 0 >> >(gpudata, result);
  97.     /*把结果从显示芯片复制回主内存*/
  98.     int sum;
  99.     //cudaMemcpy 将结果从显存中复制回内存
  100.     cudaMemcpy(&sum, result, sizeof(int), cudaMemcpyDeviceToHost);
  101.     //Free
  102.     cudaFree(gpudata);
  103.     cudaFree(result);
  104.     printf("GPUsum: %d \n", sum);
  105.     sum = 0;
  106.     for (int i = 0; i < DATA_SIZE; i++) {
  108.         sum += data[i] * data[i] * data[i];
  109.     }
  110.     printf("CPUsum: %d \n", sum);
  111.     return 0;
  112. }

总结cuda编程的流程：

  1. 申请显存空间 cudaMalloc
  2. 将数据从主机内存数据复制到设备显存
  3. 写好核函数
  4. CUDA编译器执行核函数 在GPU上完成计算操作
  5. 把显存数据复制到主机内存
  6. 释放显存空间

API：

  1. 获取 CUDA 设备数：
    函数原型：

    cudaError_t cudaGetDeviceCount( int* count )

可以通过 cudaGetDeviceCount 函数获取 CUDA 的设备数，实例如上InitCUDA()函数通过引用传递 count 值，获取当前支持的 CUDA 设备数，即返回具有计算能力的设备的数量。
以 *count 形式返回可用于执行的计算能力大于等于 1.0 的设备数量。如果不存在此类设备，将返回 1
cudaSuccess，注意，如果之前是异步启动，该函数可能返回错误码。

  1. 获取 CUDA 设备属性
    函数原型：

    cudaError_t cudaGetDeviceProperties( struct cudaDeviceProp* prop,int dev )

可以通过 cudaGetDeviceProperties 函数获取 CUDA 设备的属性，具体用法
函数通过引用传递 prop 关于属性的结构体，并且列出主设备号大于 1 的设备属性，其中设备属性通过函数 printDeviceProp 打印。
以*prop形式返回设备dev的属性。
返回值：cudaSuccess、cudaErrorInvalidDevice，注，如果之前是异步启动，该函数可能返回错误码。

  1. 设置 CUDA 设备
    通过函数 cudaSetDevice 就可以设置 CUDA 设备了，具体用法

    // set cuda device
    cudaSetDevice(i);

原型：

  1. cudaError_t cudaSetDevice(int dev)

将dev记录为活动主线程将执行设备码的设备。
cudaSuccess、cudaErrorInvalidDevice，注，如果之前是异步启动，该函数可能返回错误码。
cudaDeviceProp 结构定义如下：

  1. struct cudaDeviceProp {
  3. char name [256];
  4. size_t totalGlobalMem;
  5. size_t sharedMemPerBlock;
  6. int regsPerBlock;
  7. int warpSize;
  8. size_t memPitch;
  9. int maxThreadsPerBlock;
  10. int maxThreadsDim [3];
  11. int maxGridSize [3];
  12. size_t totalConstMem;
  13. int major;
  14. int minor;
  15. int clockRate;
  16. size_t textureAlignment;
  17. int deviceOverlap;
  18. int multiProcessorCount;
  19. }

CUDA 初始化完整代码函数为InitCUDA()
cudaDeviceProp 结构中的各个变量意义如下：
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：设备上多处理器的数量。

三、 CUDA程序优化和评估的方式

代码3：

  1. #include <stdio.h>
  2. #include <stdlib.h>
  3. #include <time.h>
  5. //CUDA RunTime API
  6. #include <cuda_runtime.h>
  8. //1024 * 1024  1M
  9. #define DATA_SIZE 1048576
  11. int data[DATA_SIZE];
  12. float clockRate = 1.0;
  14. //产生大量0-9之间的随机数
  15. void GenerateNumbers(int *number, int size)
  16. {
  18.     for (int i = 0; i < size; i++) {
  20.         number[i] = rand() % 10;
  21.     }
  22. }
  24. //打印设备属性
  25. void printDeviceProp(const cudaDeviceProp &prop)
  26. {
  28.     printf("Device Name : %s.\n", prop.name);
  29.     printf("totalGlobalMem : %d.\n", prop.totalGlobalMem);
  30.     printf("sharedMemPerBlock : %d.\n", prop.sharedMemPerBlock);
  31.     printf("regsPerBlock : %d.\n", prop.regsPerBlock);
  32.     printf("warpSize : %d.\n", prop.warpSize);
  33.     printf("memPitch : %d.\n", prop.memPitch);
  34.     printf("maxThreadsPerBlock : %d.\n", prop.maxThreadsPerBlock);
  35.     printf("maxThreadsDim[0 - 2] : %d %d %d.\n", prop.maxThreadsDim[0], prop.maxThreadsDim[1], prop.maxThreadsDim[2]);
  36.     printf("maxGridSize[0 - 2] : %d %d %d.\n", prop.maxGridSize[0], prop.maxGridSize[1], prop.maxGridSize[2]);
  37.     printf("totalConstMem : %d.\n", prop.totalConstMem);
  38.     printf("major.minor : %d.%d.\n", prop.major, prop.minor);
  39.     printf("clockRate : %d.\n", prop.clockRate);
  40.     printf("textureAlignment : %d.\n", prop.textureAlignment);
  41.     printf("deviceOverlap : %d.\n", prop.deviceOverlap);
  42.     printf("multiProcessorCount : %d.\n", prop.multiProcessorCount);
  43. }
  45. //CUDA 初始化
  46. bool InitCUDA()
  47. {
  49.     int count;
  50.     //取得支持Cuda的装置的数目
  51.     cudaGetDeviceCount(&count);
  52.     if (count == 0) {
  54.         fprintf(stderr, "There is no device.\n");
  55.         return false;
  56.     }
  57.     int i;
  58.     //取得显卡属性
  59.     for (i = 0; i < count; i++) {
  61.         cudaDeviceProp prop;
  62.         cudaGetDeviceProperties(&prop, i);
  63.         //打印gpu设备信息
  64.         printDeviceProp(prop);
  65.         //获得显卡的始终频率
  66.         clockRate = prop.clockRate;
  67.         if (cudaGetDeviceProperties(&prop, i) == cudaSuccess) {
  69.             if (prop.major >= 1) {
  71.                 break;
  72.             }
  73.         }
  74.     }
  75.     if (i == count) {
  77.         fprintf(stderr, "There is no device supporting CUDA 1.x.\n");
  78.         return false;
  79.     }
  80.     cudaSetDevice(i);
  81.     return true;
  82. }
  84. // __global__ 函数 (GPU上执行) 计算立方和
  85. __global__ static void sumOfSquares(int *num, int* result, clock_t* time)
  86. {
  88.     int sum = 0;
  89.     int i;
  90.     clock_t start = clock();
  91.     for (i = 0; i < DATA_SIZE; i++) {
  93.         sum += num[i] * num[i] * num[i];
  94.     }
  95.     *result = sum;
  96.     *time = clock() - start;
  97. }
  98. int main()
  99. {
  101.     //CUDA 初始化
  102.     if (!InitCUDA()) {
  104.         return 0;
  105.     }
  106.     //生成随机数(初始化全局的数组，指针不指定申请内存大小)
  107.     GenerateNumbers(data, DATA_SIZE);
  108.     /*把数据复制到显卡内存中*/
  109.     int* gpudata, *result;
  110.     clock_t* time;
  111.     //cudaMalloc 取得一块显卡内存 ( 其中result用来存储计算结果，time用来存储运行时间 )
  112.     cudaMalloc((void**)&gpudata, sizeof(int)* DATA_SIZE);
  113.     cudaMalloc((void**)&result, sizeof(int));
  114.     cudaMalloc((void**)&time, sizeof(clock_t));
  115.     //cudaMemcpy 将产生的随机数复制到显卡内存中
  116.     //cudaMemcpyHostToDevice - 从内存复制到显卡内存
  117.     //cudaMemcpyDeviceToHost - 从显卡内存复制到内存
  118.     cudaMemcpy(gpudata, data, sizeof(int)* DATA_SIZE, cudaMemcpyHostToDevice);
  119.     // 在CUDA 中执行函数 语法：函数名称<<<block 数目, thread 数目, shared memory 大小>>>(参数...);
  120.     sumOfSquares << <1, 1, 0 >> >(gpudata, result, time);
  121.     /*把结果从显示芯片复制回主内存*/
  122.     int sum;
  123.     clock_t time_used;
  124.     //cudaMemcpy 将结果从显存中复制回内存，结果在cpu上显示出来
  125.     cudaMemcpy(&sum, result, sizeof(int), cudaMemcpyDeviceToHost);
  126.     cudaMemcpy(&time_used, time, sizeof(clock_t), cudaMemcpyDeviceToHost);
  127.     //Free
  128.     cudaFree(gpudata);
  129.     cudaFree(result);
  130.     cudaFree(time);
  131.     printf("GPUsum: %d time_clock: %d time: %fs\n", sum, time_used, ((float)time_used / (clockRate * 1000)));
  132.     sum = 0;
  133.     for (int i = 0; i < DATA_SIZE; i++) {
  134.       //循环1M的数据
  135.         sum += data[i] * data[i] * data[i];
  136.     }
  137.     printf("CPUsum: %d \n", sum);
  138.     return 0;
  139. }

以上代码分析以及优化方式：
1M的数据（1024*1024） int类型是4个字节，即4byte数据 ，一共是的数据量就是4M
带宽：4MB/1.026178 = 3.89MB/s （4MB的数据，用时是1.026178s，计算得出3.89MB/s）
我显卡的带宽是14Gb/s左右 也就是说 完全没有实现并行计算的威力
优化：

  1. 一定要先从显存带宽开始
  2. 确定任务中并行和串行的算法
  3. 需要两层线程并行的内核函数，每个SM上面至少有6个warp和2个block
  4. 共享内存 shared memory

nvcc 编译代码
nvcc 是 CUDA 的编译工具，它可以将 .cu 文件解析出在 GPU 和 host 上执行的部分，也就是说，它会帮忙把 GPU 上执行和主机上执行的代码区分开来，不许要我们手动去做了。在 GPU 执行的部分会通过 NVIDIA 提供的 编译器编译成中介码，主机执行的部分则调用 gcc 编译。

  1. nvcc -o first_cuda first_cuda.cu

通过上述编译，生成可执行文件 first_cuda

四、 thread多线程概念的引入

cuda中GPU 的架构。它是由 grid 组成，每个 grid 又可以由 block 组成，而每个 block 又可以细分为 thread。所以，线程是我们处理的最小的单元。
接下来的例子通过修改前一个例子，把数组分割成若干个组（每个组由一个线程实现），每个组计算出一个和，然后在 CPU 中将分组的这几个和加在一起，得到最终的结果。这种思想叫做归约 。其实和分治思想差不多，就是先将大规模问题分解为小规模的问题，最后这些小规模问题整合得到最终解。
由于我的 GPU 支持的块内最大的线程数是 256个，即 cudaGetDeviceProperties 中的 maxThreadsPerBlock 属性。如何获取这个属性。
我们使用 256 个线程来实现并行加速。

代码4：thread多线程概念的引入:用时0.0244065s

  1. #include <stdio.h>
  2. #include <stdlib.h>
  3. #include <time.h>
  5. //CUDA RunTime API
  6. #include <cuda_runtime.h>
  8. //1M
  9. #define DATA_SIZE 1048576
  10. #define THREAD_NUM 256  //256线程
  12. int data[DATA_SIZE];
  13. int clockRate = 1.0;
  15. //产生大量0-9之间的随机数
  16. void GenerateNumbers(int *number, int size)
  17. {
  19.     for (int i = 0; i < size; i++) {
  21.         number[i] = rand() % 10;
  22.     }
  23. }
  25. //打印设备信息
  26. void printDeviceProp(const cudaDeviceProp &prop)
  27. {
  29.     printf("Device Name : %s.\n", prop.name);
  30.     printf("totalGlobalMem : %d.\n", prop.totalGlobalMem);
  31.     printf("sharedMemPerBlock : %d.\n", prop.sharedMemPerBlock);
  32.     printf("regsPerBlock : %d.\n", prop.regsPerBlock);
  33.     printf("warpSize : %d.\n", prop.warpSize);
  34.     printf("memPitch : %d.\n", prop.memPitch);
  35.     printf("maxThreadsPerBlock : %d.\n", prop.maxThreadsPerBlock);
  36.     printf("maxThreadsDim[0 - 2] : %d %d %d.\n", prop.maxThreadsDim[0], prop.maxThreadsDim[1], prop.maxThreadsDim[2]);
  37.     printf("maxGridSize[0 - 2] : %d %d %d.\n", prop.maxGridSize[0], prop.maxGridSize[1], prop.maxGridSize[2]);
  38.     printf("totalConstMem : %d.\n", prop.totalConstMem);
  39.     printf("major.minor : %d.%d.\n", prop.major, prop.minor);
  40.     printf("clockRate : %d.\n", prop.clockRate);
  41.     printf("textureAlignment : %d.\n", prop.textureAlignment);
  42.     printf("deviceOverlap : %d.\n", prop.deviceOverlap);
  43.     printf("multiProcessorCount : %d.\n", prop.multiProcessorCount);
  44. }
  46. //CUDA 初始化
  47. bool InitCUDA()
  48. {
  50.     int count;
  51.     //取得支持Cuda的装置的数目
  52.     cudaGetDeviceCount(&count);
  53.     if (count == 0) {
  55.         fprintf(stderr, "There is no device.\n");
  56.         return false;
  57.     }
  58.     int i;
  59.     for (i = 0; i < count; i++) {
  61.         cudaDeviceProp prop;
  62.         cudaGetDeviceProperties(&prop, i);
  63.         //打印设备信息
  64.         printDeviceProp(prop);
  65.         //获得显卡的时钟频率
  66.         clockRate = prop.clockRate;
  67.         if (cudaGetDeviceProperties(&prop, i) == cudaSuccess) {
  69.             if (prop.major >= 1) {
  71.                 break;
  72.             }
  73.         }
  74.     }
  76.     if (i == count) {
  78.         fprintf(stderr, "There is no device supporting CUDA 1.x.\n");
  79.         return false;
  80.     }
  81.     cudaSetDevice(i);
  82.     return true;
  83. }
  86. // __global__ 函数 (GPU上执行) 计算立方和
  87. __global__ static void sumOfSquares(int *num, int* result, clock_t* time)
  88. {
  90.     //表示目前的 thread 是第几个 thread（由 0 开始计算）
  91.     const int tid = threadIdx.x; //tid相当于一个向量，大小为256
  92.     //计算每个线程需要完成的量
  93.     const int size = DATA_SIZE / THREAD_NUM;
  94.     int sum = 0;
  95.     int i;
  96.     //记录运算开始的时间
  97.     clock_t start;
  98.     //只在 thread 0（即 threadIdx.x = 0 的时候）进行记录
  99.     if (tid == 0) start = clock();
  100.     for (i = tid * size; i < (tid + 1) * size; i++) {
  102.         sum += num[i] * num[i] * num[i];
  103.     }
  104.     result[tid] = sum;
  105.     //计算时间的动作，只在 thread 0（即 threadIdx.x = 0 的时候）进行
  106.     if (tid == 0) *time = clock() - start;
  107. }
  108. int main()
  109. {
  111.     //CUDA 初始化
  112.     if (!InitCUDA()) {
  114.         return 0;
  115.     }
  116.     //生成随机数
  117.     GenerateNumbers(data, DATA_SIZE);
  118.     /*把数据复制到显卡内存中*/
  119.     int* gpudata, *result;
  120.     clock_t* time;
  121.     //cudaMalloc 取得一块显卡内存 ( 其中result用来存储计算结果，time用来存储运行时间 )
  122.     cudaMalloc((void**)&gpudata, sizeof(int)* DATA_SIZE);
  123.     cudaMalloc((void**)&result, sizeof(int)*THREAD_NUM);
  124.     cudaMalloc((void**)&time, sizeof(clock_t));
  125.     //cudaMemcpy 将产生的随机数复制到显卡内存中
  126.     //cudaMemcpyHostToDevice - 从内存复制到显卡内存
  127.     //cudaMemcpyDeviceToHost - 从显卡内存复制到内存
  128.     cudaMemcpy(gpudata, data, sizeof(int)* DATA_SIZE, cudaMemcpyHostToDevice);
  129.     // 在CUDA 中执行函数 语法：函数名称<<<block 数目, thread 数目, shared memory 大小>>>(参数...);
  130.     sumOfSquares << < 1, THREAD_NUM, 0 >> >(gpudata, result, time);
  131.     /*把结果从显示芯片复制回主内存*/
  132.     int sum[THREAD_NUM];
  133.     clock_t time_use;
  134.     //cudaMemcpy 将结果从显存中复制回内存
  135.     cudaMemcpy(&sum, result, sizeof(int)* THREAD_NUM, cudaMemcpyDeviceToHost);
  136.     cudaMemcpy(&time_use, time, sizeof(clock_t), cudaMemcpyDeviceToHost);
  137.     //Free
  138.     cudaFree(gpudata);
  139.     cudaFree(result);
  140.     cudaFree(time);
  141.     int final_sum = 0;
  142.     for (int i = 0; i < THREAD_NUM; i++) {
  144.         final_sum += sum[i];
  145.     }
  146.     printf("GPUsum: %d time_clock: %d time: %fs\n", final_sum, time_use, ((float)time_use / (clockRate * 1000)));
  147.     final_sum = 0;
  148.     for (int i = 0; i < DATA_SIZE; i++) {
  150.         final_sum += data[i] * data[i] * data[i];
  151.     }
  152.     printf("CPUsum: %d \n", final_sum);
  153.     return 0;
  154. }
  155. //创建了256个线程帮我们并行计算 我们需要提前安排好每个线程计算的数据 防止线程同步的问题

重点代码分析

  1. __global__ static void sumOfSquares(int *num, int* result, clock_t* time)
  2. {
  4.  //表示目前的 thread 是第几个 thread（由 0 开始计算）
  5.  const int tid = threadIdx.x; //tid相当于一个向量，大小为256
  6.  //计算每个线程需要完成的量
  7.  const int size = DATA_SIZE / THREAD_NUM;
  8.  int sum = 0;
  9.  int i;
  10.  //记录运算开始的时间
  11.  clock_t start;
  12.  //只在 thread 0（即 threadIdx.x = 0 的时候）进行记录
  13.  if (tid == 0) start = clock();
  14.  for (i = tid * size; i < (tid + 1) * size; i++) {
  16.   sum += num[i] * num[i] * num[i];
  17.  }
  18.  result[tid] = sum;
  19.  //计算时间的动作，只在 thread 0（即 threadIdx.x = 0 的时候）进行
  20.  if (tid == 0) *time = clock() - start;
  21. }
  22. int main()
  23. {
  25.  //C

代码5：内存连续存储:用时0.0053375s
线程：0,1,2…,255, 0,1,2…255

  1. #include <stdio.h>
  2. #include <stdlib.h>
  3. #include <time.h>
  5. //CUDA RunTime API
  6. #include <cuda_runtime.h>
  8. //1M
  9. #define DATA_SIZE 1048576
  10. #define THREAD_NUM 256
  12. int data[DATA_SIZE];
  13. int clockRate = 1.0;
  15. //产生大量0-9之间的随机数
  16. void GenerateNumbers(int *number, int size)
  17. {
  19.     for (int i = 0; i < size; i++) {
  21.         number[i] = rand() % 10;
  22.     }
  23. }
  25. //打印设备信息
  26. void printDeviceProp(const cudaDeviceProp &prop)
  27. {
  29.     printf("Device Name : %s.\n", prop.name);
  30.     printf("totalGlobalMem : %d.\n", prop.totalGlobalMem);
  31.     printf("sharedMemPerBlock : %d.\n", prop.sharedMemPerBlock);
  32.     printf("regsPerBlock : %d.\n", prop.regsPerBlock);
  33.     printf("warpSize : %d.\n", prop.warpSize);
  34.     printf("memPitch : %d.\n", prop.memPitch);
  35.     printf("maxThreadsPerBlock : %d.\n", prop.maxThreadsPerBlock);
  36.     printf("maxThreadsDim[0 - 2] : %d %d %d.\n", prop.maxThreadsDim[0], prop.maxThreadsDim[1], prop.maxThreadsDim[2]);
  37.     printf("maxGridSize[0 - 2] : %d %d %d.\n", prop.maxGridSize[0], prop.maxGridSize[1], prop.maxGridSize[2]);
  38.     printf("totalConstMem : %d.\n", prop.totalConstMem);
  39.     printf("major.minor : %d.%d.\n", prop.major, prop.minor);
  40.     printf("clockRate : %d.\n", prop.clockRate);
  41.     printf("textureAlignment : %d.\n", prop.textureAlignment);
  42.     printf("deviceOverlap : %d.\n", prop.deviceOverlap);
  43.     printf("multiProcessorCount : %d.\n", prop.multiProcessorCount);
  44. }
  46. //CUDA 初始化
  47. bool InitCUDA()
  48. {
  50.     int count;
  51.     //取得支持Cuda的装置的数目
  52.     cudaGetDeviceCount(&count);
  53.     if (count == 0) {
  55.         fprintf(stderr, "There is no device.\n");
  56.         return false;
  57.     }
  58.     int i;
  59.     for (i = 0; i < count; i++) {
  61.         cudaDeviceProp prop;
  62.         cudaGetDeviceProperties(&prop, i);
  63.         //打印设备信息
  64.         printDeviceProp(prop);
  65.         //获得显卡的时钟频率
  66.         clockRate = prop.clockRate;
  67.         if (cudaGetDeviceProperties(&prop, i) == cudaSuccess) {
  69.             if (prop.major >= 1) {
  71.                 break;
  72.             }
  73.         }
  74.     }
  76.     if (i == count) {
  78.         fprintf(stderr, "There is no device supporting CUDA 1.x.\n");
  79.         return false;
  80.     }
  81.     cudaSetDevice(i);
  82.     return true;
  83. }
  86. // __global__ 函数 (GPU上执行) 计算立方和
  87. __global__ static void sumOfSquares(int *num, int* result, clock_t* time)
  88. {
  90.     //表示目前的 thread 是第几个 thread（由 0 开始计算）
  91.     const int tid = threadIdx.x;
  92.     int sum = 0;
  93.     int i;
  94.     //记录运算开始的时间
  95.     clock_t start;
  96.     //只在 thread 0（即 threadIdx.x = 0 的时候）进行记录
  97.     if (tid == 0) start = clock();
  98.     //改为连续存取（thread 0 读取第一个数字，thread 1 读取第二个数字 …）
  99.     for (i = tid; i < DATA_SIZE; i += THREAD_NUM) {
  101.         sum += num[i] * num[i] * num[i];
  102.     }
  103.     result[tid] = sum;
  104.     //计算时间的动作，只在 thread 0（即 threadIdx.x = 0 的时候）进行
  105.     if (tid == 0) *time = clock() - start;
  106. }
  107. int main()
  108. {
  110.     //CUDA 初始化
  111.     if (!InitCUDA()) {
  113.         return 0;
  114.     }
  115.     //生成随机数
  116.     GenerateNumbers(data, DATA_SIZE);
  117.     /*把数据复制到显卡内存中*/
  118.     int* gpudata, *result;
  119.     clock_t* time;
  120.     //cudaMalloc 取得一块显卡内存 ( 其中result用来存储计算结果，time用来存储运行时间 )
  121.     cudaMalloc((void**)&gpudata, sizeof(int)* DATA_SIZE);
  122.     cudaMalloc((void**)&result, sizeof(int)*THREAD_NUM);
  123.     cudaMalloc((void**)&time, sizeof(clock_t));
  124.     //cudaMemcpy 将产生的随机数复制到显卡内存中
  125.     //cudaMemcpyHostToDevice - 从内存复制到显卡内存
  126.     //cudaMemcpyDeviceToHost - 从显卡内存复制到内存
  127.     cudaMemcpy(gpudata, data, sizeof(int)* DATA_SIZE, cudaMemcpyHostToDevice);
  128.     // 在CUDA 中执行函数 语法：函数名称<<<block 数目, thread 数目, shared memory 大小>>>(参数...);
  129.     sumOfSquares << < 1, THREAD_NUM, 0 >> >(gpudata, result, time);
  130.     /*把结果从显示芯片复制回主内存*/
  131.     int sum[THREAD_NUM];
  132.     clock_t time_use;
  133.     //cudaMemcpy 将结果从显存中复制回内存
  134.     cudaMemcpy(&sum, result, sizeof(int)* THREAD_NUM, cudaMemcpyDeviceToHost);
  135.     cudaMemcpy(&time_use, time, sizeof(clock_t), cudaMemcpyDeviceToHost);
  136.     //Free
  137.     cudaFree(gpudata);
  138.     cudaFree(result);
  139.     cudaFree(time);
  140.     int final_sum = 0;
  141.     for (int i = 0; i < THREAD_NUM; i++) {
  143.         final_sum += sum[i];
  144.     }
  145.     printf("GPUsum: %d time_clock: %d time: %f\n", final_sum, time_use, ((float)time_use / (clockRate * 1000)));
  146.     final_sum = 0;
  147.     for (int i = 0; i < DATA_SIZE; i++) {
  149.         final_sum += data[i] * data[i] * data[i];
  150.     }
  151.     printf("CPUsum: %d \n", final_sum);
  152.     return 0;
  153. }
  154. //我们的程序需要尽可能连续操作内存，减少内存存取方面的时间浪费

五、block线程块概念

代码6：block线程块概念引入：用时：0.002125s
在这里插入图片描述

  1. #include <stdio.h>
  2. #include <stdlib.h>
  3. #include <time.h>
  5. //CUDA RunTime API
  6. #include <cuda_runtime.h>
  8. //1M
  9. #define DATA_SIZE 1048576  //1024*1024
  10. #define THREAD_NUM 256
  11. #define BLOCK_NUM 32  // block
  12. int clockRate = 1.0;
  13. int data[DATA_SIZE];
  15. //产生大量0-9之间的随机数
  16. void GenerateNumbers(int *number, int size)
  17. {
  19.     for (int i = 0; i < size; i++) {
  21.         number[i] = rand() % 10;
  22.     }
  23. }
  25. //打印设备信息
  26. void printDeviceProp(const cudaDeviceProp &prop)
  27. {
  29.     printf("Device Name : %s.\n", prop.name);
  30.     printf("totalGlobalMem : %d.\n", prop.totalGlobalMem);
  31.     printf("sharedMemPerBlock : %d.\n", prop.sharedMemPerBlock);
  32.     printf("regsPerBlock : %d.\n", prop.regsPerBlock);
  33.     printf("warpSize : %d.\n", prop.warpSize);
  34.     printf("memPitch : %d.\n", prop.memPitch);
  35.     printf("maxThreadsPerBlock : %d.\n", prop.maxThreadsPerBlock);
  36.     printf("maxThreadsDim[0 - 2] : %d %d %d.\n", prop.maxThreadsDim[0], prop.maxThreadsDim[1], prop.maxThreadsDim[2]);
  37.     printf("maxGridSize[0 - 2] : %d %d %d.\n", prop.maxGridSize[0], prop.maxGridSize[1], prop.maxGridSize[2]);
  38.     printf("totalConstMem : %d.\n", prop.totalConstMem);
  39.     printf("major.minor : %d.%d.\n", prop.major, prop.minor);
  40.     printf("clockRate : %d.\n", prop.clockRate);
  41.     printf("textureAlignment : %d.\n", prop.textureAlignment);
  42.     printf("deviceOverlap : %d.\n", prop.deviceOverlap);
  43.     printf("multiProcessorCount : %d.\n", prop.multiProcessorCount);
  44. }
  46. //CUDA 初始化
  47. bool InitCUDA()
  48. {
  50.     int count;
  51.     //取得支持Cuda的装置的数目
  52.     cudaGetDeviceCount(&count);
  53.     if (count == 0) {
  55.         fprintf(stderr, "There is no device.\n");
  56.         return false;
  57.     }
  58.     int i;
  59.     for (i = 0; i < count; i++) {
  61.         cudaDeviceProp prop;
  62.         cudaGetDeviceProperties(&prop, i);
  63.         //打印设备信息
  64.         printDeviceProp(prop);
  65.         //获得显卡的时钟频率
  66.         clockRate = prop.clockRate;
  67.         if (cudaGetDeviceProperties(&prop, i) == cudaSuccess) {
  69.             if (prop.major >= 1) {
  71.                 break;
  72.             }
  73.         }
  74.     }
  76.     if (i == count) {
  78.         fprintf(stderr, "There is no device supporting CUDA 1.x.\n");
  79.         return false;
  80.     }
  81.     cudaSetDevice(i);
  82.     return true;
  83. }
  86. // __global__ 函数 (GPU上执行) 计算立方和
  87. // __global__ 函数 (GPU上执行) 计算立方和
  88. __global__ static void sumOfSquares(int *num, int* result, clock_t* time)
  89. {
  91.     //表示目前的 thread 是第几个 thread（由 0 开始计算）0~255
  92.     const int tid = threadIdx.x;
  93.     //表示目前的 thread 属于第几个 block（由 0 开始计算）0~31
  94.     const int bid = blockIdx.x;
  95.     int sum = 0;
  96.     int i;
  97.     //记录运算开始的时间
  98.     clock_t start;
  99.     //只在 thread 0（即 threadIdx.x = 0 的时候）进行记录，每个 block 都会记录开始时间及结束时间
  100.     if (tid == 0) time[bid] = clock();
  101.     //thread需要同时通过tid和bid来确定，同时不要忘记保证内存连续性
  102.     for (i = bid * THREAD_NUM + tid; i < DATA_SIZE; i += BLOCK_NUM * THREAD_NUM) {
  104.         sum += num[i] * num[i] * num[i];
  105.     }
  106.     //Result的数量随之增加
  107.     result[bid * THREAD_NUM + tid] = sum;
  108.     //计算时间的动作，只在 thread 0（即 threadIdx.x = 0 的时候）进行，每个 block 都会记录开始时间及结束时间
  109.     if (tid == 0) time[bid + BLOCK_NUM] = clock();
  110. }
  111. int main()
  112. {
  114.     //CUDA 初始化
  115.     if (!InitCUDA()) {
  117.         return 0;
  118.     }
  119.     //生成随机数
  120.     GenerateNumbers(data, DATA_SIZE);
  121.     /*把数据复制到显卡内存中*/
  122.     int* gpudata, *result;
  123.     clock_t* time;
  124.     //cudaMalloc 取得一块显卡内存 ( 其中result用来存储计算结果，time用来存储运行时间 )
  125.     cudaMalloc((void**)&gpudata, sizeof(int)* DATA_SIZE);
  126.     cudaMalloc((void**)&result, sizeof(int)*THREAD_NUM* BLOCK_NUM);
  127.     cudaMalloc((void**)&time, sizeof(clock_t)* BLOCK_NUM * 2);
  128.     //cudaMemcpy 将产生的随机数复制到显卡内存中
  129.     //cudaMemcpyHostToDevice - 从内存复制到显卡内存
  130.     //cudaMemcpyDeviceToHost - 从显卡内存复制到内存
  131.     cudaMemcpy(gpudata, data, sizeof(int)* DATA_SIZE, cudaMemcpyHostToDevice);
  132.     // 在CUDA 中执行函数 语法：函数名称<<<block 数目, thread 数目, shared memory 大小>>>(参数...);
  133.     sumOfSquares << < BLOCK_NUM, THREAD_NUM, 0 >> >(gpudata, result, time);
  134.     /*把结果从显示芯片复制回主内存*/
  135.     int sum[THREAD_NUM*BLOCK_NUM];
  136.     clock_t time_use[BLOCK_NUM * 2];
  137.     //cudaMemcpy 将结果从显存中复制回内存
  138.     cudaMemcpy(&sum, result, sizeof(int)* THREAD_NUM*BLOCK_NUM, cudaMemcpyDeviceToHost);
  139.     cudaMemcpy(&time_use, time, sizeof(clock_t)* BLOCK_NUM * 2, cudaMemcpyDeviceToHost);
  140.     //Free
  141.     cudaFree(gpudata);
  142.     cudaFree(result);
  143.     cudaFree(time);
  144.     int final_sum = 0;
  145.     for (int i = 0; i < THREAD_NUM*BLOCK_NUM; i++) {
  147.         final_sum += sum[i];
  148.     }
  149.     //采取新的计时策略 把每个 block 最早的开始时间，和最晚的结束时间相减，取得总运行时间
  150.     clock_t min_start, max_end;
  151.     min_start = time_use[0];
  152.     max_end = time_use[BLOCK_NUM];
  153.     for (int i = 1; i < BLOCK_NUM; i++) {
  155.         if (min_start > time_use[i])
  156.             min_start = time_use[i];
  157.         if (max_end < time_use[i + BLOCK_NUM])
  158.             max_end = time_use[i + BLOCK_NUM];
  159.     }
  160.     printf("GPUsum: %d time_clock: %d time: %fs\n", final_sum, max_end - min_start, ((float)(max_end - min_start) / (clockRate * 1000)));
  161.     final_sum = 0;
  162.     for (int i = 0; i < DATA_SIZE; i++) {
  164.         final_sum += data[i] * data[i] * data[i];
  165.     }
  166.     printf("CPUsum: %d \n", final_sum);
  167.     return 0;
  168. }

六、共享内存与同步

代码7:共享内存机制：用时：0.003568s
看似用时比上一次多了，其实主要节约了cpu中for循环的计算用时，同时也减少了数据copy的用时。总体来说，cpu和gpu的总时间减少了。
线程同步问题。即一个block中所有的线程都操作完毕了。
一个block中的所有线程可以共享内存
通常，数据的量是大于线程的总和的

  1. #include <stdio.h>
  2. #include <stdlib.h>
  3. #include <time.h>
  5. //CUDA RunTime API
  6. #include <cuda_runtime.h>
  8. //1M
  9. #define DATA_SIZE 1048576
  10. #define THREAD_NUM 256
  11. #define BLOCK_NUM 32
  12. int clockRate = 1.0;
  13. int data[DATA_SIZE];
  14. //产生大量0-9之间的随机数
  15. void GenerateNumbers(int *number, int size)
  16. {
  18.     for (int i = 0; i < size; i++) {
  20.         number[i] = rand() % 10;
  21.     }
  22. }
  23. //打印设备信息
  24. void printDeviceProp(const cudaDeviceProp &prop)
  25. {
  27.     printf("Device Name : %s.\n", prop.name);
  28.     printf("totalGlobalMem : %d.\n", prop.totalGlobalMem);
  29.     printf("sharedMemPerBlock : %d.\n", prop.sharedMemPerBlock);
  30.     printf("regsPerBlock : %d.\n", prop.regsPerBlock);
  31.     printf("warpSize : %d.\n", prop.warpSize);
  32.     printf("memPitch : %d.\n", prop.memPitch);
  33.     printf("maxThreadsPerBlock : %d.\n", prop.maxThreadsPerBlock);
  34.     printf("maxThreadsDim[0 - 2] : %d %d %d.\n", prop.maxThreadsDim[0], prop.maxThreadsDim[1], prop.maxThreadsDim[2]);
  35.     printf("maxGridSize[0 - 2] : %d %d %d.\n", prop.maxGridSize[0], prop.maxGridSize[1], prop.maxGridSize[2]);
  36.     printf("totalConstMem : %d.\n", prop.totalConstMem);
  37.     printf("major.minor : %d.%d.\n", prop.major, prop.minor);
  38.     printf("clockRate : %d.\n", prop.clockRate);
  39.     printf("textureAlignment : %d.\n", prop.textureAlignment);
  40.     printf("deviceOverlap : %d.\n", prop.deviceOverlap);
  41.     printf("multiProcessorCount : %d.\n", prop.multiProcessorCount);
  42. }
  43. //CUDA 初始化
  44. bool InitCUDA()
  45. {
  47.     int count;
  48.     //取得支持Cuda的装置的数目
  49.     cudaGetDeviceCount(&count);
  50.     if (count == 0) {
  52.         fprintf(stderr, "There is no device.\n");
  53.         return false;
  54.     }
  55.     int i;
  56.     for (i = 0; i < count; i++) {
  58.         cudaDeviceProp prop;
  59.         cudaGetDeviceProperties(&prop, i);
  60.         //打印设备信息
  61.         printDeviceProp(prop);
  62.         //获得显卡的时钟频率
  63.         clockRate = prop.clockRate;
  64.         if (cudaGetDeviceProperties(&prop, i) == cudaSuccess) {
  66.             if (prop.major >= 1) {
  68.                 break;
  69.             }
  70.         }
  71.     }
  72.     if (i == count) {
  74.         fprintf(stderr, "There is no device supporting CUDA 1.x.\n");
  75.         return false;
  76.     }
  77.     cudaSetDevice(i);
  78.     return true;
  79. }
  80. // __global__ 函数 (GPU上执行) 计算立方和
  81. __global__ static void sumOfSquares(int *num, int* result, clock_t* time)
  82. {
  84.     //声明一块共享内存
  85.     extern __shared__ int shared[];  
  86.     //表示目前的 thread 是第几个 thread（由 0 开始计算）
  87.     const int tid = threadIdx.x;
  88.     //表示目前的 thread 属于第几个 block（由 0 开始计算）
  89.     const int bid = blockIdx.x;
  90.     shared[tid] = 0;
  91.     int i;
  92.     //记录运算开始的时间
  93.     clock_t start;
  94.     //只在 thread 0（即 threadIdx.x = 0 的时候）进行记录，每个 block 都会记录开始时间及结束时间
  95.     if (tid == 0) time[bid] = clock();
  96.     //thread需要同时通过tid和bid来确定，同时不要忘记保证内存连续性
  97.     for (i = bid * THREAD_NUM + tid; i < DATA_SIZE; i += BLOCK_NUM * THREAD_NUM) {
  99.         shared[tid] += num[i] * num[i] * num[i];
  100.     }
  101.     //同步 保证每个 thread 都已经把结果写到 shared[tid] 里面
  102.     //cuda编程里，很大程度再解决如何并行和串行不出题
  103.     __syncthreads();//共享内存的基础上
  104.     //使用线程0完成加和
  105.     if (tid == 0)
  106.     {
  108.         for (i = 1; i < THREAD_NUM; i++)
  109.         {
  111.             shared[0] += shared[i];
  112.         }
  113.         result[bid] = shared[0];
  114.     }
  115.     //计算时间的动作，只在 thread 0（即 threadIdx.x = 0 的时候）进行，每个 block 都会记录开始时间及结束时间
  116.     if (tid == 0) time[bid + BLOCK_NUM] = clock();
  117. }
  118. int main()
  119. {
  121.     //CUDA 初始化
  122.     if (!InitCUDA()) {
  124.         return 0;
  125.     }
  126.     //生成随机数
  127.     GenerateNumbers(data, DATA_SIZE);
  128.     /*把数据复制到显卡内存中*/
  129.     int* gpudata, *result;
  130.     clock_t* time;
  131.     //cudaMalloc 取得一块显卡内存 ( 其中result用来存储计算结果，time用来存储运行时间 )
  132.     cudaMalloc((void**)&gpudata, sizeof(int)* DATA_SIZE);
  133.     cudaMalloc((void**)&result, sizeof(int)* BLOCK_NUM);
  134.     cudaMalloc((void**)&time, sizeof(clock_t)* BLOCK_NUM * 2);
  135.     //cudaMemcpy 将产生的随机数复制到显卡内存中
  136.     //cudaMemcpyHostToDevice - 从内存复制到显卡内存
  137.     //cudaMemcpyDeviceToHost - 从显卡内存复制到内存
  138.     cudaMemcpy(gpudata, data, sizeof(int)* DATA_SIZE, cudaMemcpyHostToDevice);
  139.     // 在CUDA 中执行函数 语法：函数名称<<<block 数目, thread 数目, shared memory 大小>>>(参数...);
  140.     sumOfSquares << < BLOCK_NUM, THREAD_NUM, THREAD_NUM * sizeof(int) >> >(gpudata, result, time);
  141.     /*把结果从显示芯片复制回主内存*/
  142.     int sum[BLOCK_NUM];
  143.     clock_t time_use[BLOCK_NUM * 2];
  144.     //cudaMemcpy 将结果从显存中复制回内存
  145.     cudaMemcpy(&sum, result, sizeof(int)*BLOCK_NUM, cudaMemcpyDeviceToHost);
  146.     cudaMemcpy(&time_use, time, sizeof(clock_t)* BLOCK_NUM * 2, cudaMemcpyDeviceToHost);
  147.     //Free
  148.     cudaFree(gpudata);
  149.     cudaFree(result);
  150.     cudaFree(time);
  151.     int final_sum = 0;
  152.     for (int i = 0; i < BLOCK_NUM; i++) {
  154.         final_sum += sum[i];
  155.     }
  156.     //采取新的计时策略 把每个 block 最早的开始时间，和最晚的结束时间相减，取得总运行时间
  157.     clock_t min_start, max_end;
  158.     min_start = time_use[0];
  159.     max_end = time_use[BLOCK_NUM];
  160.     for (int i = 1; i < BLOCK_NUM; i++) {
  162.         if (min_start > time_use[i])
  163.             min_start = time_use[i];
  164.         if (max_end < time_use[i + BLOCK_NUM])
  165.             max_end = time_use[i + BLOCK_NUM];
  166.     }
  167.     printf("GPUsum: %d time_clock: %d time: %fs\n", final_sum, max_end - min_start, ((float)(max_end - min_start) / (clockRate * 1000)));
  168.     final_sum = 0;
  169.     for (int i = 0; i < DATA_SIZE; i++) {
  171.         final_sum += data[i] * data[i] * data[i];
  172.     }
  173.     printf("CPUsum: %d \n", final_sum);
  174.     return 0;
  175. }

七、树状求和算法加速

代码8：树状求和算法简介: 用时：0.002941s
算法优化

  1. #include <stdio.h>
  2. #include <stdlib.h>
  3. #include <time.h>
  5. //CUDA RunTime API
  6. #include <cuda_runtime.h>
  8. //1M
  9. #define DATA_SIZE 1048576
  10. #define THREAD_NUM 256
  11. #define BLOCK_NUM 32
  12. int clockRate = 1.0;
  13. int data[DATA_SIZE];
  14. //产生大量0-9之间的随机数
  15. void GenerateNumbers(int *number, int size)
  16. {
  18.     for (int i = 0; i < size; i++) {
  20.         number[i] = rand() % 10;
  21.     }
  22. }
  23. //打印设备信息
  24. void printDeviceProp(const cudaDeviceProp &prop)
  25. {
  27.     printf("Device Name : %s.\n", prop.name);
  28.     printf("totalGlobalMem : %d.\n", prop.totalGlobalMem);
  29.     printf("sharedMemPerBlock : %d.\n", prop.sharedMemPerBlock);
  30.     printf("regsPerBlock : %d.\n", prop.regsPerBlock);
  31.     printf("warpSize : %d.\n", prop.warpSize);
  32.     printf("memPitch : %d.\n", prop.memPitch);
  33.     printf("maxThreadsPerBlock : %d.\n", prop.maxThreadsPerBlock);
  34.     printf("maxThreadsDim[0 - 2] : %d %d %d.\n", prop.maxThreadsDim[0], prop.maxThreadsDim[1], prop.maxThreadsDim[2]);
  35.     printf("maxGridSize[0 - 2] : %d %d %d.\n", prop.maxGridSize[0], prop.maxGridSize[1], prop.maxGridSize[2]);
  36.     printf("totalConstMem : %d.\n", prop.totalConstMem);
  37.     printf("major.minor : %d.%d.\n", prop.major, prop.minor);
  38.     printf("clockRate : %d.\n", prop.clockRate);
  39.     printf("textureAlignment : %d.\n", prop.textureAlignment);
  40.     printf("deviceOverlap : %d.\n", prop.deviceOverlap);
  41.     printf("multiProcessorCount : %d.\n", prop.multiProcessorCount);
  42. }
  43. //CUDA 初始化
  44. bool InitCUDA()
  45. {
  47.     int count;
  48.     //取得支持Cuda的装置的数目
  49.     cudaGetDeviceCount(&count);
  50.     if (count == 0) {
  52.         fprintf(stderr, "There is no device.\n");
  53.         return false;
  54.     }
  55.     int i;
  56.     for (i = 0; i < count; i++) {
  58.         cudaDeviceProp prop;
  59.         cudaGetDeviceProperties(&prop, i);
  60.         //打印设备信息
  61.         printDeviceProp(prop);
  62.         //获得显卡的时钟频率
  63.         clockRate = prop.clockRate;
  64.         if (cudaGetDeviceProperties(&prop, i) == cudaSuccess) {
  66.             if (prop.major >= 1) {
  68.                 break;
  69.             }
  70.         }
  71.     }
  72.     if (i == count) {
  74.         fprintf(stderr, "There is no device supporting CUDA 1.x.\n");
  75.         return false;
  76.     }
  77.     cudaSetDevice(i);
  78.     return true;
  79. }
  80. // __global__ 函数 (GPU上执行) 计算立方和
  81. __global__ static void sumOfSquares(int *num, int* result, clock_t* time)
  82. {
  84.     //声明一块共享内存
  85.     extern __shared__ int shared[];
  86.     //表示目前的 thread 是第几个 thread（由 0 开始计算）
  87.     const int tid = threadIdx.x;
  88.     //表示目前的 thread 属于第几个 block（由 0 开始计算）
  89.     const int bid = blockIdx.x;
  90.     shared[tid] = 0;
  91.     int i;
  92.     //记录运算开始的时间
  93.     clock_t start;
  94.     //只在 thread 0（即 threadIdx.x = 0 的时候）进行记录，每个 block 都会记录开始时间及结束时间
  95.     if (tid == 0) time[bid] = clock();
  96.     //thread需要同时通过tid和bid来确定，同时不要忘记保证内存连续性
  97.     for (i = bid * THREAD_NUM + tid; i < DATA_SIZE; i += BLOCK_NUM * THREAD_NUM) {
  99.         shared[tid] += num[i] * num[i] * num[i];
  100.     }
  101.     //同步 保证每个 thread 都已经把结果写到 shared[tid] 里面
  102.     __syncthreads();
  103.     //树状加法
  104.     int offset = 1, mask = 1;
  105.     while (offset < THREAD_NUM)
  106.     {
  108.         if ((tid & mask) == 0)
  109.         {
  111.             shared[tid] += shared[tid + offset];
  112.         }
  113.         offset += offset;
  114.         mask = offset + mask;
  115.         __syncthreads();
  116.     }
  117.     //计算时间,记录结果，只在 thread 0（即 threadIdx.x = 0 的时候）进行，每个 block 都会记录开始时间及结束时间
  118.     if (tid == 0)
  119.     {
  121.         result[bid] = shared[0];
  122.         time[bid + BLOCK_NUM] = clock();
  123.     }
  124. }
  125. int main()
  126. {
  128.     //CUDA 初始化
  129.     if (!InitCUDA()) {
  131.         return 0;
  132.     }
  133.     //生成随机数
  134.     GenerateNumbers(data, DATA_SIZE);
  135.     /*把数据复制到显卡内存中*/
  136.     int* gpudata, *result;
  137.     clock_t* time;
  138.     //cudaMalloc 取得一块显卡内存 ( 其中result用来存储计算结果，time用来存储运行时间 )
  139.     cudaMalloc((void**)&gpudata, sizeof(int)* DATA_SIZE);
  140.     cudaMalloc((void**)&result, sizeof(int)* BLOCK_NUM);
  141.     cudaMalloc((void**)&time, sizeof(clock_t)* BLOCK_NUM * 2);
  142.     //cudaMemcpy 将产生的随机数复制到显卡内存中
  143.     //cudaMemcpyHostToDevice - 从内存复制到显卡内存
  144.     //cudaMemcpyDeviceToHost - 从显卡内存复制到内存
  145.     cudaMemcpy(gpudata, data, sizeof(int)* DATA_SIZE, cudaMemcpyHostToDevice);
  146.     // 在CUDA 中执行函数 语法：函数名称<<<block 数目, thread 数目, shared memory 大小>>>(参数...);
  147.     sumOfSquares << < BLOCK_NUM, THREAD_NUM, THREAD_NUM * sizeof(int) >> >(gpudata, result, time);
  148.     /*把结果从显示芯片复制回主内存*/
  149.     int sum[BLOCK_NUM];
  150.     clock_t time_use[BLOCK_NUM * 2];
  151.     //cudaMemcpy 将结果从显存中复制回内存
  152.     cudaMemcpy(&sum, result, sizeof(int)*BLOCK_NUM, cudaMemcpyDeviceToHost);
  153.     cudaMemcpy(&time_use, time, sizeof(clock_t)* BLOCK_NUM * 2, cudaMemcpyDeviceToHost);
  154.     //Free
  155.     cudaFree(gpudata);
  156.     cudaFree(result);
  157.     cudaFree(time);
  158.     int final_sum = 0;
  159.     for (int i = 0; i < BLOCK_NUM; i++) {
  161.         final_sum += sum[i];
  162.     }
  163.     //采取新的计时策略 把每个 block 最早的开始时间，和最晚的结束时间相减，取得总运行时间
  164.     clock_t min_start, max_end;
  165.     min_start = time_use[0];
  166.     max_end = time_use[BLOCK_NUM];
  167.     for (int i = 1; i < BLOCK_NUM; i++) {
  169.         if (min_start > time_use[i])
  170.             min_start = time_use[i];
  171.         if (max_end < time_use[i + BLOCK_NUM])
  172.             max_end = time_use[i + BLOCK_NUM];
  173.     }
  174.     printf("GPUsum: %d time_clock: %d time: %f\n", final_sum, max_end - min_start, ((float)(max_end - min_start) / (clockRate * 1000)));
  175.     final_sum = 0;
  176.     for (int i = 0; i < DATA_SIZE; i++) {
  178.         final_sum += data[i] * data[i] * data[i];
  179.     }
  180.     printf("CPUsum: %d \n", final_sum);
  181.     return 0;
  182. }

发表评论

表情：
评论列表 （有 0 条评论，101人围观）

还没有评论，来说两句吧...

相关阅读

    相关 CUDA 并行计算

    CUDA 并行计算 并行计算可以被定义为同时使用许多计算资源 (核心或计算机) 来执行并发计算，一个大的问题可以被分解成多个小问题，然后在不同的计算资源上并行处理这些小

    小咪咪小咪咪/ 0 赞/ 102 阅读