编译DPC++ Compiler（支持Nvidia显卡）

开始使用DPC++

DPC++是Intel实现的SYCL版本用来支持异构计算，类似于OpenCL。

官方安装教程

预备条件

请确保当前开发环境满足如下条件:

• git
• cmake版本需要满足3.14及以上。
• python版本3.6以上的python。
• nijia版本1.8及以上（使用1.8版本以下可以供非llvm开发使用）。
• gcc版本7.1.0以上。
• cuda版本11.2通过测试可行。

当然也可以通过docker技术来实现开发环境的配置：

docker run --name sycl_build -it -v /local/workspace/dir/:/src ghcr.io/intel/llvm/ubuntu2004_base /bin/bash

若当前开发环境中有多个版本的编译环境，需要指定编译环境：

下面命令默认在bash环境下运行

指定g++版本：

export CXX=/<path>/gcc-9.2.0/bin/g++

指定gcc版本:

export CC=/<path>/gcc-9.2.0/bin/gcc

指定libstdc++相关动态链接位置

LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/<path>/gcc-9.2.0/lib64

创建工作空间

export DPCPP_HOME=~/sycl_workspace
mkdir $DPCPP_HOME
cd $DPCPP_HOME
git clone https://github.com/intel/llvm -b sycl

构建DPC++工具链

配置

运行下面脚本可实现自动化配置：

python $DPCPP_HOME/llvm/buildbot/configure.py

configure脚本存在如下参数：

• --system-ocl使用系统自带的opencl
• --no-werror编译 llvm 时不要将警告视为错误
• --cuda使用cuda作为后端
• --hip使用hip作为后端
• --hip-platform指定hip的平台(Nvida 或 Amd)
• --shared-libs构建动态链接库
• -t构建类型（debug或release）
• -o构建路径
• --cmake-gen设置构建系统类型

若需要使用cuda，可以通过如下指令进行配置：

python $DPCPP_HOME/llvm/buildbot/configure.py --cuda -t release

需要注意python版本务必在3.6以上cmake版本在3.14以上，否则会出现很多异常。
编译需要一个较大的硬盘空间。

可能出现的错误

无法找到Ninja

错误信息

CMake Error: CMake was unable to find a build program corresponding to "Ninja".  CMAKE_MAKE_PROGRAM is not set.  You probably need to select a different build tool.
CMake Error: CMAKE_C_COMPILER not set, after EnableLanguage
CMake Error: CMAKE_CXX_COMPILER not set, after EnableLanguage
CMake Error: CMAKE_ASM_COMPILER not set, after EnableLanguage

方法1

将Ninja的绝对路径配置到PATH中：

export PATH=$PATH:/depot/ninja-1.7.1/

方法2
更改configure.py在参数中指定CMAKE_MAKE_PROGRAM的值为ninja的绝对路径。
方法3
在CMakeCache中指定Ninja的绝对路径

 CMAKE_MAKE_PROGRAM:FILEPATH=/depot/ninja-1.7.1/ninja

无法找到Cuda相关工具

方法1：修改configure.py文件

...
cmake_cmd = [
     "cmake3",
...
     "-DCUDA_TOOLKIT_ROOT_DIR=/depot/cuda/cuda-11.2/" #指定cuda路径
    ]
...

方法2
将cuda的绝对路径配置到PATH中。

编译

运行下面脚本可实现自动化编译：

python $DPCPP_HOME/llvm/buildbot/compile.py

compile参数

• -o -> 构建路径
• -t -> 构建目标
• -j -> 构建使用的线程数

可能出现的错误

libstdc++版本错误

错误信息

 /lib64/libstdc++.so.6: version `GLIBCXX_3.4.20' not found 

解决方案

修改本地LD_LIBRARY位置

export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:/depot/gcc-9.2.0/lib64/lib

python错误

错误信息

...
ILED: _deps/vc-intrinsics-build/GenXIntrinsics/include/llvm/GenXIntrinsics/GenXIntrinsicDescription.gen 
cd /remote/us01home50/yuhangli/sycl_workspace/llvm/build/_deps/vc-intrinsics-build/GenXIntrinsics/include/llvm/GenXIntrinsics && /usr/local/bin/python -B /remote/us01home50/yuhangli/sycl_workspace/llvm/build/_deps/vc-intrinsics-src/GenXIntrinsics/include/llvm/GenXIntrinsics/Intrinsics.py /remote/us01home50/yuhangli/sycl_workspace/llvm/build/_deps/vc-intrinsics-src/GenXIntrinsics/include/llvm/GenXIntrinsics/Intrinsic_definitions.py /remote/us01home50/yuhangli/sycl_workspace/llvm/build/_deps/vc-intrinsics-build/GenXIntrinsics/include/llvm/GenXIntrinsics/GenXIntrinsicDescription.gen
....

错误原因
python版本不满足要求导致的

解决方案
修改CMakeCache.txt下python的位置为指定版本Python的位置

PYTHON_EXECUTABLE:FILEPATH=/usr/bin/python3.6

运行

配置运行环境

设置oneAPI环境：

export PATH=$PATH:$DPCPP_HOME/llvm/build/bin/
export PATH=$PATH:/slowfs/fs_model5/yhli/oneAPI/llvm/build/bin/

设置oneAPI链接库：

export LD_LIBRARY_PATH=$LD_LIBRARY_PATH:$DPCPP_HOME/llvm/build/lib
export  LD_LIBRARY_PATH=/depot/gcc-9.2.0/lib64:/slowfs/fs_model5/yhli/oneAPI/llvm/build/lib

测试用例

#include <CL/sycl.hpp>
int main() {
  // Creating buffer of 4 ints to be used inside the kernel code
   cl::sycl::buffer<cl::sycl::cl_int, 1> Buffer(4);
   // Creating SYCL queue
   cl::sycl::queue Queue;
 
   // Size of index space for kernel
   cl::sycl::range<1> NumOfWorkItems{Buffer.size()};
  
  std::cout << "Selected device: " <<
  Queue.get_device().get_info<sycl::info::device::name>() << "\n";
  // Submitting command group(work) to queue
  Queue.submit([&](cl::sycl::handler &cgh) {
    // Getting write only access to the buffer on a device
    auto Accessor = Buffer.get_access<cl::sycl::access::mode::write>(cgh);
    // Executing kernel
    cgh.parallel_for<class FillBuffer>(
        NumOfWorkItems, [=](cl::sycl::id<1> WIid) {
          // Fill buffer with indexes
          Accessor[WIid] = (cl::sycl::cl_int)WIid.get(0);
        });
  });

  // Getting read only access to the buffer on the host.
  // Implicit barrier waiting for queue to complete the work.
    const auto HostAccessor = Buffer.get_access<cl::sycl::access::mode::read>();
 // Check the results
 bool MismatchFound = false;
 for (size_t I = 0; I < Buffer.size(); ++I) {
   if (HostAccessor[I] != I) {
     std::cout << "The result is incorrect for element: " << I
                 << " , expected: " << I << " , got: " << HostAccessor[I]
                << std::endl;
      MismatchFound = true;
    }
  }

  if (!MismatchFound) {
    std::cout << "The results are correct!" << std::endl;
  }

  return MismatchFound;
  return 0;
}

编译指令

clang++ -Wall -std=c++17  -fsycl --cuda-path="/depot/cuda/cuda-11.2" -fsycl-targets=nvptx64-nvidia-cuda --gcc-toolchain="/depot/gcc-9.2.0" -O3 source_code.cpp -o <application_name> 

当需要依赖MKL时：

dpcpp -Wall --gcc-toolchain="/depot/gcc-9.2.0" -DMKL_ILP64 -lmkl_sycl -lmkl_intel_ilp64 -lmkl_tbb_thread -lmkl_core -std=c++17 -O3 <source_code>.cpp -o <application_name>