CUDA入门

CUDA编程可以使程序能在GPU上跑，只需要有一台带有GeForce显卡的笔记本就能跑起来并行程序，加快程序运行速度。使用显卡的好处在于功耗和成本低，但性能好。就是GPU比CPU会跑的快！快！快！

先是安装VS和CUDA，这里据说要先装VS再装CUDA才行，反正窝是装自闭了qwq，最后给电脑重装系统才好。。心痛

然后介绍一下下如何在VS中给GPU变编程（留坑）

程序遵循以下流程：

主机端准备数据 -> 数据复制到GPU内存中 -> GPU执行核函数 -> 数据由GPU取回到主机

这里拿经典的矢量相加程序作为sample （代码注释是从另一个博主那里copy的 后面标注惹

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

#include <stdio.h>

// 接口函数： 主机代码调用GPU设备实现矢量加法 c = a + b
cudaError_t addWithCuda(int *c, const int *a, const int *b, unsigned int size);

// 核函数：每个线程负责一个分量的加法
__global__ void addKernel(int *c, const int *a, const int *b)
{
    int i = threadIdx.x; // 获取线程ID
    c[i] = a[i] + b[i];
}

int main()
{
    const int arraySize = 5;
    const int a[arraySize] = { 1, 2, 3, 4, 5 };
    const int b[arraySize] = { 10, 20, 30, 40, 50 };
    int c[arraySize] = { 0 };

    // 并行矢量相加
    cudaError_t cudaStatus = addWithCuda(c, a, b, arraySize);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "addWithCuda failed!");
        return 1;
    }

    printf("{1,2,3,4,5} + {10,20,30,40,50} = {%d,%d,%d,%d,%d}\n",
        c[0], c[1], c[2], c[3], c[4]);

    // CUDA设备重置，以便其它性能检测和跟踪工具的运行，如Nsight and Visual Profiler to show complete traces.traces.
    cudaStatus = cudaDeviceReset();
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaDeviceReset failed!");
        return 1;
    }

    return 0;
}

// 接口函数实现： 主机代码调用GPU设备实现矢量加法 c = a + b
cudaError_t addWithCuda(int *c, const int *a, const int *b, unsigned int size)
{
    int *dev_a = 0;
    int *dev_b = 0;
    int *dev_c = 0;
    cudaError_t cudaStatus;

    // 选择程序运行在哪块GPU上，(多GPU机器可以选择)
    cudaStatus = cudaSetDevice(0);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");
        goto Error;
    }

    // 依次为 c = a + b三个矢量在GPU上开辟内存 .
    cudaStatus = cudaMalloc((void**)&dev_c, size * sizeof(int));
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMalloc failed!");
        goto Error;
    }

    cudaStatus = cudaMalloc((void**)&dev_a, size * sizeof(int));
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMalloc failed!");
        goto Error;
    }

    cudaStatus = cudaMalloc((void**)&dev_b, size * sizeof(int));
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMalloc failed!");
        goto Error;
    }

    // 将矢量a和b依次copy进入GPU内存中
    cudaStatus = cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMemcpy failed!");
        goto Error;
    }

    cudaStatus = cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMemcpy failed!");
        goto Error;
    }

    // 运行核函数，运行设置为1个block，每个block中size个线程
    addKernel<<<1, size>>>(dev_c, dev_a, dev_b);

    // 检查是否出现了错误
    cudaStatus = cudaGetLastError();
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "addKernel launch failed: %s\n", cudaGetErrorString(cudaStatus));
        goto Error;
    }

    // 停止CPU端线程的执行，直到GPU完成之前CUDA的任务，包括kernel函数、数据拷贝等
    cudaStatus = cudaDeviceSynchronize();
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching addKernel!\n", cudaStatus);
        goto Error;
    }

    // 将计算结果从GPU复制到主机内存
    cudaStatus = cudaMemcpy(c, dev_c, size * sizeof(int), cudaMemcpyDeviceToHost);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaMemcpy failed!");
        goto Error;
    }

Error:
    cudaFree(dev_c);
    cudaFree(dev_a);
    cudaFree(dev_b);

    return cudaStatus;
}
--------------------- 
作者：shuzfan 
来源：CSDN 
原文：https://blog.csdn.net/shuzfan/article/details/76650947 
版权声明：本文为博主原创文章，转载请附上博文链接！

CUDA和C相比，多了几个头文件：

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

还多了一个神奇的函数，前面有__global__修饰：

__global__ void addKernel(int *c, const int *a, const int *b)
{
    int i = threadIdx.x; // 获取线程ID
    c[i] = a[i] + b[i];
}

这个函数就是在GPU上运行的函数，称为核函数(Kernel Function)，__global__为CUDA C为标准C增加的修饰符，表示该函数将会编译在GPU上而不是CPU上，所以这个函数将被交给编译设备代码的编译器（NVCC）来处理，而main函数则依旧交给主机编译器（CPU(VS)）。

在后面有一个cudaError_t，它是一个枚举类型，可以作为几乎所有CUDA函数的返回类型，用来检测错误，返回0代表成功。

其实，CUDA就是通过直接提供API接口或者在语言层面集成一些新的东西来实现在主机代码中调用设备代码。

再介绍一下GPU线程的层次结构：

GPU线程的层次结构

线程（thread）是并行运算结构中的最小单位，thread可以以一维、二维、三维的形式组织在一起。线程（thread）组成线程块（block），线程块（block）组成线程块网络（grid）

回到程序中的addKernel函数上来，这个函数会被GPU上的多个线程同时执行一次，线程间彼此没有通信，相互独立。到底会有多少个线程来分别执行核函数，是在“<<<>>>”符号里定义的。“<<<>>>”表示运行时配置符号，在本程序中的定义是<<<1，size>>>，表示分配了一个线程块（Block），每个线程块有分配了size个线程，“<<<>>>”中的 参数并不是传递给设备代码的参数，而是定义主机代码运行时如何启动设备代码。以上定义的这些线程都是一个维度上的，可以通过thredaIdx.x来获取执行当前计算任务的线程的ID号。

<<<>>>中的参数也可能是四个，第一个表示线程块数、第二个表示每个线程块中线程数，第三个表示每个block用到的共享内存大小，第四个为流对象，表示当前核函数在哪个流上运行。

 

 

参考：https://blog.csdn.net/dcrmg/article/details/54446393

https://blog.csdn.net/zjc156m/article/details/79015742

https://blog.csdn.net/shuzfan/article/details/76602565