USE_BATCHED cmake option listing target library to build. New BATCHED_REGRESSION testing mode. (#1357)

* Introduce cmake option USE_BATCHED to enabling building batched SIMD shader execution for a comma separated list of batch sizes, ISA's, and noFMA: (0, b8_AVX, b8_AVX2, b8_AVX2_noFMA, b8_AVX512, b8_AVX512_noFMA, b16_AVX512, b16_AVX512_noFMA)
When USE_BATCHED is set, -DOSL_USE_BATCHED=1 will be passed allowing for source files to conditionally compile in batched execution support.
When USE_BATCHED is set, cmake variable BUILD_BATCHED=True and projects can conditionally add batched sources to be built.
Updated OSL source with #ifdef OSL_USE_BATCHED/#endif were needed.
Improved/streamlined liboslexec/CMakeLists.txt to iterate over targets in USE_BATCHED vs. hardcoding all variants.
Updated _OSL_SUPPORTS_* defines to match USE_BATCHED values simplying creating compiler definitions in Cmake.

Added new BATCHED_REGRESSION testing mode, when a file named BATCHED_REGRESSION exists in a testsuite folder, the run.py will be executed and its output files copied to "./baseline".  The same run.py is then run again but with environment variable TESTSHADE_BATCHED=1, which cause testshade to execute in batched mode.  Then the results/image comparison happens between the batched output files and those in "./baseline".  The "./ref" dir is never used in a regression test.
Presence of BATCHED_REGRESSION prevents non-regression tests from being created for a folder.

Converted arithmetic-cov to be a BATCHED_REGRESSION test.
Removed testsuite/arithmetic-cov/ref image files.

Renamed ShadingSystem::supports_batch_execution_at(int width) to ShadingSystem::configure_batch_execution_at(int width) to better identify its behavior:  based on currently set attributes for llvm_jit_target and llvm_jit_fma, test if current machine is capable of supporting batched execution at the specified width.  If no specific llvm_jit_target was requested, sets llvm_jit_target and llvm_jit_fma to the supported configuration for the requested width.  Returns true if supported, false otherwise

Fix bug in batched Matrix*Matrix code generation where the order of arguments were being reversed (swap was intended to transform float*Matrix to a Matrix*float function call, but accidentally happened for Matrix*Matrix)

* Moved define of OSL_USE_BATCHED to oslconfig.h.in
Fixed order of includes in batched*.cpp so oslconfig.h is higher

* Renamed testsuite/arithmetic-cov to testsuite/arithmetic-reg to better represent it is a regression test.
Changed #ifdef OSL_USE_BATCHED to #if OSL_USE_BATCHED and rely on it being set to 0 or 1 by <OSL/oslconfig.h>

Signed-off-by: Alex M. Wells <alex.m.wells@intel.com>
This commit is contained in:
Alex Wells
2021-04-29 21:04:21 -07:00
committed by GitHub
parent ab7107cf23
commit 4292cc9974
238 changed files with 501 additions and 333 deletions
+8
View File
@@ -167,6 +167,10 @@ ifneq (${USE_SIMD},)
MY_CMAKE_FLAGS += -DUSE_SIMD:STRING="${USE_SIMD}"
endif
ifneq (${USE_BATCHED},)
MY_CMAKE_FLAGS += -DUSE_BATCHED:STRING="${USE_BATCHED}"
endif
ifneq (${TEST},)
TEST_FLAGS += -R ${TEST}
endif
@@ -407,6 +411,10 @@ help:
@echo " 0, sse2, sse3, ssse3, sse4.1, sse4.2, f16c,"
@echo " avx, avx2, avx512f)"
@echo " USE_OPTIX=1 Build the OptiX test renderer"
@echo " USE_BATCHED=targets Build batched SIMD execution of shaders for (comma-separated choices:"
@echo " 0, b8_AVX, b8_AVX2, b8_AVX2_noFMA,"
@echo " b8_AVX512, b8_AVX512_noFMA,"
@echo " b16_AVX512, b16_AVX512_noFMA)"
@echo " make test, extra options:"
@echo " TEST=regex Run only tests matching the regex"
@echo ""
+26
View File
@@ -283,6 +283,32 @@ if (NOT USE_SIMD STREQUAL "")
add_compile_options (${SIMD_COMPILE_FLAGS})
endif ()
###########################################################################
# Batched SIMD shader execution options.
#
# The USE_BATCHED option may be set to indicate that support for batched
# SIMD shader execution be compiled along with targe specific libraries
set (USE_BATCHED "" CACHE STRING "Build batched SIMD shader execution for (0, b8_AVX, b8_AVX2, b8_AVX2_noFMA, b8_AVX512, b8_AVX512_noFMA, b16_AVX512, b16_AVX512_noFMA)")
set (BATCHED_SUPPORT_DEFINES "")
set (BATCHED_TARGET_LIBS "")
set (BUILD_BATCHED False)
if (NOT USE_BATCHED STREQUAL "")
message (STATUS "Compiling with batched SIMD targets ${USE_BATCHED}")
if (NOT USE_BATCHED STREQUAL "0")
set (BUILD_BATCHED 1)
string (REPLACE "," ";" BATCHED_TARGET_LIST ${USE_BATCHED})
foreach (batched_target ${BATCHED_TARGET_LIST})
if (VERBOSE)
message (STATUS "BATCHED target: ${batched_target}")
endif ()
set (BATCHED_SUPPORT_DEFINES ${BATCHED_SUPPORT_DEFINES} "__OSL_SUPPORTS_${batched_target}")
endforeach()
else ()
message (STATUS "Compiling with NO batched SIMD targets, USE_BATCHED=0")
endif ()
else ()
message (STATUS "Compiling with NO batched SIMD targets, USE_BATCHED is empty")
endif ()
###########################################################################
# Preparation to test for compiler/language features
+28 -14
View File
@@ -100,7 +100,8 @@ macro ( TESTSUITE )
# Run the test unoptimized, unless it matches a few patterns that
# we don't test unoptimized (or has an OPTIMIZEONLY marker file).
if (NOT _testname MATCHES "^getattribute-shader" AND
NOT _testname MATCHES "optix" AND
NOT _testname MATCHES "optix" AND
NOT EXISTS "${_testsrcdir}/BATCHED_REGRESSION" AND
NOT EXISTS "${_testsrcdir}/OPTIMIZEONLY")
add_one_testsuite ("${_testname}" "${_testsrcdir}"
ENV TESTSHADE_OPT=0 )
@@ -109,7 +110,8 @@ macro ( TESTSUITE )
# optimization, triggered by setting TESTSHADE_OPT env variable.
# Skip OptiX-only tests and those with a NOOPTIMIZE marker file.
if (NOT _testname MATCHES "optix"
AND NOT EXISTS "${_testsrcdir}/NOOPTIMIZE")
AND NOT EXISTS "${_testsrcdir}/BATCHED_REGRESSION"
AND NOT EXISTS "${_testsrcdir}/NOOPTIMIZE")
add_one_testsuite ("${_testname}.opt" "${_testsrcdir}"
ENV TESTSHADE_OPT=2 )
endif ()
@@ -120,7 +122,8 @@ macro ( TESTSUITE )
if (USE_OPTIX
AND (EXISTS "${_testsrcdir}/OPTIX" OR test_all_optix OR _testname MATCHES "optix")
AND NOT EXISTS "${_testsrcdir}/NOOPTIX"
AND NOT EXISTS "${_testsrcdir}/NOOPTIX-FIXME")
AND NOT EXISTS "${_testsrcdir}/NOOPTIX-FIXME"
AND NOT EXISTS "${_testsrcdir}/BATCHED_REGRESSION")
# Unoptimized
if (NOT EXISTS "${_testsrcdir}/OPTIMIZEONLY")
add_one_testsuite ("${_testname}.optix" "${_testsrcdir}"
@@ -131,16 +134,27 @@ macro ( TESTSUITE )
ENV TESTSHADE_OPT=2 TESTSHADE_OPTIX=1 )
endif ()
# When building for Batched support, also run it in Batched mode
# if there is an BATCHED marker file in the directory.
# If an environment variable $TESTSUITE_BATCHED is nonzero, then
# run all tests in Batched mode, even if there's no BATCHED marker.
if ((EXISTS "${_testsrcdir}/BATCHED" OR test_all_batched OR _testname MATCHES "batched")
AND NOT EXISTS "${_testsrcdir}/NOBATCHED"
AND NOT EXISTS "${_testsrcdir}/NOBATCHED-FIXME")
# optimized for right now
add_one_testsuite ("${_testname}.batched.opt" "${_testsrcdir}"
ENV TESTSHADE_OPT=2 TESTSHADE_BATCHED=1 )
if (BUILD_BATCHED)
# When building for Batched support, also run it in Batched mode
# if there is an BATCHED marker file in the directory.
# If an environment variable $TESTSUITE_BATCHED is nonzero, then
# run all tests in Batched mode, even if there's no BATCHED marker.
if ((EXISTS "${_testsrcdir}/BATCHED" OR test_all_batched OR _testname MATCHES "batched")
AND NOT EXISTS "${_testsrcdir}/BATCHED_REGRESSION"
AND NOT EXISTS "${_testsrcdir}/NOBATCHED"
AND NOT EXISTS "${_testsrcdir}/NOBATCHED-FIXME")
# optimized for right now
add_one_testsuite ("${_testname}.batched.opt" "${_testsrcdir}"
ENV TESTSHADE_OPT=2 TESTSHADE_BATCHED=1 )
endif ()
# When building for Batched support, optionally run a regression test
# if there is an BATCHED_REGRESSION marker file in the directory.
if (EXISTS "${_testsrcdir}/BATCHED_REGRESSION")
# optimized for right now
add_one_testsuite ("${_testname}.regress.batched.opt" "${_testsrcdir}"
ENV TESTSHADE_OPT=2 OSL_REGRESSION_TEST=BATCHED )
endif ()
endif ()
endforeach ()
if (VERBOSE)
@@ -153,7 +167,7 @@ macro (osl_add_all_tests)
# List all the individual testsuite tests here, except those that need
# special installed tests.
TESTSUITE ( aastep allowconnect-err and-or-not-synonyms arithmetic
arithmetic-cov
arithmetic-reg
array array-derivs array-range array-aassign
blackbody blendmath breakcont
bug-array-heapoffsets bug-locallifetime bug-outputinit
+4
View File
@@ -4,6 +4,10 @@
#pragma once
#if !defined(OSL_USE_BATCHED) || (OSL_USE_BATCHED == 0)
# error batched_shaderglobals.h should not be included unless OSL_USE_BATCHED is defined to 1
#endif
#include <OSL/wide.h>
#include <OSL/shaderglobals.h>
+2 -1
View File
@@ -4,7 +4,6 @@
#pragma once
#include <OSL/platform.h>
#include <OSL/export.h>
#include <OSL/oslversion.h>
@@ -47,6 +46,8 @@
OSL_NAMESPACE_ENTER
#define OSL_USE_BATCHED @BUILD_BATCHED@
/// Various compile-time defaults are defined here that could, in
/// principle, be redefined if you are using OSL in some particular
/// renderer that wanted things a different way.
+9 -3
View File
@@ -23,7 +23,9 @@ struct PerThreadInfo;
class ShadingContext;
class ShaderSymbol;
class OSLQuery;
#if OSL_USE_BATCHED
template<int WidthT> struct alignas(64) BatchedShaderGlobals;
#endif
@@ -685,10 +687,14 @@ public:
const void* symbol_address (const ShadingContext &ctx,
const ShaderSymbol *sym) const;
#if OSL_USE_BATCHED
/// Based on currently set attributes for llvm_jit_target and
/// llvm_jit_fma, test if current machine is capable of supporting
/// batched execution at the specified width
bool supports_batch_execution_at(int width);
/// batched execution at the specified width. If no specific
/// target was requested, sets llvm_jit_target and llvm_jit_fma
/// to the supported configuration for the requested width.
/// Returns true if supported, false otherwise
bool configure_batch_execution_at(int width);
template<int WidthT>
class OSLEXECPUBLIC BatchedExecutor {
@@ -725,7 +731,7 @@ public:
OSL_FORCEINLINE BatchedExecutor<WidthT> batched() {
return BatchedExecutor<WidthT>(*this);
}
#endif
/// Return the statistics output as a huge string.
+161 -213
View File
@@ -11,148 +11,6 @@ set ( liboslexec_target_srcs
set ( liboslexec_override_limits
)
# Strategy is to make a copy of each cpp in liboslexec_target_srcs for
# each target batch width and ISA combination, applying compiler flags to define
# -D__OSL_WIDTH=[4|8|16]
# -D__OSL_TARGET_ISA=[AVX512|AVX2|AVX|SSE4_2]
# and compiler flags to choose correct target ISA for your compiler
# You may then add/remove the desired ${TARGET_SOURCES_B[4|8|16]_[AVX512|AVX2|AVX|SSE4_2]}
# to your add_library call
foreach(target_src ${liboslexec_target_srcs})
set(DST_B16_AVX512 "${CMAKE_CURRENT_BINARY_DIR}/${target_src}_b16_AVX512.cpp")
set(DST_B16_AVX512_NOFMA "${CMAKE_CURRENT_BINARY_DIR}/${target_src}_b16_AVX512_noFMA.cpp")
set(DST_B8_AVX512 "${CMAKE_CURRENT_BINARY_DIR}/${target_src}_b8_AVX512.cpp")
set(DST_B8_AVX512_NOFMA "${CMAKE_CURRENT_BINARY_DIR}/${target_src}_b8_AVX512_noFMA.cpp")
set(DST_B8_AVX2 "${CMAKE_CURRENT_BINARY_DIR}/${target_src}_b8_AVX2.cpp")
set(DST_B8_AVX2_NOFMA "${CMAKE_CURRENT_BINARY_DIR}/${target_src}_b8_AVX2_noFMA.cpp")
set(DST_B8_AVX "${CMAKE_CURRENT_BINARY_DIR}/${target_src}_b8_AVX.cpp")
set(SRC "${CMAKE_CURRENT_SOURCE_DIR}/${target_src}.cpp")
CONFIGURE_FILE("${SRC}" ${DST_B16_AVX512} COPYONLY)
CONFIGURE_FILE("${SRC}" ${DST_B16_AVX512_NOFMA} COPYONLY)
CONFIGURE_FILE(${SRC} ${DST_B8_AVX512} COPYONLY)
CONFIGURE_FILE(${SRC} ${DST_B8_AVX512_NOFMA} COPYONLY)
CONFIGURE_FILE(${SRC} ${DST_B8_AVX2} COPYONLY)
CONFIGURE_FILE(${SRC} ${DST_B8_AVX2_NOFMA} COPYONLY)
CONFIGURE_FILE(${SRC} ${DST_B8_AVX} COPYONLY)
list(APPEND TARGET_SOURCES_B16_AVX512 "${DST_B16_AVX512}")
list(APPEND TARGET_SOURCES_B16_AVX512_NOFMA "${DST_B16_AVX512_NOFMA}")
list(APPEND TARGET_SOURCES_B8_AVX512 "${DST_B8_AVX512}")
list(APPEND TARGET_SOURCES_B8_AVX512_NOFMA "${DST_B8_AVX512_NOFMA}")
list(APPEND TARGET_SOURCES_B8_AVX2 "${DST_B8_AVX2}")
list(APPEND TARGET_SOURCES_B8_AVX2_NOFMA "${DST_B8_AVX2_NOFMA}")
list(APPEND TARGET_SOURCES_B8_AVX "${DST_B8_AVX}")
set_property(SOURCE "${DST_B16_AVX512}" "${DST_B16_AVX512_NOFMA}" "${DST_B8_AVX512}" "${DST_B8_AVX512_NOFMA}" "${DST_B8_AVX2}" "${DST_B8_AVX2_NOFMA}" "${DST_B8_AVX}" "${DST_B4_SSE4_2}"
APPEND PROPERTY COMPILE_OPTIONS
"-I${CMAKE_CURRENT_SOURCE_DIR}/wide"
"-I${CMAKE_CURRENT_SOURCE_DIR}/../liboslnoise/wide"
)
set_property(SOURCE "${DST_B16_AVX512}" "${DST_B16_AVX512_NOFMA}"
APPEND PROPERTY COMPILE_DEFINITIONS "__OSL_WIDTH=16")
set_property(SOURCE "${DST_B8_AVX512}" "${DST_B8_AVX512_NOFMA}" "${DST_B8_AVX2}" "${DST_B8_AVX2_NOFMA}" "${DST_B8_AVX}"
APPEND PROPERTY COMPILE_DEFINITIONS "__OSL_WIDTH=8")
set_property(SOURCE "${DST_B16_AVX512}" "${DST_B8_AVX512}"
APPEND PROPERTY COMPILE_DEFINITIONS "__OSL_TARGET_ISA=AVX512")
set_property(SOURCE "${DST_B16_AVX512_NOFMA}" "${DST_B8_AVX512_NOFMA}"
APPEND PROPERTY COMPILE_DEFINITIONS "__OSL_TARGET_ISA=AVX512_noFMA")
set_property(SOURCE "${DST_B8_AVX2}"
APPEND PROPERTY COMPILE_DEFINITIONS "__OSL_TARGET_ISA=AVX2")
set_property(SOURCE "${DST_B8_AVX2_NOFMA}"
APPEND PROPERTY COMPILE_DEFINITIONS "__OSL_TARGET_ISA=AVX2_noFMA")
set_property(SOURCE "${DST_B8_AVX}"
APPEND PROPERTY COMPILE_DEFINITIONS "__OSL_TARGET_ISA=AVX")
if (CMAKE_COMPILER_IS_INTEL)
if (MSVC)
set_property(SOURCE "${DST_B16_AVX512}" "${DST_B16_AVX512_NOFMA}"
APPEND PROPERTY COMPILE_OPTIONS "/QxCORE-AVX512" "/Qopt-zmm-usage:high")
set_property(SOURCE "${DST_B8_AVX512}" "${DST_B8_AVX512_NOFMA}"
APPEND PROPERTY COMPILE_OPTIONS "/QxCORE-AVX512" "/Qopt-zmm-usage:low")
set_property(SOURCE "${DST_B8_AVX2}" "${DST_B8_AVX2_NOFMA}"
APPEND PROPERTY COMPILE_OPTIONS "/QxCORE-AVX2")
set_property(SOURCE "${DST_B8_AVX}"
APPEND PROPERTY COMPILE_OPTIONS "/QxAVX")
set_property(SOURCE "${DST_B16_AVX512_NOFMA}" "${DST_B8_AVX512_NOFMA}" "${DST_B8_AVX2_NOFMA}"
APPEND PROPERTY COMPILE_OPTIONS "/Qfma-")
if (${target_src} IN_LIST liboslexec_override_limits)
set_property(SOURCE "${DST_B16_AVX512}" "${DST_B16_AVX512_NOFMA}" "${DST_B8_AVX512}" "${DST_B8_AVX512_NOFMA}"
"${DST_B8_AVX2}" "${DST_B8_AVX2_NOFMA}" "${DST_B8_AVX}"
APPEND PROPERTY COMPILE_OPTIONS "/Qoverride-limits")
endif()
set_property(SOURCE "${DST_B16_AVX512}" "${DST_B16_AVX512_NOFMA}" "${DST_B8_AVX512}" "${DST_B8_AVX512_NOFMA}"
"${DST_B8_AVX2}" "${DST_B8_AVX2_NOFMA}" "${DST_B8_AVX}"
APPEND PROPERTY COMPILE_OPTIONS "/Qprec-div" "/Qqopt-report:5")
else ()
set_property(SOURCE "${DST_B16_AVX512}" "${DST_B16_AVX512_NOFMA}"
APPEND PROPERTY COMPILE_OPTIONS "-xCORE-AVX512" "-qopt-zmm-usage=high")
set_property(SOURCE "${DST_B8_AVX512}" "${DST_B8_AVX512_NOFMA}"
APPEND PROPERTY COMPILE_OPTIONS "-xCORE-AVX512" "-qopt-zmm-usage=low")
set_property(SOURCE "${DST_B8_AVX2}" "${DST_B8_AVX2_NOFMA}"
APPEND PROPERTY COMPILE_OPTIONS "-xCORE-AVX2")
set_property(SOURCE "${DST_B8_AVX}"
APPEND PROPERTY COMPILE_OPTIONS "-xAVX")
set_property(SOURCE "${DST_B16_AVX512_NOFMA}" "${DST_B8_AVX512_NOFMA}" "${DST_B8_AVX2_NOFMA}"
APPEND PROPERTY COMPILE_OPTIONS "-no-fma")
if (${target_src} IN_LIST liboslexec_override_limits)
set_property(SOURCE "${DST_B16_AVX512}" "${DST_B16_AVX512_NOFMA}" "${DST_B8_AVX512}" "${DST_B8_AVX512_NOFMA}"
"${DST_B8_AVX2}" "${DST_B8_AVX2_NOFMA}" "${DST_B8_AVX}"
APPEND PROPERTY COMPILE_OPTIONS "-qoverride-limits")
endif()
set_property(SOURCE "${DST_B16_AVX512}" "${DST_B16_AVX512_NOFMA}" "${DST_B8_AVX512}" "${DST_B8_AVX512_NOFMA}"
"${DST_B8_AVX2}" "${DST_B8_AVX2_NOFMA}" "${DST_B8_AVX}"
APPEND PROPERTY COMPILE_OPTIONS "-prec-div" "-qopt-report=5")
endif ()
endif ()
if (CMAKE_COMPILER_IS_CLANG OR CMAKE_COMPILER_IS_APPLECLANG)
# REQUIRES CLANG 7+ to use sucessfully vectorize with OpenMP simd
set_property(SOURCE "${DST_B16_AVX512}" "${DST_B16_AVX512_NOFMA}" "${DST_B8_AVX512}" "${DST_B8_AVX512_NOFMA}"
APPEND PROPERTY COMPILE_OPTIONS "-march=skylake-avx512")
if (CLANG_VERSION_STRING VERSION_GREATER 7.0)
set_property(SOURCE "${DST_B16_AVX512}" "${DST_B16_AVX512_NOFMA}"
APPEND PROPERTY COMPILE_OPTIONS "-mprefer-vector-width=512")
set_property(SOURCE "${DST_B8_AVX512}" "${DST_B8_AVX512_NOFMA}"
APPEND PROPERTY COMPILE_OPTIONS "-mprefer-vector-width=256")
endif ()
set_property(SOURCE "${DST_B8_AVX2}" "${DST_B8_AVX2_NOFMA}"
APPEND PROPERTY COMPILE_OPTIONS "-march=core-avx2")
set_property(SOURCE "${DST_B8_AVX}"
APPEND PROPERTY COMPILE_OPTIONS "-march=corei7-avx")
set_property(SOURCE "${DST_B16_AVX512_NOFMA}" "${DST_B8_AVX512_NOFMA}" "${DST_B8_AVX2_NOFMA}"
APPEND PROPERTY COMPILE_OPTIONS "-mno-fma")
# large SIMD function loops will exceed llvm's -inline-threshold default of 225.
# remark: loop not vectorized: call instruction cannot be vectorized [-Rpass-analysis]
# choose to increase that limit via compiler flags vs.
# workaround with __attribute__((flatten))
set_property(SOURCE "${DST_B16_AVX512}" "${DST_B16_AVX512_NOFMA}" "${DST_B8_AVX512}" "${DST_B8_AVX512_NOFMA}"
"${DST_B8_AVX2}" "${DST_B8_AVX2_NOFMA}" "${DST_B8_AVX}"
APPEND PROPERTY COMPILE_OPTIONS "-mllvm" "-inline-threshold=100000")
# For loops with small loop bodies, clang was unrolling the loop before
# #pragma omp simd
# was processed. Thus there was no loop left for the pragma to apply to
# To work around this, for these wide SIMD libraries, we will disable
# clang's loop unrolling. We really don't want to do this and wish to remove
# it as soon as clang addresses this issue.
set_property(SOURCE "${DST_B16_AVX512}" "${DST_B16_AVX512_NOFMA}" "${DST_B8_AVX512}" "${DST_B8_AVX512_NOFMA}"
"${DST_B8_AVX2}" "${DST_B8_AVX2_NOFMA}" "${DST_B8_AVX}"
APPEND PROPERTY COMPILE_OPTIONS "-fno-unroll-loops")
endif ()
endforeach(target_src)
set (local_lib oslexec)
set (lib_src
shadingsys.cpp closure.cpp
@@ -164,18 +22,22 @@ set (lib_src
opspline.cpp opstring.cpp optexture.cpp
oslexec.cpp
pointcloud.cpp rendservices.cpp
batched_rendservices.cpp
constfold.cpp runtimeoptimize.cpp typespec.cpp
lpexp.cpp lpeparse.cpp automata.cpp accum.cpp
opclosure.cpp
shadeimage.cpp
backendllvm.cpp
llvm_gen.cpp llvm_instance.cpp llvm_util.cpp
)
if (BUILD_BATCHED)
list(APPEND lib_src
batched_analysis.cpp
batched_backendllvm.cpp
batched_llvm_gen.cpp
batched_llvm_instance.cpp
)
batched_rendservices.cpp
)
endif()
if (BUILD_SHARED_LIBS)
# oslcomp symbols used in oslexec
@@ -341,77 +203,167 @@ endif ()
# For each target specific shared library being built, we need
# we need to let a couple source files know so they can build
# function mappings for those targets
set_property(SOURCE "batched_llvm_instance.cpp" "shadingsys.cpp"
APPEND PROPERTY COMPILE_DEFINITIONS
"__OSL_SUPPORTS_B16_AVX512"
"__OSL_SUPPORTS_B8_AVX512"
"__OSL_SUPPORTS_B8_AVX2"
"__OSL_SUPPORTS_B8_AVX"
#"__OSL_SUPPORTS_B16_AVX512_NOFMA"
#"__OSL_SUPPORTS_B8_AVX512_NOFMA"
#"__OSL_SUPPORTS_B8_AVX2_NOFMA"
)
set_property(SOURCE "batched_llvm_instance.cpp" "shadingsys.cpp"
APPEND PROPERTY COMPILE_DEFINITIONS ${BATCHED_SUPPORT_DEFINES})
set (_b16_AVX512_oslexec_lib _b16_AVX512_oslexec)
set (_b8_AVX512_oslexec_lib _b8_AVX512_oslexec)
set (_b8_AVX2_oslexec_lib _b8_AVX2_oslexec)
set (_b8_AVX_oslexec_lib _b8_AVX_oslexec)
set (_b16_AVX512_noFMA_oslexec_lib _b16_AVX512_noFMA_oslexec)
set (_b8_AVX512_noFMA_oslexec_lib _b8_AVX512_noFMA_oslexec)
set (_b8_AVX2_noFMA_oslexec_lib _b8_AVX2_noFMA_oslexec)
set (BATCHED_TARGET_LIBS)
foreach(batched_target ${BATCHED_TARGET_LIST})
set (batched_target_lib "_${batched_target}_oslexec")
list (APPEND BATCHED_TARGET_LIBS ${batched_target_lib})
string (REPLACE "_" ";" TARGET_OPTS ${batched_target})
list (GET TARGET_OPTS 0 TARGET_OPT_SIZE)
list (GET TARGET_OPTS 1 TARGET_OPT_ISA)
list (LENGTH TARGET_OPTS NUM_TARGET_OPTS)
set (TARGET_OPT_FMA "FMA")
set (TARGET_ISA ${TARGET_OPT_ISA})
if (NUM_TARGET_OPTS EQUAL 3)
list (GET TARGET_OPTS 2 TARGET_OPT_FMA)
set (TARGET_ISA "${TARGET_OPT_ISA}_${TARGET_OPT_FMA}")
endif()
string (SUBSTRING ${TARGET_OPT_SIZE} 1 -1 TARGET_BATCH_SIZE)
# Strategy is to make a copy of each cpp in liboslexec_target_srcs for
# each target batch width and ISA combination, applying compiler flags to define
# -D__OSL_WIDTH=[4|8|16]
# -D__OSL_TARGET_ISA=[AVX512|AVX2|AVX|SSE4_2]
# and compiler flags to choose correct target ISA for your compiler
# You may then add/remove the desired ${TARGET_SOURCES_B[4|8|16]_[AVX512|AVX2|AVX|SSE4_2]}
# to your add_library call
set (TARGET_LIB_SOURCES)
set (TARGET_CXX_DEFS)
set (TARGET_CXX_OPTS)
list (APPEND TARGET_CXX_OPTS
"-I${CMAKE_CURRENT_SOURCE_DIR}/wide"
"-I${CMAKE_CURRENT_SOURCE_DIR}/../liboslnoise/wide"
)
list (APPEND TARGET_CXX_DEFS
"__OSL_WIDTH=${TARGET_BATCH_SIZE}"
"__OSL_TARGET_ISA=${TARGET_ISA}"
)
add_library ( ${_b16_AVX512_oslexec_lib} MODULE "${TARGET_SOURCES_B16_AVX512}" )
add_library ( ${_b8_AVX512_oslexec_lib} MODULE "${TARGET_SOURCES_B8_AVX512}" )
add_library ( ${_b8_AVX2_oslexec_lib} MODULE "${TARGET_SOURCES_B8_AVX2}" )
add_library ( ${_b8_AVX_oslexec_lib} MODULE "${TARGET_SOURCES_B8_AVX}" )
#add_library ( ${_b16_AVX512_noFMA_oslexec_lib} MODULE "${TARGET_SOURCES_B16_AVX512_NOFMA}" )
#add_library ( ${_b8_AVX512_noFMA_oslexec_lib} MODULE "${TARGET_SOURCES_B8_AVX512_NOFMA}" )
#add_library ( ${_b8_AVX2_noFMA_oslexec_lib} MODULE "${TARGET_SOURCES_B8_AVX2_NOFMA}" )
set ( liboslexec_target_libs
${_b16_AVX512_oslexec_lib}
${_b8_AVX512_oslexec_lib}
${_b8_AVX2_oslexec_lib}
${_b8_AVX_oslexec_lib}
# ${_b16_AVX512_noFMA_oslexec_lib}
# ${_b8_AVX512_noFMA_oslexec_lib}
# ${_b8_AVX2_noFMA_oslexec_lib}
)
foreach(target_lib ${liboslexec_target_libs})
target_include_directories (${target_lib}
if (CMAKE_COMPILER_IS_INTEL)
if (MSVC)
if (${TARGET_OPT_ISA} STREQUAL "AVX512")
list (APPEND TARGET_CXX_OPTS "/QxCORE-AVX512")
if (${TARGET_BATCH_SIZE} STREQUAL "16")
list (APPEND TARGET_CXX_OPTS "/Qopt-zmm-usage:high")
else ()
list (APPEND TARGET_CXX_OPTS "/Qopt-zmm-usage:low")
endif ()
elseif (${TARGET_OPT_ISA} STREQUAL "AVX2")
list (APPEND TARGET_CXX_OPTS "/QxCORE-AVX2")
elseif (${TARGET_OPT_ISA} STREQUAL "AVX")
list (APPEND TARGET_CXX_OPTS "/QxAVX")
else ()
message (FATAL_ERROR "Unknown ISA=${TARGET_OPT_ISA} extract from USE_BATCHED entry ${batched_target}")
endif ()
if (${TARGET_OPT_FMA} STREQUAL "noFMA")
list (APPEND TARGET_CXX_OPTS "/Qfma-")
endif ()
list (APPEND TARGET_CXX_OPTS "/Qprec-div" "/Qqopt-report:5")
else ()
if (${TARGET_OPT_ISA} STREQUAL "AVX512")
list (APPEND TARGET_CXX_OPTS "-xCORE-AVX512")
if (${TARGET_BATCH_SIZE} STREQUAL "16")
list (APPEND TARGET_CXX_OPTS "-qopt-zmm-usage:high")
else ()
list (APPEND TARGET_CXX_OPTS "-qopt-zmm-usage:low")
endif ()
elseif (${TARGET_OPT_ISA} STREQUAL "AVX2")
list (APPEND TARGET_CXX_OPTS "-xCORE-AVX2")
elseif (${TARGET_OPT_ISA} STREQUAL "AVX")
list (APPEND TARGET_CXX_OPTS "-xAVX")
else ()
message (FATAL_ERROR "Unknown ISA=${TARGET_OPT_ISA} extract from USE_BATCHED entry ${batched_target}")
endif ()
if (${TARGET_OPT_FMA} STREQUAL "noFMA")
list (APPEND TARGET_CXX_OPTS "-no-fma")
endif ()
list (APPEND TARGET_CXX_OPTS "-prec-div" "-qopt-report=5")
endif ()
endif ()
if (CMAKE_COMPILER_IS_CLANG OR CMAKE_COMPILER_IS_APPLECLANG)
# REQUIRES CLANG 7+ to use sucessfully vectorize with OpenMP simd
if (${TARGET_OPT_ISA} STREQUAL "AVX512")
list (APPEND TARGET_CXX_OPTS "-march=skylake-avx512")
if (${TARGET_BATCH_SIZE} STREQUAL "16")
list (APPEND TARGET_CXX_OPTS "-mprefer-vector-width=512")
else ()
list (APPEND TARGET_CXX_OPTS "-mprefer-vector-width=256")
endif ()
elseif (${TARGET_OPT_ISA} STREQUAL "AVX2")
list (APPEND TARGET_CXX_OPTS "-march=core-avx2")
elseif (${TARGET_OPT_ISA} STREQUAL "AVX")
list (APPEND TARGET_CXX_OPTS "-march=corei7-avx")
else ()
message (FATAL_ERROR "Unknown ISA=${TARGET_OPT_ISA} extract from USE_BATCHED entry ${batched_target}")
endif ()
if (${TARGET_OPT_FMA} STREQUAL "noFMA")
list (APPEND TARGET_CXX_OPTS "-mno-fma")
endif ()
if (${target_src} IN_LIST liboslexec_override_limits)
list (APPEND TARGET_CXX_OPTS "-qoverride-limits")
endif()
# large SIMD function loops will exceed llvm's -inline-threshold default of 225.
# remark: loop not vectorized: call instruction cannot be vectorized [-Rpass-analysis]
# choose to increase that limit via compiler flags vs.
# workaround with __attribute__((flatten))
list (APPEND TARGET_CXX_OPTS "-mllvm" "-inline-threshold=100000")
# For loops with small loop bodies, clang was unrolling the loop before
# #pragma omp simd
# was processed. Thus there was no loop left for the pragma to apply to
# To work around this, for these wide SIMD libraries, we will disable
# clang's loop unrolling. We really don't want to do this and wish to remove
# it as soon as clang addresses this issue.
list (APPEND TARGET_CXX_OPTS "-fno-unroll-loops")
endif ()
foreach(target_src ${liboslexec_target_srcs})
set(TARGET_SRC "${CMAKE_CURRENT_BINARY_DIR}/${target_src}_${batched_target}.cpp")
set(SRC "${CMAKE_CURRENT_SOURCE_DIR}/${target_src}.cpp")
CONFIGURE_FILE("${SRC}" ${TARGET_SRC} COPYONLY)
list(APPEND TARGET_LIB_SOURCES "${TARGET_SRC}")
set_property(SOURCE ${TARGET_SRC}
APPEND PROPERTY COMPILE_DEFINITIONS ${TARGET_CXX_DEFS})
set_property(SOURCE ${TARGET_SRC}
APPEND PROPERTY COMPILE_OPTIONS ${TARGET_CXX_OPTS})
if (${target_src} IN_LIST liboslexec_override_limits)
if (CMAKE_COMPILER_IS_INTEL)
if (MSVC)
set_property(SOURCE ${TARGET_SRC} APPEND PROPERTY COMPILE_OPTIONS "/Qoverride-limits")
else ()
set_property(SOURCE ${TARGET_SRC} APPEND PROPERTY COMPILE_OPTIONS "-qoverride-limits")
endif()
endif()
endif()
endforeach(target_src)
add_library ( ${batched_target_lib} MODULE ${TARGET_LIB_SOURCES} )
target_include_directories (${batched_target_lib}
PUBLIC
${CMAKE_INSTALL_FULL_INCLUDEDIR}
${IMATH_INCLUDES}
)
target_compile_definitions (${target_lib}
PRIVATE
OSL_EXPORTS
CUDA_TARGET_ARCH="${CUDA_TARGET_ARCH}"
OSL_CUDA_VERSION="${CUDA_VERSION}"
OSL_OPTIX_VERSION="${OPTIX_VERSION}"
)
target_link_libraries (${target_lib}
target_link_libraries (${batched_target_lib}
PUBLIC
OpenImageIO::OpenImageIO
# For OpenEXR/Imath 3.x:
$<$<TARGET_EXISTS:Imath::Imath>:Imath::Imath>
$<$<TARGET_EXISTS:Imath::Half>:Imath::Half>
# For OpenEXR >= 2.4/2.5 with reliable exported targets
$<$<TARGET_EXISTS:IlmBase::Imath>:IlmBase::Imath>
$<$<TARGET_EXISTS:IlmBase::Half>:IlmBase::Half>
$<$<TARGET_EXISTS:IlmBase::IlmThread>:IlmBase::IlmThread>
$<$<TARGET_EXISTS:IlmBase::Iex>:IlmBase::Iex>
# For OpenEXR <= 2.3:
${ILMBASE_LIBRARIES}
PRIVATE
pugixml::pugixml
${Boost_LIBRARIES} ${CMAKE_DL_LIBS}
${CLANG_LIBRARIES}
${LLVM_LIBRARIES} ${LLVM_LDFLAGS} ${LLVM_SYSTEM_LIBRARIES}
)
endforeach(target_lib)
endforeach(batched_target)
add_library (${local_lib} ${lib_src})
target_include_directories (${local_lib}
@@ -483,13 +435,9 @@ endif ()
# add_dependencies (${local_lib} "${CMAKE_SOURCE_DIR}/src/build-scripts/hidesymbols.map")
install_targets ( _b16_AVX512_oslexec )
install_targets ( _b8_AVX512_oslexec )
install_targets ( _b8_AVX2_oslexec )
install_targets ( _b8_AVX_oslexec )
#install_targets ( _b16_AVX512_noFMA_oslexec )
#install_targets ( _b8_AVX512_noFMA_oslexec )
#install_targets ( _b8_AVX2_noFMA_oslexec )
foreach(batched_target_lib ${BATCHED_TARGET_LIBS})
install_targets ( ${batched_target_lib} )
endforeach(batched_target_lib)
install_targets (${local_lib})
+2
View File
@@ -13,6 +13,8 @@
#include <llvm/ADT/SmallVector.h>
#include <OSL/oslconfig.h>
#include <OSL/batched_rendererservices.h>
#include "batched_analysis.h"
+1 -1
View File
@@ -488,7 +488,7 @@ LLVMGEN (llvm_gen_mul)
func_spec.arg(Result,false,op_is_uniform);
Symbol* A_prime = &A;
Symbol* B_prime = &B;
if (B.typespec().is_matrix()) {
if (!A.typespec().is_matrix()) {
// Always pass the matrix as the 1st operand
std::swap(A_prime,B_prime);
}
+16 -14
View File
@@ -19,6 +19,8 @@
#include <OpenImageIO/sysutil.h>
#include <OpenImageIO/timer.h>
#include <OSL/oslconfig.h>
#include <OSL/batched_shaderglobals.h>
#include <OSL/batched_texture.h>
@@ -319,7 +321,7 @@ public:
// NOTE: Because builtindecl_wide_xmacro.h passes macros through the name
// parameter of DECL(name,signature), we must have a layer of indirection
// to allow those macros to expand. Thus the use of DECL_INDIRECT.
#ifdef __OSL_SUPPORTS_B16_AVX512
#ifdef __OSL_SUPPORTS_b16_AVX512
template<>
const NameAndSignature
ConcreteTargetLibraryHelper<16, TargetISA::AVX512>::library_functions[]
@@ -346,7 +348,7 @@ const char*
= "b16_AVX512_";
#endif
#ifdef __OSL_SUPPORTS_B16_AVX512_NOFMA
#ifdef __OSL_SUPPORTS_b16_AVX512_noFMA
template<>
const NameAndSignature
ConcreteTargetLibraryHelper<16, TargetISA::AVX512_noFMA>::library_functions[]
@@ -373,7 +375,7 @@ const char* ConcreteTargetLibraryHelper<
= "b16_AVX512_noFMA_";
#endif
#ifdef __OSL_SUPPORTS_B8_AVX512
#ifdef __OSL_SUPPORTS_b8_AVX512
template<>
const NameAndSignature
ConcreteTargetLibraryHelper<8, TargetISA::AVX512>::library_functions[]
@@ -400,7 +402,7 @@ const char*
= "b8_AVX512_";
#endif
#ifdef __OSL_SUPPORTS_B8_AVX512_NOFMA
#ifdef __OSL_SUPPORTS_b8_AVX512_noFMA
template<>
const NameAndSignature
ConcreteTargetLibraryHelper<8, TargetISA::AVX512_noFMA>::library_functions[]
@@ -427,7 +429,7 @@ const char* ConcreteTargetLibraryHelper<
= "b8_AVX512_noFMA_";
#endif
#ifdef __OSL_SUPPORTS_B8_AVX2
#ifdef __OSL_SUPPORTS_b8_AVX2
template<>
const NameAndSignature
ConcreteTargetLibraryHelper<8, TargetISA::AVX2>::library_functions[]
@@ -454,7 +456,7 @@ const char*
= "b8_AVX2_";
#endif
#ifdef __OSL_SUPPORTS_B8_AVX2_NOFMA
#ifdef __OSL_SUPPORTS_b8_AVX2_noFMA
template<>
const NameAndSignature
ConcreteTargetLibraryHelper<8, TargetISA::AVX2_noFMA>::library_functions[]
@@ -481,7 +483,7 @@ const char*
= "b8_AVX2_noFMA_";
#endif
#ifdef __OSL_SUPPORTS_B8_AVX
#ifdef __OSL_SUPPORTS_b8_AVX
template<>
const NameAndSignature
ConcreteTargetLibraryHelper<8, TargetISA::AVX>::library_functions[]
@@ -518,12 +520,12 @@ BatchedBackendLLVM::TargetLibraryHelper::build(ShadingContext* context,
switch (vector_width) {
case 16:
switch (target_isa) {
#ifdef __OSL_SUPPORTS_B16_AVX512
#ifdef __OSL_SUPPORTS_b16_AVX512
case TargetISA::AVX512:
return RetType(
new ConcreteTargetLibraryHelper<16, TargetISA::AVX512>());
#endif
#ifdef __OSL_SUPPORTS_B16_AVX512_NOFMA
#ifdef __OSL_SUPPORTS_b16_AVX512_noFMA
case TargetISA::AVX512_noFMA:
return RetType(
new ConcreteTargetLibraryHelper<16, TargetISA::AVX512_noFMA>());
@@ -534,27 +536,27 @@ BatchedBackendLLVM::TargetLibraryHelper::build(ShadingContext* context,
break;
case 8:
switch (target_isa) {
#ifdef __OSL_SUPPORTS_B8_AVX512
#ifdef __OSL_SUPPORTS_b8_AVX512
case TargetISA::AVX512:
return RetType(
new ConcreteTargetLibraryHelper<8, TargetISA::AVX512>());
#endif
#ifdef __OSL_SUPPORTS_B8_AVX512_NOFMA
#ifdef __OSL_SUPPORTS_b8_AVX512_noFMA
case TargetISA::AVX512_noFMA:
return RetType(
new ConcreteTargetLibraryHelper<8, TargetISA::AVX512_noFMA>());
#endif
#ifdef __OSL_SUPPORTS_B8_AVX2
#ifdef __OSL_SUPPORTS_b8_AVX2
case TargetISA::AVX2:
return RetType(
new ConcreteTargetLibraryHelper<8, TargetISA::AVX2>());
#endif
#ifdef __OSL_SUPPORTS_B8_AVX2_NOFMA
#ifdef __OSL_SUPPORTS_b8_AVX2_noFMA
case TargetISA::AVX2_noFMA:
return RetType(
new ConcreteTargetLibraryHelper<8, TargetISA::AVX2_noFMA>());
#endif
#ifdef __OSL_SUPPORTS_B8_AVX
#ifdef __OSL_SUPPORTS_b8_AVX
case TargetISA::AVX:
return RetType(
new ConcreteTargetLibraryHelper<8, TargetISA::AVX>());
+18 -3
View File
@@ -11,7 +11,12 @@
#include <OpenImageIO/timer.h>
#include <OpenImageIO/thread.h>
#include <OSL/oslconfig.h>
#if OSL_USE_BATCHED
#include <OSL/batched_shaderglobals.h>
#endif
#include <OSL/mask.h>
#include <OSL/wide.h>
@@ -203,6 +208,7 @@ ShadingContext::execute (ShaderGroup &sgroup, ShaderGlobals &ssg, bool run)
return result;
}
#if OSL_USE_BATCHED
template<int WidthT>
bool
ShadingContext::Batched<WidthT>::execute_init
@@ -360,7 +366,7 @@ ShadingContext::Batched<WidthT>::execute(ShaderGroup &sgroup, int batch_size, Ba
}
return result;
}
#endif
void
ShadingContext::record_error (ErrorHandler::ErrCode code,
@@ -423,6 +429,7 @@ ShadingContext::process_errors () const
// interleaved with other threads.
lock_guard lock (buffered_errors_mutex);
#if OSL_USE_BATCHED
if (execution_is_batched()) {
OSL_DASSERT(batch_size_executed <= MaxSupportedSimdLaneCount);
// Process each data lane separately and in the correct order
@@ -435,7 +442,9 @@ ShadingContext::process_errors () const
return mask.is_on(lane_mask);
});
}
} else {
} else
#endif
{
// Non-batch errors: ignore the mask, just print them out once
OSL_INTEL_PRAGMA(noinline)
process_errors_helper(shadingsys(), m_buffered_errors, 0, nerrors,
@@ -462,6 +471,7 @@ const void *
ShadingContext::symbol_data (const Symbol &sym) const
{
const ShaderGroup &sgroup (*group());
#if OSL_USE_BATCHED
if (execution_is_batched()) {
if (! sgroup.batch_jitted())
return NULL; // can't retrieve symbol if we didn't optimize & batched jit
@@ -470,7 +480,9 @@ ShadingContext::symbol_data (const Symbol &sym) const
// lives on the heap
return m_heap.get() + sym.wide_dataoffset();
}
} else {
} else
#endif
{
if (! sgroup.jitted())
return NULL; // can't retrieve symbol if we didn't optimize & jit
@@ -548,8 +560,11 @@ osl_incr_layers_executed (ShaderGlobals *sg)
ctx->incr_layers_executed ();
}
#if OSL_USE_BATCHED
// Explicit template instantiation for supported batch sizes
template class ShadingContext::Batched<16>;
template class ShadingContext::Batched<8>;
#endif
OSL_NAMESPACE_EXIT
+29 -3
View File
@@ -109,7 +109,9 @@ typedef std::shared_ptr<ShaderInstance> ShaderInstanceRef;
class Dictionary;
class RuntimeOptimizer;
class BackendLLVM;
#if OSL_USE_BATCHED
class BatchedBackendLLVM;
#endif
struct ConnectedParam;
void print_closure (std::ostream &out, const ClosureColor *closure, ShadingSystemImpl *ss);
@@ -117,27 +119,40 @@ void print_closure (std::ostream &out, const ClosureColor *closure, ShadingSyste
/// Signature of the function that LLVM generates to run the shader
/// group.
typedef void (*RunLLVMGroupFunc)(void* /* shader globals */, void*);
#if OSL_USE_BATCHED
typedef void (*RunLLVMGroupFuncWide)(void* /* batched shader globals */, void*, int run_mask_value);
#endif
/// Signature of a constant-folding method
typedef int (*OpFolder) (RuntimeOptimizer &rop, int opnum);
/// Signature of an LLVM-IR-generating method
typedef bool (*OpLLVMGen) (BackendLLVM &rop, int opnum);
#if OSL_USE_BATCHED
typedef bool (*OpLLVMGenWide) (BatchedBackendLLVM &rop, int opnum);
#endif
struct OpDescriptor {
ustring name; // name of op
OpLLVMGen llvmgen; // llvm-generating routine
#if OSL_USE_BATCHED
OpLLVMGenWide llvmgenwide; // wide version of llvm-generating routine
#endif
OpFolder folder; // constant-folding routine
bool simple_assign; // wholly overwrites arg0, no other writes,
// no side effects
int flags; // other flags
OpDescriptor () { }
OpDescriptor (const char *n, OpLLVMGen ll, OpLLVMGenWide llw,
OpDescriptor (const char *n, OpLLVMGen ll,
#if OSL_USE_BATCHED
OpLLVMGenWide llw,
#endif
OpFolder fold=NULL, bool simple=false, int flags=0)
: name(n), llvmgen(ll), llvmgenwide(llw), folder(fold),
: name(n), llvmgen(ll),
#if OSL_USE_BATCHED
llvmgenwide(llw),
#endif
folder(fold),
simple_assign(simple), flags(flags)
{}
@@ -663,6 +678,7 @@ public:
ocio_transform (StringParam fromspace, StringParam tospace,
const Color& C, Color& Cout);
#if OSL_USE_BATCHED
// Group all batched methods behind a templated interface
// so we can support multiple widths
template<int WidthT>
@@ -684,6 +700,7 @@ public:
OSL_FORCEINLINE Batched<WidthT> batched() {
return Batched<WidthT>(*this);
}
#endif
private:
void printstats () const;
@@ -914,7 +931,9 @@ private:
friend class ShaderInstance;
friend class RuntimeOptimizer;
friend class BackendLLVM;
#if OSL_USE_BATCHED
friend class BatchedBackendLLVM;
#endif
};
@@ -1511,6 +1530,7 @@ public:
m_llvm_compiled_layers[layer] = func;
}
#if OSL_USE_BATCHED
// Hold onto wide versions of llvm functions side by side with scalar
RunLLVMGroupFuncWide llvm_compiled_wide_version() const {
return m_llvm_compiled_wide_version;
@@ -1533,7 +1553,7 @@ public:
if (layer < nlayers())
m_llvm_compiled_wide_layers[layer] = func;
}
#endif
// Is this shader group equivalent to ret void?
bool does_nothing() const {
return m_does_nothing;
@@ -1621,9 +1641,11 @@ private:
RunLLVMGroupFunc m_llvm_compiled_version = nullptr;
RunLLVMGroupFunc m_llvm_compiled_init = nullptr;
std::vector<RunLLVMGroupFunc> m_llvm_compiled_layers;
#if OSL_USE_BATCHED
RunLLVMGroupFuncWide m_llvm_compiled_wide_version = nullptr;
RunLLVMGroupFuncWide m_llvm_compiled_wide_init = nullptr;
std::vector<RunLLVMGroupFuncWide> m_llvm_compiled_wide_layers;
#endif
std::vector<ShaderInstanceRef> m_layers;
ustring m_name;
int m_exec_repeat = 1; ///< How many times to execute group
@@ -1660,7 +1682,9 @@ private:
friend class OSL::pvt::ShadingSystemImpl;
friend class OSL::pvt::BackendLLVM;
#if OSL_USE_BATCHED
friend class OSL::pvt::BatchedBackendLLVM;
#endif
friend class ShadingContext;
};
@@ -1695,6 +1719,7 @@ public:
/// layer, and cleanup. (See similarly named method of ShadingSystem.)
bool execute (ShaderGroup &group, ShaderGlobals &globals, bool run=true);
#if OSL_USE_BATCHED
// Group all batched methods behind a templated interface
// so we can support multiple widths
template<int WidthT>
@@ -1762,6 +1787,7 @@ public:
OSL_FORCEINLINE Batched<WidthT> batched() {
return Batched<WidthT>(*this);
}
#endif
ClosureComponent * closure_component_allot(int id, size_t prim_size, const Color3 &w) {
// Allocate the component and the mul back to back
+2
View File
@@ -2769,6 +2769,7 @@ RuntimeOptimizer::post_optimize_instance ()
add_useparam (allsymptrs);
#if OSL_USE_BATCHED
// We must identify which symbols are uniform before
// trying to coalesce to avoid merging a varying with
// a uniform symbol or forced_llvm_bool with an integer
@@ -2776,6 +2777,7 @@ RuntimeOptimizer::post_optimize_instance ()
BatchedAnalysis batched_analysis(shadingsys(), group());
batched_analysis.analyze_layer(inst());
}
#endif
if (optimize() >= 1 && m_opt_coalesce_temps)
coalesce_temporaries ();
+64 -18
View File
@@ -12,7 +12,9 @@
#include "oslexec_pvt.h"
#include <OSL/genclosure.h>
#include "backendllvm.h"
#if OSL_USE_BATCHED
#include "batched_backendllvm.h"
#endif
#include <OSL/oslquery.h>
#include <OpenImageIO/filesystem.h>
@@ -304,6 +306,7 @@ ShadingSystem::execute_layer (ShadingContext &ctx, ShaderGlobals &globals,
return layernumber >= 0 ? ctx.execute_layer (globals, layernumber) : false;
}
#if OSL_USE_BATCHED
template<int WidthT>
bool
ShadingSystem::BatchedExecutor<WidthT>::execute (ShadingContext &ctx, ShaderGroup &group,
@@ -348,7 +351,7 @@ ShadingSystem::BatchedExecutor<WidthT>::execute_layer (ShadingContext &ctx, int
int layernumber = sym->layer();
return layernumber >= 0 ? ctx.batched<WidthT>().execute_layer (batch_size, globals_batch, layernumber) : false;
}
#endif
bool
ShadingSystem::execute_cleanup (ShadingContext &ctx)
@@ -439,8 +442,9 @@ ShadingSystem::symbol_address (const ShadingContext &ctx,
return ctx.symbol_data (*(const Symbol *)sym);
}
#if OSL_USE_BATCHED
bool
ShadingSystem::supports_batch_execution_at(int width)
ShadingSystem::configure_batch_execution_at(int width)
{
auto requestedISA = LLVM_Util::lookup_isa_by_name(m_impl->llvm_jit_target());
OSL_MAYBE_UNUSED bool target_requested = (requestedISA != TargetISA::UNKNOWN);
@@ -457,22 +461,28 @@ ShadingSystem::supports_batch_execution_at(int width)
case TargetISA::UNKNOWN:
// fallthrough
case TargetISA::AVX512:
#ifdef __OSL_SUPPORTS_B16_AVX512
if (jit_fma) {
#ifdef __OSL_SUPPORTS_b16_AVX512
if (LLVM_Util::supports_isa(TargetISA::AVX512)) {
if (!target_requested)
m_impl->attribute("llvm_jit_target", LLVM_Util::target_isa_name(TargetISA::AVX512));
return true;
}
#endif
if (target_requested) { break; }
}
#endif
// fallthrough
case TargetISA::AVX512_noFMA:
#ifdef __OSL_SUPPORTS_B16_AVX512_NOFMA
#ifdef __OSL_SUPPORTS_b16_AVX512_noFMA
if (LLVM_Util::supports_isa(TargetISA::AVX512_noFMA)) {
if (!target_requested) {
m_impl->attribute("llvm_jit_target", LLVM_Util::target_isa_name(TargetISA::AVX512_noFMA));
}
m_impl->attribute ("llvm_jit_fma", 0);
return true;
}
if (target_requested) { break; }
#endif
if (target_requested) { break; }
// fallthrough
default:
return false;
@@ -484,48 +494,63 @@ ShadingSystem::supports_batch_execution_at(int width)
case TargetISA::UNKNOWN:
// fallthrough
case TargetISA::AVX512:
#ifdef __OSL_SUPPORTS_B8_AVX512
if (jit_fma) {
#ifdef __OSL_SUPPORTS_b8_AVX512
if (LLVM_Util::supports_isa(TargetISA::AVX512)) {
if (!target_requested)
m_impl->attribute("llvm_jit_target", LLVM_Util::target_isa_name(TargetISA::AVX512));
return true;
}
#endif
if (target_requested) { break; }
}
#endif
// fallthrough
case TargetISA::AVX512_noFMA:
#ifdef __OSL_SUPPORTS_B8_AVX512_NOFMA
#ifdef __OSL_SUPPORTS_b8_AVX512_noFMA
if (LLVM_Util::supports_isa(TargetISA::AVX512_noFMA)) {
if (!target_requested) {
m_impl->attribute("llvm_jit_target", LLVM_Util::target_isa_name(TargetISA::AVX512_noFMA));
}
m_impl->attribute ("llvm_jit_fma", 0);
return true;
}
if (target_requested) { break; }
#endif
if (target_requested) { break; }
// fallthrough
case TargetISA::AVX2:
#ifdef __OSL_SUPPORTS_B8_AVX2
if (jit_fma) {
#ifdef __OSL_SUPPORTS_b8_AVX2
if (LLVM_Util::supports_isa(TargetISA::AVX2)) {
if (!target_requested)
m_impl->attribute("llvm_jit_target", LLVM_Util::target_isa_name(TargetISA::AVX2));
return true;
}
#endif
if (target_requested) { break; }
}
#endif
// fallthrough
case TargetISA::AVX2_noFMA:
#ifdef __OSL_SUPPORTS_B8_AVX2_NOFMA
#ifdef __OSL_SUPPORTS_b8_AVX2_noFMA
if (LLVM_Util::supports_isa(TargetISA::AVX2_noFMA)) {
if (!target_requested)
m_impl->attribute("llvm_jit_target", LLVM_Util::target_isa_name(TargetISA::AVX2_noFMA));
m_impl->attribute ("llvm_jit_fma", 0);
return true;
}
if (target_requested) { break; }
#endif
if (target_requested) { break; }
// fallthrough
case TargetISA::AVX:
#ifdef __OSL_SUPPORTS_B8_AVX
#ifdef __OSL_SUPPORTS_b8_AVX
if (LLVM_Util::supports_isa(TargetISA::AVX)) {
if (!target_requested)
m_impl->attribute("llvm_jit_target", LLVM_Util::target_isa_name(TargetISA::AVX));
// AVX doesn't support FMA
m_impl->attribute ("llvm_jit_fma", 0);
return true;
}
if (target_requested) { break; }
#endif
if (target_requested) { break; }
// fallthrough
default:
return false;
@@ -535,6 +560,7 @@ ShadingSystem::supports_batch_execution_at(int width)
return false;
}
}
#endif
std::string
ShadingSystem::getstats (int level) const
@@ -689,6 +715,7 @@ ShadingSystem::optimize_group (ShaderGroup *group,
optimize_group (group, ctx, do_jit);
}
#if OSL_USE_BATCHED
template<int WidthT>
void
ShadingSystem::BatchedExecutor<WidthT>::jit_group (ShaderGroup *group, ShadingContext *ctx)
@@ -707,7 +734,7 @@ ShadingSystem::BatchedExecutor<WidthT>::jit_all_groups (int nthreads)
// Explicitly instantiate
template class ShadingSystem::BatchedExecutor<16>;
template class ShadingSystem::BatchedExecutor<8>;
#endif
static TypeDesc TypeFloatArray2 (TypeDesc::FLOAT, 2);
@@ -879,8 +906,12 @@ ShadingSystemImpl::ShadingSystemImpl (RendererServices *renderer,
m_opt_fold_getattribute(true),
m_opt_middleman(true), m_opt_texture_handle(true),
m_opt_seed_bblock_aliases(true),
#if OSL_USE_BATCHED
m_opt_batched_analysis((renderer->batched(WidthOf<16>()) != nullptr) |
(renderer->batched(WidthOf<8>()) != nullptr)),
#else
m_opt_batched_analysis(false),
#endif
m_llvm_jit_fma(false),
m_llvm_jit_aggressive(false),
m_optimize_nondebug(false),
@@ -1011,6 +1042,7 @@ ShadingSystemImpl::ShadingSystemImpl (RendererServices *renderer,
static void
shading_system_setup_op_descriptors (ShadingSystemImpl::OpDescriptorMap& op_descriptor)
{
#if OSL_USE_BATCHED
#define OP2(alias,name,ll,fold,simp,flag) \
extern bool llvm_gen_##ll (BatchedBackendLLVM &rop, int opnum); \
extern bool llvm_gen_##ll (BackendLLVM &rop, int opnum); \
@@ -1018,6 +1050,14 @@ shading_system_setup_op_descriptors (ShadingSystemImpl::OpDescriptorMap& op_desc
op_descriptor[ustring(#alias)] = OpDescriptor(#name, llvm_gen_##ll, \
llvm_gen_##ll, \
constfold_##fold, simp, flag);
#else
#define OP2(alias,name,ll,fold,simp,flag) \
extern bool llvm_gen_##ll (BackendLLVM &rop, int opnum); \
extern int constfold_##fold (RuntimeOptimizer &rop, int opnum); \
op_descriptor[ustring(#alias)] = OpDescriptor(#name, llvm_gen_##ll, \
constfold_##fold, simp, flag);
#endif
#define OP(name,ll,fold,simp,flag) OP2(name,name,ll,fold,simp,flag)
#define TEX OpDescriptor::Tex
#define SIDE OpDescriptor::SideEffects
@@ -3322,6 +3362,7 @@ ShadingSystemImpl::optimize_group (ShaderGroup &group, ShadingContext *ctx, bool
m_groups_to_compile_count -= 1;
}
#if OSL_USE_BATCHED
template <int WidthT>
void
ShadingSystemImpl::Batched<WidthT>::jit_group (ShaderGroup &group, ShadingContext *ctx)
@@ -3395,18 +3436,21 @@ ShadingSystemImpl::Batched<WidthT>::jit_group (ShaderGroup &group, ShadingContex
m_ssi.m_stat_instances_compiled += group.nlayers();
m_ssi.m_groups_to_compile_count -= 1;
}
#endif
static void optimize_all_groups_wrapper (ShadingSystemImpl *ss, int mythread, int totalthreads, bool do_jit)
{
ss->optimize_all_groups (1, mythread, totalthreads, do_jit);
}
#if OSL_USE_BATCHED
template<int WidthT>
static void batched_jit_all_groups_wrapper (ShadingSystemImpl *ss, int mythread, int totalthreads)
{
ss->batched<WidthT>().jit_all_groups(1, mythread, totalthreads);
}
#endif
void
ShadingSystemImpl::optimize_all_groups (int nthreads, int mythread, int totalthreads, bool do_jit)
@@ -3452,6 +3496,7 @@ ShadingSystemImpl::optimize_all_groups (int nthreads, int mythread, int totalthr
destroy_thread_info(threadinfo);
}
#if OSL_USE_BATCHED
template<int WidthT>
void
ShadingSystemImpl::Batched<WidthT>::jit_all_groups (int nthreads, int mythread, int totalthreads)
@@ -3501,6 +3546,7 @@ ShadingSystemImpl::Batched<WidthT>::jit_all_groups (int nthreads, int mythread,
// machine as well, start with just the batch size
template class pvt::ShadingSystemImpl::Batched<16>;
template class pvt::ShadingSystemImpl::Batched<8>;
#endif
int
ShadingSystemImpl::merge_instances (ShaderGroup &group, bool post_opt)
+2 -1
View File
@@ -13,10 +13,11 @@
#include <cmath>
#include <OSL/oslconfig.h>
#include <OSL/batched_shaderglobals.h>
#include <OSL/dual.h>
#include <OSL/dual_vec.h>
#include <OSL/oslconfig.h>
#include <OSL/sfmath.h>
#include <OSL/wide.h>
+11 -2
View File
@@ -3,8 +3,17 @@
# https://github.com/AcademySoftwareFoundation/OpenShadingLanguage
# The 'testshade' executable
set ( testshade_srcs testshade.cpp simplerend.cpp batched_simplerend.cpp optixgridrender.cpp
../testrender/optix_stringtable.cpp )
set ( testshade_srcs
testshade.cpp
simplerend.cpp
optixgridrender.cpp
../testrender/optix_stringtable.cpp )
if (BUILD_BATCHED)
list(APPEND testshade_srcs
batched_simplerend.cpp
)
endif()
if (CMAKE_COMPILER_IS_INTEL)
# Make sure our shader globals are calculated consistently
+1
View File
@@ -11,6 +11,7 @@
#include <OpenImageIO/imagebuf.h>
#include <OpenImageIO/ustring.h>
#include <OSL/oslconfig.h>
#include <OSL/batched_rendererservices.h>
OSL_NAMESPACE_ENTER
+2
View File
@@ -115,8 +115,10 @@ void register_closures(OSL::ShadingSystem* shadingsys) {
SimpleRenderer::SimpleRenderer ()
#if OSL_USE_BATCHED
: m_batch_16_simple_renderer(*this)
, m_batch_8_simple_renderer(*this)
#endif
{
Matrix44 M; M.makeIdentity();
camera_params (M, u_perspective, 90.0f,
+9 -1
View File
@@ -12,8 +12,11 @@
#include <OpenImageIO/imagebuf.h>
#include <OSL/oslexec.h>
#include <OSL/rendererservices.h>
#include "batched_simplerend.h"
#if OSL_USE_BATCHED
# include "batched_simplerend.h"
#endif
OSL_NAMESPACE_ENTER
@@ -115,12 +118,17 @@ public:
OIIO::ParamValueList options;
OIIO::ParamValueList userdata;
#if OSL_USE_BATCHED
virtual BatchedRendererServices<16> * batched(WidthOf<16>) { return &m_batch_16_simple_renderer; }
virtual BatchedRendererServices<8> * batched(WidthOf<8>) { return &m_batch_8_simple_renderer; }
#endif
protected:
#if OSL_USE_BATCHED
BatchedSimpleRenderer<16> m_batch_16_simple_renderer;
BatchedSimpleRenderer<8> m_batch_8_simple_renderer;
#endif
// Camera parameters
Matrix44 m_world_to_camera;
+23 -9
View File
@@ -30,7 +30,9 @@
#include <OSL/oslexec.h>
#include <OSL/oslcomp.h>
#include <OSL/oslquery.h>
#include <OSL/batched_shaderglobals.h>
#if OSL_USE_BATCHED
# include <OSL/batched_shaderglobals.h>
#endif
#include "optixgridrender.h"
#include "simplerend.h"
@@ -193,20 +195,21 @@ set_shadingsys_options ()
batch_size = atoi(opt_env);
if (batched) {
// We are only building FMA versions, so force it on
#if OSL_USE_BATCHED
// Allow FMA if build of OSL supports it
shadingsys->attribute ("llvm_jit_fma", 1);
bool batch_size_requested = (batch_size != -1);
// Not really looping, just emulating goto behavior using break
for(;;) {
if (!batch_size_requested || batch_size == 16) {
if (shadingsys->supports_batch_execution_at(16)) {
if (shadingsys->configure_batch_execution_at(16)) {
batch_size = 16;
break;
}
}
if (!batch_size_requested || batch_size == 8) {
if (shadingsys->supports_batch_execution_at(8)) {
if (shadingsys->configure_batch_execution_at(8)) {
batch_size = 8;
break;
}
@@ -225,6 +228,9 @@ set_shadingsys_options ()
batched = false;
break;
}
#else
batched = false;
#endif
}
if (!batched) {
@@ -864,6 +870,7 @@ setup_output_images (SimpleRenderer *rend, ShadingSystem *shadingsys,
// We also choose to JIT it now during timing for setup
OSL::PerThreadInfo *thread_info = shadingsys->create_thread_info();
ShadingContext *ctx = shadingsys->get_context(thread_info);
#if OSL_USE_BATCHED
if (batched) {
// jit_group will optimize the group if necesssary
if (batch_size == 16) {
@@ -872,7 +879,9 @@ setup_output_images (SimpleRenderer *rend, ShadingSystem *shadingsys,
ASSERT((batch_size == 8) && "Unsupport batch size");
shadingsys->batched<8>().jit_group (shadergroup.get(), ctx);
}
} else {
} else
#endif
{
shadingsys->optimize_group (shadergroup.get(), ctx, true /*do_jit*/);
}
@@ -1006,6 +1015,7 @@ save_outputs (SimpleRenderer *rend, ShadingSystem *shadingsys,
}
}
#if OSL_USE_BATCHED
// For batch of pixels (bx[WidthT], by[WidthT]) that was just shaded
// by the given shading context, save each of the requested outputs
// to the corresponding output ImageBuf.
@@ -1147,8 +1157,8 @@ batched_save_outputs (SimpleRenderer *rend, ShadingSystem *shadingsys, ShadingCo
std::cout << oStreams[batchIndex]->str();
}
}
}
#endif
static void
test_group_attributes (ShaderGroup *group)
@@ -1318,7 +1328,8 @@ shade_region (SimpleRenderer *rend, ShaderGroup *shadergroup,
}
// Set up the ShaderGlobals fields for pixel (x,y).
#if OSL_USE_BATCHED
// Set up the uniform portion of BatchedShaderGlobals fields
template<int WidthT>
static void
setup_uniform_shaderglobals (BatchedShaderGlobals<WidthT> &bsg, ShadingSystem *shadingsys)
@@ -1539,7 +1550,7 @@ batched_shade_region (SimpleRenderer *rend, ShaderGroup *shadergroup, OIIO::ROI
shadingsys->release_context (ctx);
shadingsys->destroy_thread_info(thread_info);
}
#endif
static void synchio() {
// Synch all writes to stdout & stderr now (mostly for Windows)
@@ -1793,6 +1804,7 @@ test_shade (int argc, const char *argv[])
#if 0
shade_region (rend, shadergroup.get(), roi, save);
#else
#if OSL_USE_BATCHED
if (batched) {
if (batch_size == 16) {
OIIO::ImageBufAlgo::parallel_image (roi, num_threads,
@@ -1806,7 +1818,9 @@ test_shade (int argc, const char *argv[])
batched_shade_region<8> (rend, shadergroup.get(), sub_roi, save);
});
}
} else {
} else
#endif
{
OIIO::ImageBufAlgo::parallel_image (roi, num_threads,
[&](OIIO::ROI sub_roi)->void {
shade_region (rend, shadergroup.get(), sub_roi, save);
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.
Binary file not shown.

Some files were not shown because too many files have changed in this diff Show More