• 我的NVIDIA开发者之旅——CUDA编程基础——并行矩阵乘法


    一个相对简单的CUDA编程Demo(Hello World):

    1.CUDA编程首先呢是分配thread以及block

    1. #include<stdio.h>
    2. #include<time.h>
    3. #include<cuda_runtime.h> //cuda运行时间接口
    4. #define Thread_Num 256 //每一block包含的线程数
    5. #define Matrix_Size 10
    6. const int block_num=(Matrix_Size+Thread_Num-1)/Thread_Num;

    2.然后是两个基本的函数:

    • 打印设备信息

    1. void printDeviceProp(const cudaDeviceProp &prop)
    2. {
    3. printf("Device Name : %s.\n", prop.name);
    4. printf("totalGlobalMem : %d.\n", prop.totalGlobalMem);
    5. printf("sharedMemPerBlock : %d.\n", prop.sharedMemPerBlock);
    6. printf("regsPerBlock : %d.\n", prop.regsPerBlock);
    7. printf("warpSize : %d.\n", prop.warpSize);
    8. printf("memPitch : %d.\n", prop.memPitch);
    9. printf("maxThreadsPerBlock : %d.\n", prop.maxThreadsPerBlock);
    10. printf("maxThreadsDim[0 - 2] : %d %d %d.\n", prop.maxThreadsDim[0], prop.maxThreadsDim[1], prop.maxThreadsDim[2]);
    11. printf("maxGridSize[0 - 2] : %d %d %d.\n", prop.maxGridSize[0], prop.maxGridSize[1], prop.maxGridSize[2]);
    12. printf("totalConstMem : %d.\n", prop.totalConstMem);
    13. printf("major.minor : %d.%d.\n", prop.major, prop.minor);
    14. printf("clockRate : %d.\n", prop.clockRate);
    15. printf("textureAlignment : %d.\n", prop.textureAlignment);
    16. printf("deviceOverlap : %d.\n", prop.deviceOverlap);
    17. printf("multiProcessorCount : %d.\n", prop.multiProcessorCount);
    18. }
    • 初始化cuda

    1. bool InitCUDA()
    2. {
    3. int count;
    4. cudaGetDeviceCount(&count);
    5. if(count==0){
    6. fprintf(stderr,"three is no device.\n");
    7. return false;
    8. }
    9. int i;
    10. for(i=0;i<count;i++)
    11. {
    12. cudaDeviceProp prop;
    13. cudaGetDeviceProperties(&prop,i);
    14. printDeviceProp(prop);
    15. if(cudaGetDeviceProperties(&prop,i)==cudaSuccess){
    16. if(prop.major>=1){break;}
    17. }
    18. }
    19. if (i == count) {
    20. fprintf(stderr, "There is no device supporting CUDA 1.x.\n");
    21. return false;
    22. }
    23. cudaSetDevice(i);
    24. return true;
    25. }

    3.接着随机生成两个矩阵

    1. void matGen(float* a, int n)
    2. {
    3. int i,j;
    4. for(i=0;i<n;i++)
    5. {
    6. for(j=0;j<n;j++)
    7. {
    8. a[i*n+j]=(float)rand()/(float)RAND_MAX+1.00;
    9. }
    10. }
    11. }

    4.并行矩阵乘法函数,最主要的一部分

    1. __global__ static void matMultCuda(const float* a,const float* b,float* c,int n,clock_t* time)
    2. {
    3. const int tid=threadIdx.x;
    4. const int bid=blockIdx.x;
    5. //从 bid 和 tid 计算出这个 thread 应该计算的 row 和 column
    6. const int idx = bid * Thread_Num + tid;
    7. const int row = idx / n;
    8. const int column = idx % n;
    9. int i;
    10. //clock_t start;
    11. //每个block开始时记录
    12. if(tid==0) time[bid]=clock();
    13. //计算矩阵乘法
    14. if(row < n && column < n)
    15. {
    16. float t=0;
    17. for(i=0;i<n;i++)
    18. {
    19. t=t+a[row*n+i]+a[i*n+column];
    20. }
    21. c[row*n+column]=t;
    22. }
    23. //计算时间,记录结果,只在 thread 0(即 threadIdx.x = 0 的时候)进行,每个 block 都会记录开始时间及结束时间
    24. if (tid == 0) time[bid + block_num] = clock();
    25. }

    5.运算完后打印出矩阵

    1. void printMatrix(const float *A, const int n) {
    2. for(int i = 0; i < n; i++){
    3. for(int j = 0; j < n; j++){
    4. printf("%.2f" ,A[i*n+j]);
    5. printf(" ");
    6. }
    7. printf("\n");
    8. }
    9. printf("\n");
    10. }

    6.最后我们来看一下主函数

    1. int main()
    2. {
    3. if(!InitCUDA()) return 0;
    4. float *a,*b,*c;
    5. int n=Matrix_Size;
    6. //分配内存
    7. a=(float*)malloc(sizeof(float)*n*n);
    8. b=(float*)malloc(sizeof(float)* n*n);
    9. c=(float*)malloc(sizeof(float)* n*n);
    10. //设置随机种子
    11. srand(0);
    12. //随机生成两个矩阵
    13. matGen(a,n);
    14. matGen(b,n);
    15. float *cuda_a,*cuda_b,*cuda_c;
    16. clock_t* time;
    17. //cudaMalloc 获取一块显卡内存
    18. cudaMalloc((void**)&cuda_a, sizeof(float)* n*n);
    19. cudaMalloc((void**)&cuda_b, sizeof(float)* n*n);
    20. cudaMalloc((void**)&cuda_c, sizeof(float)* n*n);
    21. cudaMalloc((void**)&time, sizeof(clock_t)* block_num*2);
    22. //cudaMemcpy 将产生的矩阵复制到显卡内存中
    23. //cudaMemcpyHostToDevice - 从内存复制到显卡内存
    24. //cudaMemcpyDeviceToHost - 从显卡内存复制到内存
    25. cudaMemcpy(a,cuda_a,sizeof(float)* n*n,cudaMemcpyHostToDevice);
    26. cudaMemcpy(b,cuda_b,sizeof(float)* n*n,cudaMemcpyHostToDevice);
    27. //printMatrix(cuda_a, n);
    28. //printMatrix(cuda_b, n);
    29. // 在CUDA 中执行函数 语法:函数名称<<<block 数目, thread 数目, shared memory 大小>>>(参数...);
    30. matMultCuda<<<block_num, Thread_Num, 0>>>(cuda_a,cuda_b,cuda_c,n,time);
    31. //把结果复制回内存中
    32. clock_t time_use[block_num*2];
    33. cudaMemcpy(c,cuda_c,sizeof(float)* n*n,cudaMemcpyDeviceToHost);
    34. cudaMemcpy(&time_use, time, sizeof(clock_t)* block_num * 2, cudaMemcpyDeviceToHost);
    35. printMatrix(a, n);
    36. printMatrix(b, n);
    37. printMatrix(c, n);
    38. //释放资源
    39. cudaFree(cuda_a);
    40. cudaFree(cuda_b);
    41. cudaFree(cuda_c);
    42. cudaFree(time);
    43. //把每个 block 最早的开始时间,和最晚的结束时间相减,取得总运行时间
    44. clock_t min_start, max_end;
    45. min_start = time_use[0];
    46. max_end = time_use[block_num];
    47. for (int i = 1; i < block_num; i++)
    48. {
    49. if (min_start > time_use[i]) min_start = time_use[i];
    50. if (max_end < time_use[i + block_num]) max_end = time_use[i + block_num];
    51. }
    52. //核函数运行时间
    53. clock_t final_time = max_end - min_start;
    54. printf("gputime: %d\n", final_time);
    55. return 0;
    56. }

    最后感谢大家的观看啦,一起学习一起进步鸭! 

  • 相关阅读:
    图解哈希表及其原理
    orb-slam2 从单目开始的简单学习(8):LocalMapping
    【云原生-k8s】Linux服务器搭建单机版kubernetes服务
    3分钟带你了解前端缓存-HTTP缓存
    Logstash同步MySQL数据到ElasticSearch
    android全局捕获异常错误
    Http代理与socks5代理有何区别?如何选择?(二)
    NNDL 作业11:优化算法比较
    初学者,应该如何学好 C/C++语言?
    this与super
  • 原文地址:https://blog.csdn.net/qq_53904578/article/details/125549127