【高性能计算】CUDA编程之环境配置与小试牛刀（教程与代码-1）

Carlos橘子柠檬

已于 2022-03-04 20:37:52 修改

阅读量1.4k

点赞数 1

分类专栏： CUDA 文章标签： vscode cuda gpu

于 2022-02-23 21:13:27 首次发布

本文链接：https://blog.csdn.net/heroybc/article/details/123099324

版权

CUDA 专栏收录该内容

4 篇文章 0 订阅

订阅专栏

前言

最近在学习linux环境下的CUDA编程，力求实现AI算法加速应用，整理些许笔记如下，代码经过调试，可以跑通。基础环境为mac本地，通过vscode远程控制ubuntu的开发板。

1 CUDA开发环境

环境配置

检查clang，检查所使用的mac是否自带了clang，敲命令行：

clang -v // or clang --version

MacOS为了照顾用户习惯，也可以执行gcc命令，但MacOS10.12以后系统不自带gcc。
如果没有的话则去安装，敲命令行：

xcode-select --install

按照步骤安装完成

Linux

gcc
C++环境配置与测试
打开VS Code按快捷键 shift+command+X ，并搜索C++
安装Code Runner
创建一个demo文件
新建一个C文件(文件后缀为.c)并保存至文件夹。

/* 文件hello.c */
#include <stdio.h>
int main()
{
    printf("hello world");
    return 0;
}

点击打开.c文件
点击IDE上方菜单，Run -> Start Debugging -> 选择C/C++: clang build active file 得到 launch.json
点击IDE上方菜单，Terminal-> Configure Tasks-> 选择 C/C++:gcc build active file 得到 tasks.json
Run Code ，终端会输出Hello World
CUDA环境配置与测试
使用代码
必须保存为.cu扩展名

#include <iostream>
#include <stdio.h>
__global__ void myfirstkernel (void) {  // 内核调用，在主机程序中运行设备代码
}
int main (void) {
    myfirstkernel << < 1, 1>> >();
    printf( "Hello, CUDA!\n" );
    return 0;
}

命令行输入
编译 nvcc filename.cu -o filename (-o后面是可执行文件名)。
执行 sudo ./filename

2 使用CUDA C进行编程

2.1 CUDA程序结构

算术逻辑单元（ALU）
可以将GPU看作多个块（Block）的组合，每个块可以执行多个线程。每个块绑定到GPU上的不同流多处理器。来自同一块的线程可以相互通信。
CUDA C程序的开发步骤如下：
1）为主机和设备显存中的数据分配内存。
2）将数据从主机内存复制到设备显存。
3）通过指定并行度来启动内核。
4）所有线程完成后，将数据从设备显存复制回主机内存。
5）释放主机和设备上使用的所有内存。
主机代码是由标准C编译器编译的，设备代码是由NVIDIA GPU编译器来执行。
使用ANSI C关键字和CUDA扩展关键字编写的设备代码称为内核。它是主机代码（Host Code）通过内核调用的方式来启动的。简单地说，内核调用的含义是我们从主机代码启动设备代码。
内核调用通常会生成大量的块（Block）和线程（Thread）来在GPU上并行地处理数据。
作为参数传递给内核的指针应该仅指向设备显存。
每个块在流多处理器上运行，一个块中的线程可以通过共享内存（Shared Memory）彼此通信。程序员无法选定哪个流多处理器将执行特定的块，也无法选定块和线程以何种顺序执行。
使用设备指针从设备函数中读取和写入设备显存，并且应该使用主机指针从主机函数中读取和写入主机内存。
双值相加代码：

#include <iostream>
#include <cuda.h>
#include <cuda_runtime.h>
#include <stdio.h>
//Definition of kernel function to add two variables
__global__ void gpuAdd(int d_a, int d_b, int *d_c) {
    *d_c = d_a + d_b;
}
//main function
int main(void) {
    //Defining host variable to store answer
    int h_c;
    //Defining device pointer
    int *d_c;
    //Allocating memory for device pointer
    cudaMalloc((void**)&d_c, sizeof(int));
    //Kernel call by passing 1 and 4 as inputs and storing answer 
    //<< <1,1> >> means 1 block is executed with 1 thread per block
    gpuAdd << <1, 1 >> > (1, 4, d_c);
    //Copy result from device memory to host memory
    cudaMemcpy(&h_c, d_c, sizeof(int), cudaMemcpyDeviceToHost);
    printf("1 + 4 = %d\n", h_c);
    //Free up memory
    cudaFree(d_c);
    return 0;
}

以上代码可以改为指针变量形式

#include <iostream>
#include <cuda.h>
#include <cuda_runtime.h>
#include <stdio.h>
//Definition of kernel function to add two variables
__global__ void gpuAdd(int *d_a, int *d_b, int *d_c) {
    *d_c = *d_a + *d_b;
}
//main function
int main(void) {
    //Defining host variable to store answer
    int h_c, h_a = 1, h_b = 4;
    //Defining device pointer
    int *d_c, *d_a, *d_b;
    //Allocating memory for device pointer
    cudaMalloc((void**)&d_c, sizeof(int));
    cudaMalloc((void**)&d_a, sizeof(int));
    cudaMalloc((void**)&d_b, sizeof(int));
    cudaMemcpy(d_a, &h_a, sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, &h_b, sizeof(int), cudaMemcpyHostToDevice);
    //Kernel call by passing 1 and 4 as inputs and storing answer
    //<< <1,1> >> means 1 block is executed with 1 thread per block
    gpuAdd <<<1, 1 >>> (d_a, d_b, d_c);
    //Copy result from device memory to host memory
    cudaMemcpy(&h_c, d_c, sizeof(int), cudaMemcpyDeviceToHost);
    printf("%d + %d = %d\n", h_a, h_b, h_c);
    //Free up memory
    cudaFree(d_c);
    cudaFree(d_b);
    cudaFree(d_a);
    return 0;
}

打印CUDA基本信息
打印块ID
打印设备数

#include <memory>
#include <iostream>
#include <cuda_runtime.h>
#include <stdio.h>
__global__ void myfirstkernel(void) {
    //blockIdx.x gives the block number of current kernel
    printf("Hello!!!I'm thread in block: %d\n", blockIdx.x);
}
// Main Program 
int main(void)
{
    // 打印设备数
    int device_Count = 0;
    cudaGetDeviceCount(&device_Count);
    // This function returns count of number of CUDA enable devices 
    // and 0 if there are no CUDA capable devices.
    if (device_Count == 0)
    {
        printf("There are no available device that support CUDA\n");
    }
    else
    {
        printf("Detected %d CUDA Capable device(s)\n", device_Count);
    }
    // 打印执行线程
    //A kernel call with 16 blocks and 1 thread per block
    myfirstkernel << <16,1>> >();
    //Function used for waiting for all kernels to finish
    cudaDeviceSynchronize();
    printf("All threads are finished!\n");
    return 0;
}

blockIdx.x内置变量读取到块
threadIdx.x内置变量线程
CUDA函数调用：cudaDeviceSynchronize()。为何要加这句？这是因为启动内核是一个异步操作，只要发布了内核启动命令，不等内核执行完成，控制权就会立刻返回给调用内核的CPU线程。
GPU上的内存为分层结构。它可以分为L1缓存、L2缓存、全局内存、纹理内存和共享内存。使用GPU多线程计算简单例子

#include "stdio.h"
#include<iostream>
//Defining Number of elements in Array
#define N   5
//Defining vector addition function for CPU
void cpuAdd(int *h_a, int *h_b, int *h_c) {
    int tid = 0;    
    while (tid < N)
    {
        h_c[tid] = h_a[tid] + h_b[tid];
        tid += 1;
    }
}
// 使用GPU相加计算
__global__ void gpuAdd(int *d_a, int *d_b, int *d_c) {
    //Getting block index of current kernel
    int tid = blockIdx.x;   // handle the data at this index
    if (tid < N)
        d_c[tid] = d_a[tid] + d_b[tid];
}
int main(void) {
    int h_a[N], h_b[N], h_c[N];
    //Defining device pointers
    int *d_a, *d_b, *d_c;
    // allocate the memory
    cudaMalloc((void**)&d_a, N * sizeof(int));
    cudaMalloc((void**)&d_b, N * sizeof(int));
    cudaMalloc((void**)&d_c, N * sizeof(int));
    //Initializing two arrays for addition
    for (int i = 0; i < N; i++) {
        h_a[i] = 2 * i*i;
        h_b[i] = i;
    }
    //Calling CPU function for vector addition
    cpuAdd (h_a, h_b, h_c);
    //Printing Answer
    printf("Vector addition on CPU\n");
    for (int i = 0; i < N; i++) {
        printf("CPU -> sum of %d is %d + %d = %d\n", i, h_a[i], h_b[i], h_c[i]);
    }
    // Copy input arrays from host to device memory
    cudaMemcpy(d_a, h_a, N * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, N * sizeof(int), cudaMemcpyHostToDevice);
    //  Calling kernels with N blocks and one thread per block, 
    //  passing device pointers as parameters
    gpuAdd << <N, 1 >> >(d_a, d_b, d_c);
    //Copy result back to host memory from device memory
    cudaMemcpy(h_c, d_c, N * sizeof(int), cudaMemcpyDeviceToHost);
        printf("Vector addition on GPU \n");
    //Printing result on console
    for (int i = 0; i < N; i++) {
        printf("GPU -> sum of %d is %d + %d = %d\n", i, h_a[i], h_b[i], h_c[i]);
    }
    //Free up memory
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
    return 0;
}