CUDA编程之入门

CUDA代码中的函数

• CUDA程序文件后缀为.cu。一个.cu文件内既包含CPU程序（称为主机程序），也包含GPU程序（称为设备程序）。凡是挂有“__global__”或者“__device__”前缀的函数，都是在GPU上运行的设备程序，不同的是__global__设备程序可被主机程序调用，而__device__设备程序则只能被设备程序调用。没有挂任何前缀的函数，都是主机程序。主机程序显示声明可以用__host__前缀。

• 主机程序主要完成设备环境初始化，数据传输等必备过程，设备程序只负责计算。

• 设备程序需要由NVCC进行编译，而主机程序只需要由主机编译器

• 主机程序中，有一些“cuda”开头的函数，这些都是CUDA Runtime API，即运行时函数，主要负责完成设备的初始化、内存分配、内存拷贝等任务。

cudaMalloc()分配内存

• 行为与C语言中的 malloc()，用 cudaFree()释放内存，与free()相似。

• 主机指针只能访问主机代码中的内存，设备指针只能访问设备代码中的内存。

• 想要访问设备内存可以有两种常见方式：
  （1）在设备代码中使用设备指针；
  （2）调用cudaMemcpy()

cudaMemcpy()传输数据

• 内存复制，第一个参数是目的指针，第二个参数是源指针，第三个参数是复制内存的大小，第四个参数告诉运行时源指针，这是一个什么类型的指针，即把内存从哪里复制到哪里。第四个参数可以选用以下的形式：

（1）cudaMemcpyHostToDevice：从主机复制到设备；

（2）cudaMemcpyDeviceToHost：从设备复制到主机；

（3）cudaMemcpyDeviceToDevice：从设备复制到设备；

（4）cudaMemcpyHostToHost：从主机复制到主机；

（从主机复制到主机直接用标准C的memcpy()）

该代码在设备上进行操作时，反复用 cudaError_t 判断是否有异常，有则打印相应的异常情况，并释放内存退出设备代码

• 代码主要运行步骤：

（1）选择了一个GPU设备，然后生成设备所需内存，再用cudaMemcpy把主机的内存复制到设备内存中。

（2）addKernel<<<1, size>>>(dev_c, dev_a, dev_b)：开启GPU中一个核，各使用一个线程为每个元素进行加法运算，一共开了5个GPU线程。

（3）cudaDeviceSynchronize()：等待核完成运算。

（4）把计算后得到的设备内存复制回主机内存。

（5）释放内存。

（6）最后使用cudaDeviceReset()以便配置和跟踪工具显示完整的跟踪。

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

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;
    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 };

// Add vectors in parallel.
    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]);

// cudaDeviceReset must be called before exiting in order for profiling and tracing tools such as Nsight and Visual Profiler to show complete traces.
    cudaStatus = cudaDeviceReset();
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaDeviceReset failed!");
        return 1;
    }
    return 0;
}

// Helper function for using CUDA to add vectors in parallel.
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;

// Choose which GPU to run on, change this on a multi-GPU system.
    cudaStatus = cudaSetDevice(0);
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaSetDevice failed!  Do you have a CUDA-capable GPU installed?");
        goto Error;
    }

 // Allocate GPU buffers for three vectors (two input, one output) .
    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;
    }

// Copy input vectors from host memory to GPU buffers.
    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;
    }

// Launch a kernel on the GPU with one thread for each element.
    addKernel<<<1, size>>>(dev_c, dev_a, dev_b);

// Check for any errors launching the kernel
    cudaStatus = cudaGetLastError();
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "addKernel launch failed: %s\n", cudaGetErrorString(cudaStatus));
        goto Error;
    }
    
// cudaDeviceSynchronize waits for the kernel to finish, and returns any errors encountered during the launch.
    cudaStatus = cudaDeviceSynchronize();
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "cudaDeviceSynchronize returned error code %d after launching addKernel!\n", cudaStatus);
        goto Error;
    }

// Copy output vector from GPU buffer to host memory.
    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;
}