当前位置: 首页 > news >正文

深度学习部署(十六): CUDA RunTime API _vector-add 使用cuda核函数实现向量加法

1. 知识点

  1. nthreads的取值,不能大于block能取值的最大值。一般可以直接给512、256,性能就是比较不错的
    • (input_size + block_size - 1) / block_size;是向上取整
  2. 对于一维数组时,采用只定义layout的x维度,若处理的是二维,则可以考虑定义x、y维度,例如处理的是图像
  3. 关于把数据视作一维时,索引的计算
    • 以下是通用的计算公式
    Pseudo code:
    position = 0
    for i in range(6):position *= dims[i]position += indexs[i]
    
    • 例如当只使用x维度时,实际上dims = [1, 1, gd, 1, 1, bd],indexs = [0, 0, bi, 0, 0, ti]
      • 因为0和1的存在,上面的循环则可以简化为:idx = threadIdx.x + blockIdx.x * blockDim.x
      • 即:idx = ti + bi * bd

2. main.cpp文件

#include <cuda_runtime.h>
#include <stdio.h>#define checkRuntime(op)  __check_cuda_runtime((op), #op, __FILE__, __LINE__)bool __check_cuda_runtime(cudaError_t code, const char* op, const char* file, int line){if(code != cudaSuccess){    const char* err_name = cudaGetErrorName(code);    const char* err_message = cudaGetErrorString(code);  printf("runtime error %s:%d  %s failed. \n  code = %s, message = %s\n", file, line, op, err_name, err_message);   return false;}return true;
}void vector_add(const float* a, const float* b, float* c, int ndata);int main(){const int size = 3;float vector_a[size] = {2, 3, 2};float vector_b[size] = {5, 3, 3};float vector_c[size] = {0};float* vector_a_device = nullptr;float* vector_b_device = nullptr;float* vector_c_device = nullptr;checkRuntime(cudaMalloc(&vector_a_device, size * sizeof(float)));checkRuntime(cudaMalloc(&vector_b_device, size * sizeof(float)));checkRuntime(cudaMalloc(&vector_c_device, size * sizeof(float)));checkRuntime(cudaMemcpy(vector_a_device, vector_a, size * sizeof(float), cudaMemcpyHostToDevice));checkRuntime(cudaMemcpy(vector_b_device, vector_b, size * sizeof(float), cudaMemcpyHostToDevice));vector_add(vector_a_device, vector_b_device, vector_c_device, size);checkRuntime(cudaMemcpy(vector_c, vector_c_device, size * sizeof(float), cudaMemcpyDeviceToHost));for(int i = 0; i < size; ++i){printf("vector_c[%d] = %f\n", i, vector_c[i]);}checkRuntime(cudaFree(vector_a_device));checkRuntime(cudaFree(vector_b_device));checkRuntime(cudaFree(vector_c_device));return 0;
}

先定义三个数组: a, b, c 再用cudaMalloc()在GPU上开辟三个内存,在GPU上让a + b 并且让结果存储进c上,再把c的内存从GPU上放到Host上输出

3. 案例.cu文件

#include <stdio.h>
#include <cuda_runtime.h>__global__ void vector_add_kernel(const float* a, const float* b, float* c, int ndata){int idx = threadIdx.x + blockIdx.x * blockDim.x;if(idx >= ndata) return;/*    dims                 indexsgridDim.z            blockIdx.zgridDim.y            blockIdx.ygridDim.x            blockIdx.xblockDim.z           threadIdx.zblockDim.y           threadIdx.yblockDim.x           threadIdx.xPseudo code:position = 0for i in 6:position *= dims[i]position += indexs[i]*/c[idx] = a[idx] + b[idx];
}void vector_add(const float* a, const float* b, float* c, int ndata){const int nthreads = 512;int block_size = ndata < nthreads ? ndata : nthreads;  // 如果ndata < nthreads 那block_size = ndata就够了int grid_size = (ndata + block_size - 1) / block_size; // 其含义是我需要多少个blocks可以处理完所有的任务printf("block_size = %d, grid_size = %d\n", block_size, grid_size);vector_add_kernel<<<grid_size, block_size, 0, nullptr>>>(a, b, c, ndata);// 在核函数执行结束后,通过cudaPeekAtLastError获取得到的代码,来知道是否出现错误// cudaPeekAtLastError和cudaGetLastError都可以获取得到错误代码// cudaGetLastError是获取错误代码并清除掉,也就是再一次执行cudaGetLastError获取的会是success// 而cudaPeekAtLastError是获取当前错误,但是再一次执行cudaPeekAtLastError或者cudaGetLastErro拿到的还是那个错cudaError_t code = cudaPeekAtLastError();if(code != cudaSuccess){    const char* err_name    = cudaGetErrorName(code);    const char* err_message = cudaGetErrorString(code);  printf("kernel error %s:%d  test_print_kernel failed. \n  code = %s, message = %s\n", __FILE__, __LINE__, err_name, err_message);   }
}

两个注意的点

  1. 像这个案例他就三个数相加,其实启动三个线程就足够了,但是一般block给的是512, 256,所以要设定一下,如果数组的长度小于256/512, 就直接用数组的长度的线程数就好。这里就是3个线程

  2. 如果线程索引大于了数组的长度就直接返回了,不然就访问了不知道在哪里的内存了

http://www.lryc.cn/news/39690.html

相关文章:

  • 堆结构的两个应用
  • Java中的 static
  • 基于Vision Transformer的图像去雾算法研究与实现(附源码)
  • 服务器相关常用的命令
  • 今天是国际数学日,既是爱因斯坦的生日又是霍金的忌日
  • Qt Quick - StackLayout 堆布局
  • C/C++网络编程笔记Socket
  • RK3568平台开发系列讲解(网络篇)什么是Socket套接字
  • 2022年全国职业院校技能大赛(中职组)网络安全竞赛试题——渗透测试解析(详细)
  • 尚融宝03-mybatis-plus基本CRUD和常用注解
  • vue多行显示文字展开
  • SpringBoot:SpringBoot 的底层运行原理解析
  • 哪些场景会产生OOM?怎么解决?
  • 金三银四、金九银十 面试宝典 Spring、MyBatis、SpringMVC面试题 超级无敌全的面试题汇总(超万字的面试题,让你的SSM框架无可挑剔)
  • JAVA开发(Spring框架详解)
  • 自学大数据第八天~HDFS命令(二)
  • 贪心算法(几种常规样例)
  • 【数据结构】基础知识总结
  • 宣布推出 .NET 社区工具包 8.1!
  • ChatGPT解开了我一直以来对自动化测试的疑惑
  • 十大经典排序算法(上)
  • 如何从 MySQL 读取 100w 数据进行处理
  • 【数据降维-第2篇】核主成分分析(KPCA)快速理解,及MATLAB实现
  • Python+ChatGPT实战之进行游戏运营数据分析
  • Java每日一练(20230313)
  • 国内ChatGPT日趋成熟后,可以优先解决的几个日常小问题
  • 业内人士真心话,软件测试是没有前途的,我慌了......
  • 哈佛与冯诺依曼结构
  • 传输安全HTTPS
  • Docker--(六)--Docker资源限制