CUDA编程--Square函数

写在前面的话： 上一篇文章我们编写了‘Hello World’, 这篇文章我们再进一步！

首先，我们理一下CUDA的编程思路：

host: cpu

device: gpu

1.分配host内存, 并进行数据初始化;
2.分配device内存, 并从host将数据拷贝到device上;
3.调用CUDA的核函数在device上完成指定的运算;
4.将device 上的运算结果拷贝到host上 (性能)
5.释放device和host上分配的内存

简单点说，就是把CPU变量搬到GPU上，把函数改写到GPU上，在GPU上运算，再把结果搬回CPU。 函数改写到GPU只需要注意线程和块，其他基本一致。 我们下面用数组平方的实例来进一步了解。

实战：数组平方

这一次的代码我们不像之前那样在终端进行编译，我们使用IDE进行编译和运行，毕竟稍微长一点的代码还是IDE用着比较舒服。我使用的IDE是Clion，可以自行百度安装。 用IDE的话，我们需要建立4个文件： CMakelists.txt, main.cpp, square.cu, square.h 。使用Clion新建项目 CMakelists.txt, main.cpp会自动生成。剩余两个文件需要我们新建生成：

CMakelists.txt： 用于配置编译选项，例如文件路径，cuda选项，可执行文件等

main.cpp： 主函数cpp ，不多说

square.cu : Cuda 编程文件，在该文件中编写Cuda程序

square.h: 头文件，用于申明函数，使square中的函数可以被cpp调用

完整的代码可以在我的github中获得：GitHub - weiguangzhao/cuda_demohttps://github.com/weiguangzhao/cuda_demo

square.cu

# include "square.h"
# include <stdio.h>


__global__ void square(float * d_out, float * d_in){
    int idx = threadIdx.x;
    float f =d_in[idx];
    d_out[idx] = f*f;
};

void square_test(){

    // 定义数组长度和位数
    int array_size = 64;
    int array_bytes = array_size * sizeof(float);

    // 产生数组
    float h_in[array_size];  //输入数组
    for (int  i=0; i< array_size; i++){
        h_in[i] = float(i);
    }
    float h_out[array_size]; //存储结果

    // 定义GPU内存指针
    float * d_in;
    float * d_out;

    // 分配GPU内存
    cudaMalloc((void **) &d_in, array_bytes);
    cudaMalloc((void **) &d_out, array_bytes);

    // 把CPU数据搬到GPU上
    cudaMemcpy(d_in, h_in, array_bytes, cudaMemcpyHostToDevice);

    // 运行cuda内核开始计算
    square<<<1, array_size>>>(d_out, d_in);

    // 把运算结果搬回CPU
    cudaMemcpy(h_out, d_out, array_bytes, cudaMemcpyDeviceToHost);

    // 打印结果
    for (int i =0; i< array_size; i++){
        printf("%f \n", h_out[i]);
    }

    // 释放GPU内存
    cudaFree(d_in);
    cudaFree(d_out);
}

从square文件中，可以清楚的看到CUDA编程逻辑，先定义CPU变量h_in, h_out (h指host)，初始化CPU变量值，再分配GPU空间，再将CPU数据搬移到GPU上。Square函数，我们利用GPU线程索引去实现并行运算。通过以下代码，我们设置了一个块和64个线程去执行square函数，这也就意味着，一次并行运算便可以实现64维数组的平方计算。

// 运行cuda内核开始计算
    square<<<1, array_size>>>(d_out, d_in);

运算完成后， 我们将结果（d_out） 搬回CPU 并进行打印。

square.h

头文件，我们需要申明CPU函数， 而不是带有__global__ 的GPU函数

#ifndef DEMO1_SQUARE_SQUARE_H
#define DEMO1_SQUARE_SQUARE_H

//void square(float * d_out, float * d_in);
void square_test();

#endif //DEMO1_SQUARE_SQUARE_H

main.cpp

主函数更简单了， 就是在main中调用头文件申明的函数即可。

#include <iostream>

#include "square.h"

int main() {
    std::cout << "Hello, World!" << std::endl;
    square_test();
    std::cout << "Hello, World!" << std::endl;
    return 0;
}

CMakelists.txt

编译文件和C++的略有不同，生成的文件为cuda可执行文件，还需要检测cuda包是否存在，以及Cuda配置选项。其中Cuda的算力arch=sm_75可以在该网址查询：CUDA GPU | NVIDIA Developer

cmake_minimum_required(VERSION 3.19)
project(demo1_square)

set(CMAKE_CXX_STANDARD 14)

# packages
find_package(CUDA REQUIRED)
if(${CUDA_FOUND})
    set(CUDA_SOURCE_PROPERTY_FORMAT OBJ)
    set(CUDA_SEPARABLE_COMPILATION ON)
    include_directories(${CUDA_INCLUDE_DIRS})
    set(CUDA_PROPAGATE_HOST_FLAGS OFF)
    set(CUDA_NVCC_FLAGS -arch=sm_75;-O3;-G;-g;-std=c++11)# 根据具体GPU性能更改算力参数
    #SET(CUDA_NVCC_FLAGS ${CUDA_NVCC_FLAGS};-gencode arch=compute_61,code=sm_61;-std=c++11;-O3;-G;-g)
    link_directories($ENV{CUDA_PATH}/lib/x64)
    MESSAGE(STATUS "found cuda")
else(${CUDA_FOUND})
    MESSAGE(STATUS "cuda not found!")
endif(${CUDA_FOUND})

file(GLOB CURRENT_SOURCES  *.cpp *.cu)
file(GLOB_RECURSE CURRENT_HEADERS  *.h *.hpp *.cuh)

source_group("Source" FILES ${CURRENT_SOURCES})
source_group("Include" FILES ${CURRENT_HEADERS})

CUDA_ADD_EXECUTABLE(demo1_square ${CURRENT_HEADERS} ${CURRENT_SOURCES})

编写好文件后，我们只需要右键Reload Cmakelists.txt文件，便可以点击右上角运行按钮进行运行

 结果如下：

 参考链接：

WIN10下查看CUDA版本 & 显卡计算能力_boniu2019的博客-CSDN博客

CUDA CMakeLists_fb_941219的博客-CSDN博客

CUDA学习第１课：用CUDA做数的平方_哔哩哔哩_bilibili