1、GPUfusion简介

为了能够支持CUDA编程环境，设计了CUDA兼容模块“GPUfusion”。通过该模块，用户可以无感知地在DTK平台编写CUDA程序，并快速地将已有原生CUDA应用在DTK上进行部署，简化原生CUDA程序的移植流程。

GPUfusion已内置在DTK工具包内，使用前需要设置相应的环境变量。在DTK的cuda目录下面提供了设置相应环境变量的脚本env.sh，设置方法如下：

1）首先加载DTK目录下的env.sh，用于设置DTK环境变量：

source dtk/env.sh

2）然后加载DTK的cuda目录下的env.sh，用于设置CUDA环境变量：

source dtk/cuda/env.sh

环境变量加载成功后，则可以按照CUDA的用法，进行程序构建、编译、运行和测试。目前主要覆盖了接口功能、编译选项、构建等方面的兼容，从而支持在DTK平台编译运行CUDA程序。

2． CUDA程序移植示例

基本流程

使用GPUfusion进行CUDA程序移植，一般包括程序构建、程序编译、程序运行、运行结果检查四个阶段。

l 程序构建

1. 基于Makefile

对于以Makefile方式构建的应用来说，相关的CUDA路径已经固化在Makefile中，一般会以“/usr/local/cuda”为路径，因此需要用户手动进行相关CUDA路径的替换。

1. 基于CMake

GPUfusion支持find_package(CUDA)和enable_language(CUDA)的CMake特性。可以在不修改任何构建系统的基础上，引导CMake将相关目录指向GPUfusion的文件夹，使用方法与原生CUDA一致。

l 程序编译

在程序整体编译过程中如果出现错误，建议先使用单线程输出详细信息，以输出真正错误处的编译命令，方便问题分析。指令如下：

make VERBOSE=1

l 程序运行

GPUfusion为了程序最高的效率，默认线程块的最大值为256，在程序运行中如果出现以下提示，则需要调整默认值的大小。

Launch params (1024, 1, 1) are larger than
launch bounds (256) for kernel _ZL14timedReductionPKfPfPl please add
launch_bounds to kernel define or use --gpu-max-threads-per-block recompile
program !

修改的方法有两种：

1）在编译命令中添加如下参数，然后重新编译：

--gpu-max-threads-per-block=1024

2）在核函数头部添加：launch_bounds (1024)。示例如下：

launch_bounds (1024)

global static void timedReduction(const
float *input, float *output) {

l 运行结果检查

如果程序运行结果不正确，此时需要检查程序中warp size的设置。默认原生CUDA的warp size为32，GPUfusion适用的warp size为64。用户需要确保项目中所有的warp size修改为64，以保证程序的正确性。

示例

这里使用一个简单的axpy.cu原生CUDA程序，来演示对于四种不同CUDA构建编译方式的支持，默认已加载了GPUfusion的环境变量，CMake使用3.16.2版本。

#include <iostream> 

global void axpy(float a, float* x,
float* y) {

y[threadIdx.x] = a * x[threadIdx.x];

}

int main(int argc, char* argv[]) {

const int kDataLen = 4;

float a = 2.0f;

float host_x[kDataLen] = {1.0f, 2.0f, 3.0f, 4.0f};

float host_y[kDataLen];

float host_y_res[kDataLen];

//
Copy input data to device.

float* device_x;

float* device_y;

cudaMalloc((void**)&device_x, kDataLen * sizeof(float));

cudaMalloc((void**)&device_y, kDataLen * sizeof(float));

cudaMemcpy(device_x, host_x, kDataLen * sizeof(float),

cudaMemcpyHostToDevice);

//
Launch the kernel.

axpy<<<1, kDataLen>>>(a, device_x, device_y);

//
Copy output data to host.

cudaDeviceSynchronize();

cudaMemcpy(host_y, device_y, kDataLen * sizeof(float),

cudaMemcpyDeviceToHost);

for(int i=0; i< kDataLen; ++i) {

host_y_res[i] = a * host_x[i];

}

int
errors = 0;

double eps = 1.0e-6;

for(int i=0; i< kDataLen; ++i) {

if

(std::abs(host_y_res[i] - host_y[i]) > eps)

errors++;

}

if
(errors != 0) {

std::cout << "FAILED: " << errors << "
errors" << std::endl;

}

else {

std::cout << "PASSED!"<< std::endl;

}

cudaDeviceReset();

return 0;

}

l 直接编译CUDA程序

$ nvcc axpy.cu -o axpy

$ ./axpy

PASSED!

l 使用FindCUDA的CMake项目

1.
CMakeLists.txt内容

cmake_minimum_required(VERSION 3.0)

project(MyApp)

find_package(CUDA)

cuda_add_executable(axpy axpy.cu)

2.
CMake输出

-- The C compiler identification is GNU
7.3.1

-- The CXX compiler identification is GNU
7.3.1

-- Check for working C compiler:
/opt/rh/devtoolset-7/root/usr/bin/cc

-- Check for working C compiler:
/opt/rh/devtoolset-7/root/usr/bin/cc -- works

-- Detecting C compiler ABI info

-- Detecting C compiler ABI info - done

-- Detecting C compile features

-- Detecting C compile features - done

-- Check for working CXX compiler:
/opt/rh/devtoolset-7/root/usr/bin/c++

-- Check for working CXX compiler:
/opt/rh/devtoolset-7/root/usr/bin/c++ -- works

-- Detecting CXX compiler ABI info

-- Detecting CXX compiler ABI info - done

-- Detecting CXX compile features

-- Detecting CXX compile features - done

-- Found CUDA: cuda (found version
"10.2")

-- Configuring done

-- Generating done

-- Build files have been written to:
/root/document/build

3.
make及其执行结果

$ make

[ 50%] Building NVCC (Device) object
CMakeFiles/axpy.dir/axpy_generated_axpy.cu.o

Scanning dependencies of target axpy

[100%] Linking CXX executable axpy

[100%] Built target axpy

$ ./axpy

PASSED!

l 使用enable_language的CMake项目

1.
CMakeLists.txt内容

cmake_minimum_required(VERSION 3.0)

project(MyApp)

enable_language(CUDA)

add_executable(axpy axpy.cu)

2.
CMake输出

-- The C compiler identification is GNU
7.3.1

-- The CXX compiler identification is GNU
7.3.1

-- Check for working C compiler: /opt/rh/devtoolset-7/root/usr/bin/cc

-- Check for working C compiler:
/opt/rh/devtoolset-7/root/usr/bin/cc -- works

-- Detecting C compiler ABI info

-- Detecting C compiler ABI info - done

-- Detecting C compile features

-- Detecting C compile features - done

-- Check for working CXX compiler:
/opt/rh/devtoolset-7/root/usr/bin/c++

-- Check for working CXX compiler:
/opt/rh/devtoolset-7/root/usr/bin/c++ -- works

-- Detecting CXX compiler ABI info

-- Detecting CXX compiler ABI info - done

-- Detecting CXX compile features

-- Detecting CXX compile features - done

-- The CUDA compiler identification is
NVIDIA 10.2.89

-- Check for working CUDA compiler: /usr/local/cuda/bin/nvcc

-- Check for working CUDA compiler:
/usr/local/cuda/bin/nvcc -- works

-- Detecting CUDA compiler ABI info

-- Detecting CUDA compiler ABI info - done

-- Configuring done

-- Generating done

-- Build files have been written to:
/root/document/build

3.
make及其执行结果

$ make

Scanning dependencies of target axpy

[ 50%] Building CUDA object CMakeFiles/axpy.dir/axpy.cu.o

[100%] Linking CUDA executable axpy

[100%] Built target axpy

$ ./axpy

PASSED!

l 使用project设置语言的CMake项目

1.
CMakeLists.txt内容

cmake_minimum_required(VERSION 3.0)

project(MyApp CUDA)

add_executable(axpy axpy.cu)

2.
CMake输出

-- The CUDA compiler identification is
NVIDIA 10.2.89

-- Check for working CUDA compiler:
/usr/local/cuda/bin/nvcc

-- Check for working CUDA compiler:
/usr/local/cuda/bin/nvcc -- works

-- Detecting CUDA compiler ABI info

-- Detecting CUDA compiler ABI info - done

-- Configuring done

-- Generating done

-- Build files have been written to:
/root/document/build

3.
make及其执行结果

$ make

Scanning dependencies of target axpy

[ 50%] Building CUDA object
CMakeFiles/axpy.dir/axpy.cu.o

[100%] Linking CUDA executable axpy

[100%] Built target axpy

$ ./axpy

PASSED!