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

cuda编程模型

  1. host和device:
    • host:即CPU,CPU所关联的内存就叫host memory
    • device:即GPU,GPU内的内存就叫device memory
    • 运行CUDA程序主要有三步:1)host-to-device transfer:将数据从host memory拷到device memory;2)加载GPU程序并执行,将数据缓存到on-chip上读取更快;3)device-to-host transfer:将device memory中的结果拷到host memory
  2. 【逻辑上】CUDA kernel:就是在GPU上执行的函数
    • start with a __global__ declaration specifier
    • a kernel is executed as a grid of blocks of threads
  3. 【物理上】一个CUDA block在一个SM(streaming multiprocessor)上执行;一个SM可以运行多个并发的CUDA block;一个kernel在一个deivce上执行,一个device上可以同时运行多个kernels。如下图是逻辑和硬件资源的映射关系:
    请添加图片描述
    • 其中,SM的基本执单元是包含32个线程的线程数,所以block大小一般设置为32的倍数
    • 每个线程都包含自己的指令地址计数器和寄存器状态,也有自己独立的执行路径。所以尽管一个warp中的线程同时从同一程序地址执行,但可能具有不同的行为。比如遇到了分支结构,一些线程可能进入这个分支,但另外一些可能并不用执行,只能进入死等(等待warp内其他未执行完的线程执行完毕),因为GPU规定线程束中所有线程在同一周期执行相同的指令,线程束分化会导致性能下降。
  4. CUDA中的index问题
    • 几个built-in变量:gridDim , blockDim用于指示自己的维度;blockIdx, threadIdx用于指示自己在上一层中是第几个。如下图的例子中:
      • gridDim.x=3(表示grid的x维度有3个blocks),gridDim.y=2(grid的y维度有2个blocks)
      • blockDim.x=4(表示block的x维度有4个threads),blockDim.y=3(表示block的y维度有3个threads)
      • blockIdx.x=0, blockIdx.y=1 表示的是 block(0, 1)
      • threadIdx.x=2, threadIdx.y=1表示的是 thread(2, 1)
        请添加图片描述
    • thread indexing: 为每个thread分配一个唯一的id
      • 1D grid of 1D blocks: threadId = blockIdx.x * blockDim.x + threadIdx.x
        请添加图片描述
      • 1D grid of 2D blocks: threadId = blockIdx.x * blockDim.x * blockDim * y + threadIdx.y * blockDim.x + threadIdx.x;
        请添加图片描述
      • 2D grid of 1D blocks: threadId = blockIdx.y * gridDim.x * blockDim.x + blockIdx.x * blockDim.x + threadIdx.x
        • blockId = blockIdx.y * gridDim.x + blockIdx.x
        • threadId = blockId * blockDim.x + threadIdx.x
          请添加图片描述
      • 2D grid of 2D blocks: threadId = blockIdx.y * gridDim.x * blockDim.x * blockDim.y + blockIdx.x * blockDim.x * blockDim.y + theadIdx.y * blockDim.x + threadIdx.x
        请添加图片描述
      • 总结:二维下,[小Id]就是[小Idx.y]*[大Dim.x]+[小Idx.x],即若位于第n行,前面就有n*上一级列数个,然后再加上是第几列(因为索引值从零开始,所以直接加即可)
  5. GPU存储
    在这里插入图片描述
    在这里插入图片描述
    上图中:
    L1/SMEM:指L1 cache/Shared memory,是每个SM独有的。Shared memory可以由用户写代码进行数据的读写控制,L1则不行;
    Read only:只读缓存;
    L2 Cache:所有SM都可以访问
    Global Memory:全局内存,是所有线程都能访问的内存,也是和CPU内存进行数据传递的地方。通常说的显存就是global memory
    请添加图片描述
    每个SM都拥有自己的shared memory,而这些SM们都位于同一块芯片上(on the same chip),这块芯片通过PCB电路连接内存芯片(DRAM)
    SP(cuda core、流处理器,一个thread占用一个SP)对shared memory的访问属于片上访问,可以立刻获得数据
    SP对内存芯片(DRAM)的访问需要通过请求内存控制器等一系列操作,然后才能得到数据。

    下图是一个速度问题:其中s、t、u是本地内存(local memory)中的变量,a、b、c是shared memory中的变量,所以t=s最快
    请添加图片描述
  6. __syncthreads()
    • 确保这行代码之前,同一个block内的所有线程都完成了各自的工作(如将数据从全局内存加载到了共享内存中)。只同步一个线程块中的线程,其他线程块不受影响

http://www.mrgr.cn/news/51610.html

相关文章:

  • GStreamer 简明教程(七):实现管道的动态数据流
  • 自定义多级联动选择器指南(uni-app)
  • 力扣之1398.购买了产品A和产品B却没有购买产品C顾客
  • 飞牛NAS未识别到网卡
  • 【优选算法】——双指针(下篇)!
  • 有关 C#多表查询学习
  • 一步步优化Redis实现分布式锁
  • 【IRV2】Deepfake video detection using InceptionResnetV2
  • Java基础概览和常用知识(七)
  • F5-TTS开源项目详解:非自回归语音合成技术革新与应用场景
  • 【ROS实操六】launch的使用
  • ThreadLocal的应用场景
  • Linux高阶——1013—正则表达式练习
  • 在JasperReports中自动生成序列号
  • 小马哥飞控硬件学习
  • 10.16Python基础-函数
  • 1、HTML笔记
  • AI周报(10.6-10.12)
  • java-day11
  • 通过华为鲲鹏认证的软件产品如何助力信创产业