cuda 编程记录(1)

本篇介绍

从硬件层面有2种思路可以提升性能,一种是提供更强的计算单元,缓存,分支预测,这样可以尽可能降低指令的执行耗时,也就是面向低延时的优化,代表方向就是intel的CPU,一种是提供更多的硬件线程,也就是计算单元,同时提高访问内存的带宽,也就是提高更高的吞吐量,也就是面向高吞吐的优化,代表方向是nvidia的GPU。对比效果如下:


10d260618f1976ece201668f642543f0.png

对于软件,利用CPU和GPU组合来提升性能就成了一个自然的思路。软件一般是由串行和可并行部分构成的,串行部分的优化可以借助CPU实现,可并行部分就可以利用GPU,按照amudahl定理就可以评估出优化可并行部分带来的性能提升。
cuda就是一种可以由CPU分配GPU并行的计算统一设备架构(compute unified device architecture) 。利用cuda 就可以让可并行部分在GPU上执行。

这时候可能会有一个疑问,CPU也支持并行,比如一核多线程,也有SIMD,那和GPU有什么区别呢?GPU的思路也是让线程尽可能多,而且会非常多,比如CPU高配几百个线程,GPU起步就几千几万个线程。至于SIMD,GPU是SPMD,GPU不要求同一时刻所有线程执行同样的指令,只要求所有线程处理同样的数据,这点会在cuda 编程中体会到。
本系列接下来部分就会详细介绍cuda部分。

cuda 示例

官方API上可以看到cuda的相关API,本次demo 只涉及到了基础的内存管理。

内存分配

内存分配

在cuda中,host指的是cpu侧,device是gpu侧,为了区分cpu或gpu侧的代码,需要有关键字标识,如下所示:


b5825bb6-e2fb-4b19-b1ea-eb5d0296f094.png

在cuda运行时,系统会启动一个线程网格,网格内部有多个线程块,每个线程块包含的线程是一样的。内置变量blockIdx标识线程块索引,是三维的,blockDim 是线程内部的组织结构,也是三维的。threadIdx是线程索引,也是三维的。

在调用核函数时就可以设置网格和线程块维度。

有如上信息后,如下的cuda代码就比较清楚了:

#include <iostream>
#include <cmath>

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

__global__ void vecAddKernel(float *a, float *b, float *c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        c[i] = a[i] + b[i];
    }
}

void vecAdd(float *a, float *b, float *c, int n) {
    float *a_d,*b_d,*c_d;
    int size = n * sizeof(float);
    cudaError_t err =  cudaMalloc((void **)(&a_d), size);
    if (err != cudaSuccess) {
        std::cout << "cudaMalloc failed: " << cudaGetErrorString(err) << std::endl;
        return;
    }
    cudaMalloc((void **)(&b_d), size);
    cudaMalloc((void **)(&c_d), size);

    cudaMemcpy(a_d, a, size, cudaMemcpyHostToDevice);
    cudaMemcpy(b_d, b, size, cudaMemcpyHostToDevice);
    
    vecAddKernel<<<ceil(n/256.0), 256>>>(a_d, b_d, c_d, n);
    cudaMemcpy(c, c_d, size, cudaMemcpyDeviceToHost);
    
    cudaFree(a_d);
    cudaFree(b_d);
    cudaFree(c_d);
}

针对如上代码添加主函数调用,用nvcc 编译就可以得到可执行文件。nvcc 内部会区分cpu和gpu代码,对于gpu代码会编译成ptx, 最终nvcc也会将ptx汇编后打包到可执行文件中,如下所示:


nvcc 编译

本篇总结

本篇甚至该系列都是参考kirk 的<<Programing Massively parallel processors>>,这本书对GPU的并行编程介绍的很清晰务实。

©著作权归作者所有,转载或内容合作请联系作者
【社区内容提示】社区部分内容疑似由AI辅助生成,浏览时请结合常识与多方信息审慎甄别。
平台声明:文章内容(如有图片或视频亦包括在内)由作者上传并发布,文章内容仅代表作者本人观点,简书系信息发布平台,仅提供信息存储服务。

相关阅读更多精彩内容

友情链接更多精彩内容