[Bf-blender-cvs] [fee6dd19724] cycles_oneapi: Cycles: add intel/llvm sycl compiler support
Xavier Hallade
noreply at git.blender.org
Thu Apr 21 10:14:30 CEST 2022
Commit: fee6dd197242e5bd9b46f7873ce5fd08281e857f
Author: Xavier Hallade
Date: Tue Apr 19 14:44:10 2022 +0200
Branches: cycles_oneapi
https://developer.blender.org/rBfee6dd197242e5bd9b46f7873ce5fd08281e857f
Cycles: add intel/llvm sycl compiler support
to Intel GPUs backend on Windows.
tested with release 2021-12: https://github.com/intel/llvm/releases/tag/2021-12
with the below workaround for intel/llvm headers:
sed -i 's/(detail::declptr/(::sycl::detail::declptr/g' ./build/install/include/sycl/CL/sycl/nd_item.hpp
===================================================================
M build_files/cmake/Modules/FindSYCL.cmake
M intern/cycles/kernel/CMakeLists.txt
M intern/cycles/kernel/device/oneapi/compat.h
M intern/cycles/kernel/device/oneapi/kernel.cpp
===================================================================
diff --git a/build_files/cmake/Modules/FindSYCL.cmake b/build_files/cmake/Modules/FindSYCL.cmake
index 0b42da8cf4f..6a35ff360c4 100644
--- a/build_files/cmake/Modules/FindSYCL.cmake
+++ b/build_files/cmake/Modules/FindSYCL.cmake
@@ -4,7 +4,7 @@
# - Find SYCL library
# Find the native SYCL header and libraries needed by oneAPI implementation
# This module defines
-# SYCL_DPCPP_COMPILER, compiler from oneAPI toolkit, which will be used for compilation of SYCL code
+# SYCL_COMPILER, compiler which will be used for compilation of SYCL code
# SYCL_LIBRARY, libraries to link against in order to use SYCL.
# SYCL_INCLUDE_DIR, directories where SYCL headers can be found
# SYCL_ROOT_DIR, The base directory to search for SYCL files.
@@ -23,7 +23,7 @@ SET(_sycl_search_dirs
C:/Program\ Files\ \(x86\)/Intel/oneAPI/compiler/latest/windows
)
-FIND_PROGRAM(SYCL_DPCPP_COMPILER
+FIND_PROGRAM(SYCL_COMPILER
NAMES
dpcpp
HINTS
@@ -32,7 +32,7 @@ FIND_PROGRAM(SYCL_DPCPP_COMPILER
bin
)
-FIND_LIBRARY(_SYCL_LIBRARY
+FIND_LIBRARY(SYCL_LIBRARY
NAMES
sycl
HINTS
@@ -41,7 +41,7 @@ FIND_LIBRARY(_SYCL_LIBRARY
lib64 lib
)
-FIND_PATH(_SYCL_INCLUDE_DIR
+FIND_PATH(SYCL_INCLUDE_DIR
NAMES
CL/sycl.hpp
HINTS
@@ -53,23 +53,15 @@ FIND_PATH(_SYCL_INCLUDE_DIR
INCLUDE(FindPackageHandleStandardArgs)
-FIND_PACKAGE_HANDLE_STANDARD_ARGS(SYCL DEFAULT_MSG _SYCL_LIBRARY _SYCL_INCLUDE_DIR)
+FIND_PACKAGE_HANDLE_STANDARD_ARGS(SYCL DEFAULT_MSG SYCL_LIBRARY SYCL_INCLUDE_DIR)
IF(SYCL_FOUND)
- SET(SYCL_LIBRARY ${_SYCL_LIBRARY})
-
- get_filename_component(_SYCL_INCLUDE_PARENT_DIR ${_SYCL_INCLUDE_DIR} DIRECTORY)
-
- SET(SYCL_INCLUDE_DIR ${_SYCL_INCLUDE_DIR} ${_SYCL_INCLUDE_PARENT_DIR})
+ get_filename_component(_SYCL_INCLUDE_PARENT_DIR ${SYCL_INCLUDE_DIR} DIRECTORY)
+ SET(SYCL_INCLUDE_DIR ${SYCL_INCLUDE_DIR} ${_SYCL_INCLUDE_PARENT_DIR})
ELSE()
SET(SYCL_SYCL_FOUND FALSE)
ENDIF()
MARK_AS_ADVANCED(
- SYCL_LIBRARY
- SYCL_INCLUDE_DIR
- SYCL_DPCPP_COMPILER
- _SYCL_INCLUDE_DIR
_SYCL_INCLUDE_PARENT_DIR
- _SYCL_LIBRARY
)
diff --git a/intern/cycles/kernel/CMakeLists.txt b/intern/cycles/kernel/CMakeLists.txt
index f2e68287f4c..1464f66e512 100644
--- a/intern/cycles/kernel/CMakeLists.txt
+++ b/intern/cycles/kernel/CMakeLists.txt
@@ -719,7 +719,7 @@ if(WITH_CYCLES_DEVICE_ONEAPI)
${SRC_UTIL_HEADERS}
)
- set(dpcpp_flags
+ set(sycl_compiler_flags
${CMAKE_CURRENT_SOURCE_DIR}/${SRC_KERNEL_DEVICE_ONEAPI}
-fsycl
-fsycl-unnamed-lambda
@@ -736,25 +736,25 @@ if(WITH_CYCLES_DEVICE_ONEAPI)
${ONEAPI_DPCPP_FLAGS}
)
if (WITH_CYCLES_ONEAPI_SYCL_HOST_ENABLED)
- list(APPEND dpcpp_flags -DWITH_ONEAPI_SYCL_HOST_ENABLED)
+ list(APPEND sycl_compiler_flags -DWITH_ONEAPI_SYCL_HOST_ENABLED)
endif()
if (WITH_CYCLES_ONEAPI_BINARIES)
set (CYCLES_ONEAPI_GPU_COMPILATION_OPTIONS "-internal_options '-ze-opt-large-register-file -ze-opt-regular-grf-kernel integrator_intersect'")
- list(APPEND dpcpp_flags
+ list(APPEND sycl_compiler_flags
-fsycl-targets=spir64_gen
-Xsycl-target-backend=spir64_gen "-device ${CYCLES_ONEAPI_AOT_TARGETS} ${CYCLES_ONEAPI_GPU_COMPILATION_OPTIONS}")
endif()
if(WITH_NANOVDB)
- list(APPEND dpcpp_flags
+ list(APPEND sycl_compiler_flags
-DWITH_NANOVDB
-I"${NANOVDB_INCLUDE_DIR}")
endif()
- get_filename_component(dpcpp_root ${SYCL_DPCPP_COMPILER} DIRECTORY)
+ get_filename_component(sycl_compiler_root ${SYCL_COMPILER} DIRECTORY)
+ get_filename_component(sycl_compiler_compiler_name ${SYCL_COMPILER} NAME_WE)
if(WIN32)
- list(APPEND dpcpp_flags
- /EHsc
+ list(APPEND sycl_compiler_flags
-fms-extensions
-fms-compatibility
-D_WINDLL
@@ -763,16 +763,31 @@ if(WITH_CYCLES_DEVICE_ONEAPI)
-D_WINDOWS
-DONEAPI_EXPORT)
- add_custom_command(
- OUTPUT ${cycles_kernel_oneapi_lib}
- COMMAND "${dpcpp_root}/../../env/vars.bat"
- COMMAND ${SYCL_DPCPP_COMPILER} ${dpcpp_flags}
- DEPENDS ${cycles_oneapi_kernel_sources})
+ if(sycl_compiler_compiler_name MATCHES "dpcpp")
+ add_custom_command(
+ OUTPUT ${cycles_kernel_oneapi_lib}
+ COMMAND "${sycl_compiler_root}/../../env/vars.bat"
+ COMMAND ${SYCL_COMPILER} ${sycl_compiler_flags}
+ DEPENDS ${cycles_oneapi_kernel_sources})
+ else()
+ string(REPLACE /Redist/ /Tools/ MSVC_TOOLS_DIR ${MSVC_REDIST_DIR})
+ list(APPEND sycl_compiler_flags
+ -L "${MSVC_TOOLS_DIR}/lib/x64"
+ -L "${WINDOWS_KITS_DIR}/Lib/${CMAKE_VS_WINDOWS_TARGET_PLATFORM_VERSION}/um/x64"
+ -L "${WINDOWS_KITS_DIR}/Lib/${CMAKE_VS_WINDOWS_TARGET_PLATFORM_VERSION}/ucrt/x64")
+ add_custom_command(
+ OUTPUT ${cycles_kernel_oneapi_lib}
+ COMMAND ${CMAKE_COMMAND} -E env
+ "LIB=${sycl_compiler_root}/../lib;${LIB}"
+ "PATH=${sycl_compiler_root}/../lib/ocloc;${sycl_compiler_root};${PATH}"
+ ${SYCL_COMPILER} ${sycl_compiler_flags}
+ DEPENDS ${cycles_oneapi_kernel_sources})
+ endif()
else()
- list(APPEND dpcpp_flags -fPIC)
+ list(APPEND sycl_compiler_flags -fPIC)
add_custom_command(
OUTPUT ${cycles_kernel_oneapi_lib}
- COMMAND bash -c \"source ${dpcpp_root}/../../env/vars.sh&&${SYCL_DPCPP_COMPILER} ${dpcpp_flags}\"
+ COMMAND bash -c \"source ${sycl_compiler_root}/../../env/vars.sh&&${SYCL_COMPILER} ${sycl_compiler_flags}\"
DEPENDS ${cycles_oneapi_kernel_sources})
endif()
diff --git a/intern/cycles/kernel/device/oneapi/compat.h b/intern/cycles/kernel/device/oneapi/compat.h
index 0c0ec827a71..d737f5a8029 100644
--- a/intern/cycles/kernel/device/oneapi/compat.h
+++ b/intern/cycles/kernel/device/oneapi/compat.h
@@ -192,6 +192,7 @@ ccl_always_inline float3 make_float3(float x)
#include "util/half.h"
#include "util/types.h"
+
// NOTE(sirgienko) Declaring these functions after types headers is very important because they
// include oneAPI headers, which transitively include math.h headers which will cause redefintions
// of the math defines because math.h also uses them and having them defined before math.h include
diff --git a/intern/cycles/kernel/device/oneapi/kernel.cpp b/intern/cycles/kernel/device/oneapi/kernel.cpp
index a80e04edc87..8d480566dea 100644
--- a/intern/cycles/kernel/device/oneapi/kernel.cpp
+++ b/intern/cycles/kernel/device/oneapi/kernel.cpp
@@ -10,6 +10,7 @@
# include <set>
# include <level_zero/ze_api.h>
+# include <CL/sycl.hpp>
# include <ext/oneapi/backend/level_zero.hpp>
# include "kernel/device/oneapi/compat.h"
@@ -185,7 +186,7 @@ bool oneapi_trigger_runtime_compilation(SyclQueue *queue_)
try {
queue->submit([&](sycl::handler &cgh) {
- sycl::accessor A_acc(A, cgh, sycl::read_only, sycl::no_init);
+ sycl::accessor A_acc(A, cgh, sycl::read_only);
sycl::accessor B_acc(B, cgh, sycl::write_only, sycl::no_init);
cgh.parallel_for(N, [=](sycl::id<1> idx) { B_acc[idx] = A_acc[idx] + idx.get(0); });
More information about the Bf-blender-cvs
mailing list