cuda编程入门了解

AIchiNiurou

已于 2022-04-04 00:26:46 修改

阅读量950

点赞数 2

分类专栏：机器学习+数学学习工具及问题记录文章标签： c++

于 2021-09-11 16:53:48 首次发布

本文链接：https://blog.csdn.net/weixin_44523062/article/details/120239596

版权

学习工具及问题记录同时被 2 个专栏收录

180 篇文章 2 订阅

订阅专栏

机器学习+数学

41 篇文章 1 订阅

订阅专栏

及好的cuda精简

异构host device 异步并行 cuda加速计算编程简化步骤，按这个逻辑学习代码流程

一般CUDA程序分成下面这些步骤：

分配GPU内存（CUmaloc）
拷贝内存到设备
调用CUDA内核函数来执行计算
把计算完成数据拷贝回主机端
内存销毁
上面的hello world只到第三步，没有内存交换。

cuda软件体系driver（cuda.h），runtime（cudart.so）不可混用，cuda lib（cuBlas，cuFFT, cuDPP）

driver函数前缀CU*，调用时必须初始化 CUcontext ,cuCtxCreate手动管理，堆栈方式push/pop
runtime函数前缀cuda*， cudaGetDeviceCount ,自动管理context，cuDevicePrimaryCtxRetain 以runtime api为基础
api：
CUFFT：傅里叶变换
CUBLAS：基本（线性代数程序库basic linear algebra subprograms）矩阵与向量的运算库
CYDPP：常用的并行操作函数库
CUDNN：利用cuda进行深度卷积神经网络
https://developer.nvidia.com/gpu-accelerated-libraries

每个线程都有一个栈结构存储context，上下文关联对GPU的所有操作
栈顶是当前使用的context，context与一块显卡关联，一个显卡可以被多个context关联
context只是为了方便控制device的一种手段而提出来，，栈的存在是为了方便控制多个设备

大部分api依赖context，例如查询显卡名称、参数、内存分配、释放等，runtime 自己cuInit，并使用cuDevicePrimaryCtxRetain进行第一个context的api调用创建和关联context

核函数

global
调用
kernel_function<<<grid,block,sharemem>>>（）定义多少线程执行
一个kernel所启动的所有线程称为一个网格（grid），同一个网格上的线程共享相同的全局内存空间，grid是线程结构的第一层次，而网格又可以分为很多线程块（block），一个线程块里面包含很多线程，这是第二个层次，如图
dim3 grid(3, 2);
dim3 block(5, 3);
kernel_fun<<< grid, block >>>(prams…);

grid和block都是定义为dim3类型的变量，dim3可以看成是包含三个无符号整数（x，y，z）成员的结构体变量，在定义时，缺省值初始化为1。因此grid和block可以灵活地定义为1-dim，2-dim以及3-dim结构，对于图中结构（主要水平方向为x轴），定义的grid和block如下所示，kernel在调用时也必须通过执行配置<<<grid, block>>>来指定kernel所使用的线程数及结构。
在这里插入图片描述
所以，一个线程需要两个内置的坐标变量（blockIdx，threadIdx）来唯一标识，它们都是dim3类型变量，其中blockIdx指明线程所在grid中的位置，而threaIdx指明线程所在block中的位置，如图中的Thread (1,1)满足：
threadIdx.x = 1
threadIdx.y = 1
blockIdx.x = 1
blockIdx.y = 1

__global__ void test_print_kernel(const float* pdata, int ndata){
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    /*    dims                 indexs
        gridDim.z            blockIdx.z
        gridDim.y            blockIdx.y
        gridDim.x            blockIdx.x
        blockDim.z           threadIdx.z
        blockDim.y           threadIdx.y
        blockDim.x           threadIdx.x

        Pseudo code:
        position = 0
        for i in 6:
            position *= dims[i]
            position += indexs[i]
    */
        printf("Element[%d] = %f, threadIdx.x=%d, blockIdx.x=%d, blockDim.x=%d\n", idx, pdata[idx], threadIdx.x, blockIdx.x, blockDim.x);
}

一个线程块上的线程是放在同一个流式多处理器（SM)上的，但是单个SM的资源有限，这导致线程块中的线程数是有限制的，现代GPUs的线程块可支持的线程数可达1024个。有时候，我们要知道一个线程在blcok中的全局ID，此时就必须还要知道block的组织结构，这是通过线程的内置变量blockDim来获得。它获取线程块各个维度的大小。对于一个2-dim的block [公式] ，线程 [公式] 的ID值为 [公式] ，如果是3-dim的block [公式] ，线程 [公式] 的ID值为 [公式] 。另外线程还有内置变量gridDim，用于获得网格块各个维度的大小。

kernel的这种线程组织结构天然适合vector,matrix等运算，如我们将利用上图2-dim结构实现两个矩阵的加法，每个线程负责处理每个位置的两个元素相加，代码如下所示。线程块大小为(16, 16)，然后将N*N大小的矩阵均分为不同的线程块来执行加法运算。

// Kernel定义
__global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N]) 
{ 
    int i = blockIdx.x * blockDim.x + threadIdx.x; 
    int j = blockIdx.y * blockDim.y + threadIdx.y; 
    if (i < N && j < N) 
        C[i][j] = A[i][j] + B[i][j]; 
}
int main() 
{ 
    ...
    // Kernel 线程配置
    dim3 threadsPerBlock(16, 16); 
    dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);
    // kernel调用
    MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C); 
    ...
}

#include <stdio.h>
#include <cuda_runtime.h>

#include <math.h>
//host 函数调用必须使用__device__
__device__ float sigmoid(float x){
    return 1/(1+exp(-x));
}

__global__ void test_print_kernel(const float* pdata, int ndata){

    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    /*    dims                 indexs
        gridDim.z            blockIdx.z
        gridDim.y            blockIdx.y
        gridDim.x            blockIdx.x
        blockDim.z           threadIdx.z
        blockDim.y           threadIdx.y
        blockDim.x           threadIdx.x

        Pseudo code:
        position = 0
        for i in 6:
            position *= dims[i]
            position += indexs[i]
    */
    float y=sigmoid(0.5f);
    printf("Element[%d] = %f, threadIdx.x=%d, blockIdx.x=%d, blockDim.x=%d\n", idx, pdata[idx], threadIdx.x, blockIdx.x, blockDim.x);
}

在这里插入图片描述

配置cuda编程环境 ubuntu clion nvcc

在file->setting->build->cmake，在cmake options中写入-DCMAKE_CUDA_COMPILER:PATH=/usr/local/cuda/bin/nvcc
在这里插入图片描述

CPU架构和GPU架构异构

在这里插入图片描述
ALU：四核（算术逻辑arithmetic logic unit）
control：控制器
cache：缓存在片上
DRAM：内存总线上

GPU
小绿色块：ALUs，公用一组 control cache 相当于多了许多ALUs的cpu核
控制能力减弱，但是计算能力提升

主机代码he 设备代码
主机代码在主机端运行，被编译成主机架构的机器码，设备端的在设备上执行，被编译成设备架构的机器码，所以主机端的机器码和设备端的机器码是隔离的，自己执行自己的，没办法交换执行。主机端代码主要是控制设备，完成数据传输等控制类工作，设备端主要的任务就是计算。因为当没有GPU的时候CPU也能完成这些计算，只是速度会慢很多，所以可以把GPU看成CPU的一个加速设备。NVIDIA目前的计算平台（不是架构）有：Tegra 嵌入式
Geforce 打游戏
Quadro
Tesla 计算

CUDA核心数，内存内存带宽
nvidia自己有一套描述GPU计算能力的代码，其名字就是“计算能力”，主要区分不同的架构，早其架构的计算能力不一定比新架构的计算能力强

计算能力架构名
1.x Tesla
2.x Fermi
3.x Kepler
4.x Maxwell
5.x Pascal
6.x Volta

低并行、逻辑复杂的程序适合用CPU
高并行、逻辑简单的大数据计算适合GPU

CPU和GPU线程的区别：CUDA：一种异构计算平台

CPU线程是重量级实体，操作系统交替执行线程，线程上下文切换花销很大
GPU线程是轻量级的，GPU应用一般包含成千上万的线程，多数在排队状态，线程之间切换基本没有开销。
CPU的核被设计用来尽可能减少一个或两个线程运行时间的延迟，而GPU核则是大量线程，最大幅度提高吞吐量

CUDA平台不是单单指软件或者硬件，而是建立在Nvidia GPU上的一整套平台，并扩展出多语言支持
在这里插入图片描述
CUDA C 是标准ANSI C语言的扩展，扩展出一些语法和关键字来编写设备端代码，而且CUDA库本身提供了大量API来操作设备完成计算。

在这里插入图片描述
驱动API是低级的API，使用相对困难，运行时API是高级API使用简单，其实现基于驱动API。
这两种API是互斥的，也就是你只能用一个，两者之间的函数不可以混合调用，只能用其中的一个库。

一个CUDA应用通常可以分解为两部分，

CPU 主机端代码
GPU 设备端代码
CUDA nvcc编译器会自动分离你代码里面的不同部分，如图中主机代码用C写成，使用本地的C语言编译器编译，设备端代码，也就是核函数，用CUDA C编写，通过nvcc编译，链接阶段，在内核程序调用或者明显的GPU设备操作时，添加运行时库。

在这里插入图片描述

#include<stdio.h>
__global__ void hello_world(void)
{
  printf("GPU: Hello world!\n");
}
int main(int argc,char **argv)
{
  printf("CPU: Hello world!\n");
  hello_world<<<1,10>>>();
  cudaDeviceReset();//if no this line ,it can not output hello world from gpu
  return 0;
}

https://github.com/Tony-Tan/CUDA_Freshman

关键词

__global__
是告诉编译器这个是个可以在设备上执行的核函数

hello_world<<<1,10>>>();  //核函数 griddim, blockdim, sharemem
这句话C语言中没有’<<<>>>’是对设备进行配置的参数，也是CUDA扩展出来的部分。

cudaDeviceReset();
这句话如果没有，则不能正常的运行，因为这句话包含了隐式同步，GPU和CPU执行程序是**异步**的，核函数调用后成立刻会到主机线程继续，而不管GPU端核函数是否执行完毕，所以上面的程序就是GPU刚开始执行，CPU已经退出程序了，所以我们要等GPU执行完了，再退出主机线程。

cuda难吗

CPU与GPU的编程主要区别在于对GPU架构的熟悉程度，理解机器的结构是对编程效率影响非常大的一部分，了解你的机器，才能写出更优美的代码，而目前计算设备的架构决定了局部性将会严重影响效率。
数据局部性分两种

空间局部性
时间局部性
这个两个性质告诉我们，当一个数据被使用，其附近的数据将会很快被使用，当一个数据刚被使用，则随着时间继续其被再次使用的可能性降低，数据可能被重复使用。

CUDA中有两个模型是决定性能的：
内存层次结构
线程层次结构
CUDA C写核函数的时候我们只写一小段串行代码，但是这段代码被成千上万的线程执行，所有线程执行的代码都是相同的，CUDA编程模型提供了一个层次化的组织线程，直接影响GPU上的执行顺序。

CUDA抽象了硬件实现：

线程组的层次结构
内存的层次结构
障碍同步
这些都是我们后面要研究的，线程，内存是主要研究的对象，我们能用到的工具相当丰富，NVIDIA为我们提供了：
Nvidia Nsight集成开发环境
CUDA-GDB 命令行调试器
性能分析可视化工具
CUDA-MEMCHECK工具
GPU设备管理工具

AIchiNiurou

关注

2
点赞
踩
2

收藏

觉得还不错? 一键收藏
1
评论
cuda编程入门了解

知乎推荐的学习资料入门cudahttps://face2ai.com/CUDA-F-1-1-%E5%BC%82%E6%9E%84%E8%AE%A1%E7%AE%97-CUDA/CPU架构和GPU架构异构ALU：四核（算术逻辑arithmetic logic unit）control：控制器cache：缓存在片上DRAM：内存总线上GPU小绿色块：ALUs，公用一组 control cache 相当于多了许多ALUs的cpu核控制能力减弱，但是计算能力提升主机代码he 设备代
复制链接

扫一扫