BRL-CAD/ospray
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Start running OSPRay as DPC++ on the GPU

Browse Source 
Use an enum + switch statements to avoid issues/limitations with
function ptrs.

Renable volumes with OpenVKL when running on CPU w/ enable volumes

Put both CPU and GPU modules behind CMake flags

workaround odd bug with AO renderer and intervals

Add spheres geometry and port to use Embree sphere point

Push math constants for RNG to USM instead of in global constant

WA for `Material_dispatch_getTransparency` hang, working path tracer

Re-enable subdivision surfaces for CPU-only

Port Texture2D to use switch statement dispatch

Migrate all materials to switch dispatch

Add option for large grf, does improve performance

Workaround for TextureParam get3f bug

Update headers and APIs for L0 interop header updates/deprecations
This commit is contained in:
Will Usher
2023-01-11 13:07:49 +01:00
committed by Miroslaw Pawlowski
parent 73f4c42d9b
commit 078c589aaa
376 changed files with 8097 additions and 1932 deletions
Download Patch File Download Diff File Expand all files Collapse all files
.gitignore
+1
View File
@@ -35,3 +35,4 @@ premake.local.*
cmake-build*/
.clangd
compile_commands.json
.cache/
CMakeLists.txt
+1 -1
View File
@@ -11,7 +11,7 @@ set(CMAKE_DISABLE_SOURCE_CHANGES ON)
set(CMAKE_DISABLE_IN_SOURCE_BUILD ON)
set(CMAKE_C_STANDARD 99)
set(CMAKE_CXX_STANDARD 11)
set(CMAKE_CXX_STANDARD 17)
set(CMAKE_C_STANDARD_REQUIRED ON)
set(CMAKE_CXX_STANDARD_REQUIRED ON)
apps/common/ospray_testing/builders/CornellBox.cpp
+8 -1
View File
@@ -229,9 +229,16 @@ cpp::Group CornellBox::buildGroup() const
{
cpp::Geometry quadMesh("mesh");
// TODO: Temporarily try triangulating the scene
std::vector<vec3ui> triangleIndices;
for (const auto &idx : indices) {
triangleIndices.push_back(vec3ui(idx.x, idx.y, idx.z));
triangleIndices.push_back(vec3ui(idx.z, idx.w, idx.x));
}
quadMesh.setParam("vertex.position", cpp::CopiedData(vertices));
quadMesh.setParam("vertex.color", cpp::CopiedData(colors));
quadMesh.setParam("index", cpp::CopiedData(indices));
quadMesh.setParam("index", cpp::CopiedData(triangleIndices));
quadMesh.commit();
cpp::GeometricModel model(quadMesh);
apps/ospExamples/GLFWOSPRayWindow.cpp
+1 -1
View File
@@ -565,7 +565,7 @@ void GLFWOSPRayWindow::buildUI()
rendererTypeStr = g_renderers[whichRenderer];
if (rendererType == OSPRayRendererType::DEBUGGER)
whichDebuggerType = 0; // reset UI if switching away from debug renderer
whichDebuggerType = 1; // reset UI if switching away from debug renderer
if (rendererTypeStr == "scivis")
rendererType = OSPRayRendererType::SCIVIS;
apps/ospExamples/GLFWOSPRayWindow.h
+3 -3
View File
@@ -100,12 +100,12 @@ class GLFWOSPRayWindow
float turbidity{3.f};
float horizonExtension{0.1f};
std::string scene{"boxes_lit"};
std::string scene{"cornell_box"};
std::string curveVariant{"bspline"};
OSPRayRendererType rendererType{OSPRayRendererType::SCIVIS};
std::string rendererTypeStr{"scivis"};
OSPRayRendererType rendererType{OSPRayRendererType::PATHTRACER};
std::string rendererTypeStr{"pathtracer"};
std::string pixelFilterTypeStr{"gaussian"};
apps/ospTestSuite/CMakeLists.txt
+7
View File
@@ -29,6 +29,13 @@ PUBLIC
$<BUILD_INTERFACE:${CMAKE_CURRENT_LIST_DIR}>
)
if (OSPRAY_ENABLE_VOLUMES)
target_compile_definitions(ospray_gtest_utils
PUBLIC
-DOSPRAY_ENABLE_VOLUMES
)
endif()
add_executable(ospTestSuite
${OSPRAY_RESOURCE}
test_geometry.cpp
apps/ospTestSuite/test_appearance.cpp
+2
View File
@@ -425,11 +425,13 @@ void PTBackgroundRefraction::SetUp()
// Test Instantiations //////////////////////////////////////////////////////
#if 0
INSTANTIATE_TEST_SUITE_P(Transparency,
FromOsprayTesting,
::testing::Combine(::testing::Values("transparency"),
::testing::Values("scivis", "pathtracer", "ao"),
::testing::Values(16)));
#endif
TEST_P(RendererMaterialList, material_list)
{
apps/ospTestSuite/test_enums.cpp
+5 -2
View File
@@ -2,10 +2,12 @@
// SPDX-License-Identifier: Apache-2.0
#include <gtest/gtest.h>
#include "embree3/rtcore.h"
#include "openvkl/openvkl.h"
#include "embree4/rtcore.h"
#include "ospray/OSPEnums.h"
#ifdef OSPRAY_ENABLE_VOLUMES
#include "openvkl/openvkl.h"
TEST(Enums, VKLLogLevel)
{
ASSERT_LE(sizeof(OSPLogLevel), sizeof(VKLLogLevel));
@@ -128,6 +130,7 @@ TEST(Enums, VKLError)
ASSERT_EQ(OSP_OUT_OF_MEMORY, VKL_OUT_OF_MEMORY);
ASSERT_EQ(OSP_UNSUPPORTED_CPU, VKL_UNSUPPORTED_CPU);
}
#endif
TEST(Enums, RTCSubdivisionMode)
{
apps/ospTestSuite/test_geometry.cpp
+21 -7
View File
@@ -172,14 +172,20 @@ TEST_P(FromOsprayTesting, test_scenes)
INSTANTIATE_TEST_SUITE_P(TestScenesGeometry,
FromOsprayTesting,
::testing::Combine(::testing::Values("cornell_box",
#ifdef OSPRAY_ENABLE_VOLUMES
"gravity_spheres_isosurface",
#endif
"empty",
"random_spheres",
"streamlines",
"subdivision_cube",
"planes",
"planes"
#ifdef OSPRAY_ENABLE_VOLUMES
,
"unstructured_volume_isosurface",
"instancing"),
"instancing"
#endif
),
::testing::Values("scivis", "pathtracer", "ao"),
::testing::Values(16)));
@@ -223,9 +229,13 @@ INSTANTIATE_TEST_SUITE_P(TestScenesClipping,
"clip_with_meshes",
"clip_with_subdivisions",
"clip_with_linear_curves",
"clip_with_bspline_curves",
"clip_with_bspline_curves"
#ifdef OSPRAY_ENABLE_VOLUMES
,
"clip_gravity_spheres_volume",
"clip_perlin_noise_volumes"),
"clip_perlin_noise_volumes"
#endif
),
::testing::Values("scivis", "pathtracer", "ao"),
::testing::Values(16)));
@@ -241,9 +251,13 @@ TEST_P(FromOsprayTestingMaxDepth, test_scenes)
INSTANTIATE_TEST_SUITE_P(TestScenesMaxDepth,
FromOsprayTestingMaxDepth,
::testing::Combine(
::testing::Values(
"cornell_box", "clip_with_spheres", "clip_gravity_spheres_volume"),
::testing::Combine(::testing::Values("cornell_box",
"clip_with_spheres"
#ifdef OSPRAY_ENABLE_VOLUMES
,
"clip_gravity_spheres_volume"
#endif
),
::testing::Values("ao"),
::testing::Values(16)));
apps/ospTestSuite/test_motionblur.cpp
+2 -1
View File
@@ -290,6 +290,7 @@ void MotionCamera::SetUp()
// Test Instantiations //////////////////////////////////////////////////////
#if 0
TEST_P(MotionBlurBoxes, instance_mb)
{
PerformRenderTest();
@@ -318,7 +319,7 @@ INSTANTIATE_TEST_SUITE_P(CameraStereoRollingShutter,
::testing::Values(OSP_STEREO_TOP_BOTTOM),
::testing::Values(OSP_SHUTTER_ROLLING_DOWN),
::testing::Values(0.f)));
#endif
TEST_P(MotionCamera, camera_mb)
{
PerformRenderTest();
apps/ospTestSuite/test_volumetric.cpp
+2
View File
@@ -335,6 +335,7 @@ void UnstructuredVolume::SetUp()
// Test Instantiations //////////////////////////////////////////////////////
#ifdef OSPRAY_ENABLE_VOLUMES
INSTANTIATE_TEST_SUITE_P(TestScenesVolumes,
FromOsprayTesting,
::testing::Combine(::testing::Values("gravity_spheres_volume",
@@ -391,5 +392,6 @@ INSTANTIATE_TEST_SUITE_P(Renderers,
vec4f(1.f),
vec4f(0.f, 0.f, 0.f, 1.f),
vec4f(1.f, 0.f, 0.f, 0.5f))));
#endif
} // namespace OSPRayTestScenes
cmake/compiler/dpcpp.cmake
+98 -4
View File
@@ -62,7 +62,101 @@ if (APPLE)
endif()
# enable -static-intel
set(CMAKE_EXE_LINKER_FLAGS "-static-intel ${CMAKE_EXE_LINKER_FLAGS}")
set(CMAKE_SHARED_LINKER_FLAGS "-static-intel ${CMAKE_SHARED_LINKER_FLAGS}")
set(CMAKE_STATIC_LINKER_FLAGS "-static-intel ${CMAKE_STATIC_LINKER_FLAGS}")
set(CMAKE_MODULE_LINKER_FLAGS "-static-intel ${CMAKE_MODULE_LINKER_FLAGS}")
# Note: this doesn't seem to be used/relevant to the DPC++ compiler? Maybe just the
# DPCPP release one? The nightly clang compiler doesn't recognize this flag
#set(CMAKE_EXE_LINKER_FLAGS "-static-intel ${CMAKE_EXE_LINKER_FLAGS}")
#set(CMAKE_SHARED_LINKER_FLAGS "-static-intel ${CMAKE_SHARED_LINKER_FLAGS}")
#set(CMAKE_STATIC_LINKER_FLAGS "-static-intel ${CMAKE_STATIC_LINKER_FLAGS}")
#set(CMAKE_MODULE_LINKER_FLAGS "-static-intel ${CMAKE_MODULE_LINKER_FLAGS}")
# TODO: These flags aren't recognized by IntelLLVM (the nightly dpcpp compiler)
# but at least the nan one does show up in the released dpcpp compiler.
# Are they being removed? or added?
#set(CMAKE_CXX_FLAGS "-fhonor-nan-compares -fhonor-infinities ${CMAKE_CXX_FLAGS}")
# SYCL flags to match Embree
list(APPEND OSPRAY_CXX_FLAGS_SYCL
-Wno-mismatched-tags
-Wno-pessimizing-move
-Wno-reorder
-Wno-unneeded-internal-declaration
-Wno-delete-non-abstract-non-virtual-dtor
-Wno-dangling-field
-Wno-unknown-pragmas
-Wno-logical-op-parentheses
-fsycl
-fsycl-unnamed-lambda
-Xclang -fsycl-allow-func-ptr)
# FIXME: debug information generation takes forever in SYCL
# TODO: Still true?
list(APPEND OSPRAY_CXX_FLAGS_SYCL -g0)
# FIXME: assertion still not working in SYCL
# TODO: Still true?
#list(APPEND OSPRAY_CXX_FLAGS_SYCL -UDEBUG -DNDEBUG)
# IGC options from Embree
# Enable __noinline
list(APPEND OSPRAY_IGC_OPTIONS "EnableOCLNoInlineAttr=0")
# This works around some IGC bug in spill compression
# TODO: Still true?
list(APPEND OSPRAY_IGC_OPTIONS "VISAOptions=-scratchAllocForStackInKB 128 -nospillcompression")
# Allow printf inside indirectly callable function, right now I have this in all for testing
# TODO: Should only enable for debug builds, and this needs to be done using a generator expression
# if we want to support it in VS
list(APPEND OSPRAY_IGC_OPTIONS "ForceInlineStackCallWithImplArg=0" "EnableGlobalStateBuffer=1")
# This significantly improves compile times on 17028 and up, though also impacts performance some
option(OSPRAY_IGC_FAST_COMPILE
"Pass flags to improve compilation speed at the cost of some optimization" OFF)
if (OSPRAY_IGC_FAST_COMPILE)
list(APPEND OSPRAY_IGC_OPTIONS "PartitionUnit=1")
list(APPEND OSPRAY_IGC_OPTIONS "UnitSizeThreshold=20000")
endif()
#list(APPEND OSPRAY_IGC_OPTIONS "ForceOCLSIMDWidth=8")
# Development option to dump shaders, when we compile AOT this has to be done at build time
option(OSPRAY_IGC_DUMP_SHADERS "Dump IGC shaders during build" OFF)
if (OSPRAY_IGC_DUMP_SHADERS)
list(APPEND OSPRAY_IGC_OPTIONS
"ShaderDumpEnable=1"
"ShowFullVectorsInShaderDumps=1"
"DumpToCustomDir=${PROJECT_BINARY_DIR}/ospray_igc_shader_dump")
endif()
# enables support for buffers larger than 4GB
list(APPEND OSPRAY_OCL_OPTIONS -cl-intel-greater-than-4GB-buffer-required)
list(APPEND OSPRAY_OCL_OTHER_OPTIONS
-cl-intel-force-global-mem-allocation
-cl-intel-no-local-to-generic)
# Large GRF mode
option(OSPRAY_DPCPP_LARGEGRF "Enable DPC++ Large GRF Support" OFF)
if (OSPRAY_DPCPP_LARGEGRF)
list(APPEND CMAKE_OCL_OPTIONS "-internal_options -cl-intel-256-GRF-per-thread")
endif()
# SYCL options for AOT/JIT compilation
set(OSPRAY_DPCPP_AOT_DEVICES "none" CACHE STRING
"SYCL devices to use for AOT compilation. Selecting none will compile to SPV for JIT compilation")
set_property(CACHE OSPRAY_DPCPP_AOT_DEVICES PROPERTY STRINGS none dg2 pvc)
# TODO: Is this revision info going to be visible to end users? In the end the public release
# of the code should probably just have one revision it targets right? The final consumer release rev.
set(OSPRAY_DPCPP_AOT_DEVICE_REVISION "8" CACHE STRING "AOT target device revision")
string(REPLACE ";" "," OSPRAY_IGC_OPTIONS_STR "${OSPRAY_IGC_OPTIONS}")
set(OSPRAY_OCL_OPTIONS_STR "${OSPRAY_OCL_OPTIONS}")
string(REPLACE ";" " " OSPRAY_OCL_OPTIONS_STR "${OSPRAY_OCL_OPTIONS}")
set(OSPRAY_OCL_OTHER_OPTIONS_STR "${OSPRAY_OCL_OTHER_OPTIONS}")
string(REPLACE ";" " " OSPRAY_OCL_OTHER_OPTIONS_STR "${OSPRAY_OCL_OTHER_OPTIONS}")
if (OSPRAY_DPCPP_AOT_DEVICES STREQUAL "none")
set(OSPRAY_SYCL_TARGET spir64)
else()
set(OSPRAY_SYCL_TARGET spir64_gen)
endif()
cmake/ospray_macros.cmake
+16 -3
View File
@@ -331,7 +331,11 @@ macro(ospray_configure_compiler)
elseif ("${CMAKE_CXX_COMPILER_ID}" STREQUAL "Clang" OR
"${CMAKE_CXX_COMPILER_ID}" STREQUAL "AppleClang")
set(OSPRAY_COMPILER_CLANG TRUE)
include(clang)
#include(clang)
# TODO WILL: We need to distinguish "clang" being real clang
# and "clang" being dpcpp nightly
message(WARNING "TODO WILL: We need to distinguish 'clang' = real clang vs. nightly dpcpp")
include(dpcpp)
elseif ("${CMAKE_CXX_COMPILER_ID}" STREQUAL "MSVC")
set(OSPRAY_COMPILER_MSVC TRUE)
include(msvc)
@@ -377,10 +381,12 @@ endfunction()
function(ospray_verify_embree_features)
ospray_check_embree_feature(ISPC_SUPPORT ISPC)
ospray_check_embree_feature(FILTER_FUNCTION "intersection filter")
# TODO: Having Function pointers enabled has a massive performance impact on GPU
#ospray_check_embree_feature(FILTER_FUNCTION "intersection filter")
ospray_check_embree_feature(GEOMETRY_TRIANGLE "triangle geometries")
ospray_check_embree_feature(GEOMETRY_CURVE "spline curve geometries")
ospray_check_embree_feature(GEOMETRY_USER "user geometries")
# TODO: Having Function pointers enabled has a massive performance impact on GPU
#ospray_check_embree_feature(GEOMETRY_USER "user geometries")
ospray_check_embree_feature(RAY_PACKETS "ray packets")
ospray_check_embree_feature(BACKFACE_CULLING "backface culling" OFF)
endfunction()
@@ -395,12 +401,19 @@ macro(ospray_find_embree EMBREE_VERSION_REQUIRED)
" set the 'embree_DIR' variable to the installation (or build)"
" directory.")
endif()
# Get Embree CPU info
get_target_property(EMBREE_INCLUDE_DIRS embree
INTERFACE_INCLUDE_DIRECTORIES)
get_target_property(CONFIGURATIONS embree IMPORTED_CONFIGURATIONS)
list(GET CONFIGURATIONS 0 CONFIGURATION)
get_target_property(EMBREE_LIBRARY embree
IMPORTED_LOCATION_${CONFIGURATION})
# Get Embree SYCL info
get_target_property(CONFIGURATIONS embree_sycl IMPORTED_CONFIGURATIONS)
list(GET CONFIGURATIONS 0 CONFIGURATION)
get_target_property(EMBREE_SYCL_LIBRARY embree_sycl
IMPORTED_LOCATION_${CONFIGURATION})
message(STATUS "Found Embree v${embree_VERSION}: ${EMBREE_LIBRARY}")
endmacro()
cmake/ospray_options.cmake
+4 -2
View File
@@ -13,7 +13,7 @@ set(OSPRAY_CMAKECONFIG_DIR
set(ISPC_VERSION_REQUIRED 1.18.0)
set(RKCOMMON_VERSION_REQUIRED 1.10.0)
set(EMBREE_VERSION_REQUIRED 3.13.1)
set(EMBREE_VERSION_REQUIRED 4.0.0)
set(OPENVKL_VERSION_REQUIRED 1.3.0)
set(CMAKE_RUNTIME_OUTPUT_DIRECTORY ${PROJECT_BINARY_DIR})
@@ -38,7 +38,9 @@ ospray_verify_embree_features()
ospray_determine_embree_isa_support()
# Open VKL
ospray_find_openvkl(${OPENVKL_VERSION_REQUIRED})
if (OSPRAY_ENABLE_VOLUMES)
ospray_find_openvkl(${OPENVKL_VERSION_REQUIRED})
endif()
# OpenImageDenoise
if (OSPRAY_MODULE_DENOISER)
modules/cpu/CMakeLists.txt
+321 -126
View File
@@ -10,15 +10,37 @@ cmake_dependent_option(
)
mark_as_advanced(OSPRAY_MODULE_CPU)
if (NOT OSPRAY_MODULE_CPU)
cmake_dependent_option(
OSPRAY_MODULE_GPU
"Build the GPU module"
ON
OSPRAY_ENABLE_MODULES
OFF
)
mark_as_advanced(OSPRAY_MODULE_GPU)
if ((NOT OSPRAY_MODULE_CPU) AND (NOT OSPRAY_MODULE_GPU))
return()
endif()
cmake_dependent_option(
OSPRAY_ENABLE_VOLUMES
"Enable volume rendering using OpenVKL in the CPU module"
ON
OSPRAY_MODULE_CPU
OFF
)
include(${ISPCRT_DIR}/interop.cmake)
if (NOT "${CMAKE_CXX_COMPILER}" STREQUAL "${DPCPP_COMPILER}")
message(FATAL_ERROR "CMAKE_CXX_COMPILER should be set to '${DPCPP_COMPILER}' for DPC++")
endif()
# -------------------------------------------------------
# Setup module include directories and source files
# -------------------------------------------------------
set(OSPRAY_ISPC_DEVICE_SOURCES
set(OSPRAY_CPP_SOURCES
${OSPRAY_RESOURCE}
ModuleInit.cpp
@@ -26,27 +48,19 @@ set(OSPRAY_ISPC_DEVICE_SOURCES
ISPCDevice.h
ISPCDevice.cpp
ISPCDevice.ispc
common/Data.ih
common/Data.cpp
common/Future.cpp
common/Group.cpp
common/Instance.cpp
common/MotionTransform.cpp
common/Managed.cpp
common/OSPCommon.ispc
common/World.ih
common/World.ispc
common/World.cpp
fb/FrameBuffer.ispc
fb/FrameBuffer.cpp
fb/FrameBufferView.cpp
fb/LocalFB.ispc
fb/LocalFB.cpp
fb/SparseFB.cpp
fb/SparseFB.ispc
fb/ImageOp.cpp
fb/TaskError.cpp
fb/registration.cpp
@@ -55,154 +69,192 @@ set(OSPRAY_ISPC_DEVICE_SOURCES
fb/frame_ops/Debug.cpp
fb/frame_ops/Depth.cpp
fb/frame_ops/SSAO.cpp
fb/frame_ops/SSAO.ispc
fb/pixel_ops/ToneMapper.cpp
fb/pixel_ops/ToneMapper.ispc
pf/PixelFilter.ih
pf/PixelFilter.ispc
pf/LUTPixelFilter.ih
pf/LUTPixelFilter.ispc
pf/PixelFilter.cpp
camera/Camera.cpp
camera/PerspectiveCamera.ispc
camera/PerspectiveCamera.cpp
camera/OrthographicCamera.ispc
camera/OrthographicCamera.cpp
camera/PanoramicCamera.ispc
camera/PanoramicCamera.cpp
camera/registration.cpp
geometry/Geometry.cpp
geometry/GeometricModel.ispc
geometry/GeometricModel.cpp
geometry/Boxes.ispc
geometry/Boxes.cpp
geometry/Subdivision.ispc
geometry/Subdivision.cpp
geometry/Curves.cpp
geometry/Curves.ispc
geometry/Spheres.cpp
geometry/Spheres.ispc
geometry/Isosurfaces.ispc
geometry/Isosurfaces.cpp
geometry/Planes.ispc
geometry/Planes.cpp
geometry/Mesh.cpp
geometry/Mesh.ispc
geometry/registration.cpp
lights/Light.ispc
lights/Light.cpp
lights/AmbientLight.ispc
lights/AmbientLight.cpp
lights/DirectionalLight.cpp
lights/DirectionalLight.ispc
lights/PointLight.cpp
lights/PointLight.ispc
lights/IntensityDistribution.cpp
lights/SpotLight.cpp
lights/SpotLight.ispc
lights/QuadLight.cpp
lights/QuadLight.ispc
lights/CylinderLight.cpp
lights/CylinderLight.ispc
lights/HDRILight.cpp
lights/HDRILight.ispc
lights/SunSkyLight.cpp
lights/sky_model/sky_model.cpp
lights/registration.cpp
math/Distribution1D.ispc
math/Distribution2D.ispc
math/sobol.ih
math/sobol.ispc
math/spectrum.h
math/spectrum.ih
math/spectrum.ispc
math/random.ih
math/halton.ih
math/halton.ispc
math/halton.cpp
math/MathConstants.cpp
render/LoadBalancer.cpp
render/Material.cpp
render/Material.ispc
render/Renderer.ispc
render/Renderer.cpp
render/util.ispc
render/registration.cpp
render/debug/DebugRenderer.cpp
render/debug/DebugRenderer.ispc
render/ao/AORenderer.cpp
render/ao/AORenderer.ispc
render/ao/surfaces.ispc
render/ao/volumes.ispc
render/scivis/SciVis.cpp
render/scivis/SciVis.ispc
render/scivis/SciVisData.cpp
render/pathtracer/PathTracer.cpp
render/pathtracer/PathTracerData.cpp
render/materials/OBJ.cpp
render/materials/Principled.cpp
render/materials/CarPaint.cpp
render/materials/Velvet.cpp
render/materials/Metal.cpp
render/materials/Alloy.cpp
render/materials/ThinGlass.cpp
render/materials/Glass.cpp
render/materials/MetallicPaint.cpp
render/materials/Plastic.cpp
render/materials/Luminous.cpp
render/materials/Mix.cpp
texture/Texture.cpp
texture/Texture2D.cpp
texture/registration.cpp
)
set(OSPRAY_ISPC_SOURCES
ISPCDevice.ispc
common/OSPCommon.ispc
common/World.ispc
fb/FrameBuffer.ispc
fb/FrameBufferType.ispc
fb/LocalFB.ispc
fb/SparseFB.ispc
fb/frame_ops/SSAO.ispc
fb/pixel_ops/ToneMapper.ispc
pf/PixelFilter.ispc
pf/LUTPixelFilter.ispc
camera/CameraType.ispc
camera/PerspectiveCamera.ispc
camera/OrthographicCamera.ispc
camera/PanoramicCamera.ispc
geometry/GeometryType.ispc
geometry/GeometricModel.ispc
geometry/Boxes.ispc
geometry/Curves.ispc
geometry/Spheres.ispc
geometry/Planes.ispc
geometry/Mesh.ispc
lights/LightType.ispc
lights/Light.ispc
lights/AmbientLight.ispc
lights/DirectionalLight.ispc
lights/PointLight.ispc
lights/SpotLight.ispc
lights/QuadLight.ispc
lights/CylinderLight.ispc
lights/HDRILight.ispc
math/Distribution1D.ispc
math/Distribution2D.ispc
math/sobol.ispc
math/spectrum.ispc
render/Material.ispc
render/MaterialType.ispc
render/Renderer.ispc
render/RendererType.ispc
render/util.ispc
render/debug/DebugRenderer.ispc
render/ao/AORenderer.ispc
render/ao/surfaces.ispc
render/scivis/SciVis.ispc
render/scivis/surfaces.ispc
render/scivis/volumes.ispc
render/scivis/lightAlpha.ispc
render/pathtracer/PathTracer.ispc
render/pathtracer/PathTracer.cpp
render/pathtracer/PathTracerData.cpp
render/pathtracer/PathSampler.ispc
render/pathtracer/GeometryLight.ispc
render/pathtracer/VirtualLight.ispc
render/pathtracer/TransparentShadow.ispc
render/pathtracer/ShadowCatcher.ispc
render/pathtracer/NextEventEstimation.ispc
render/pathtracer/volumes/VolumeSampler.ispc
render/bsdfs/BSDF.ispc
render/bsdfs/MicrofacetAlbedoTables.ispc
render/materials/OBJ.ispc
render/materials/OBJ.cpp
render/materials/Principled.ispc
render/materials/Principled.cpp
render/materials/CarPaint.ispc
render/materials/CarPaint.cpp
render/materials/Velvet.ispc
render/materials/Velvet.cpp
render/materials/Metal.ispc
render/materials/Metal.cpp
render/materials/Alloy.ispc
render/materials/Alloy.cpp
render/materials/ThinGlass.ispc
render/materials/ThinGlass.cpp
render/materials/Glass.ispc
render/materials/Glass.cpp
render/materials/MetallicPaint.ispc
render/materials/MetallicPaint.cpp
render/materials/Plastic.ispc
render/materials/Plastic.cpp
render/materials/Luminous.ispc
render/materials/Luminous.cpp
render/materials/Mix.ispc
render/materials/Mix.cpp
texture/Texture.cpp
texture/Texture2D.cpp
texture/Texture2D.ispc
texture/TextureVolume.cpp
texture/TextureVolume.ispc
texture/registration.cpp
volume/Volume.ispc
volume/Volume.cpp
volume/VolumetricModel.cpp
volume/transferFunction/LinearTransferFunction.ispc
volume/transferFunction/LinearTransferFunction.cpp
volume/transferFunction/TransferFunction.cpp
volume/transferFunction/registration.cpp
texture/TextureType.ispc
)
if (OSPRAY_ENABLE_VOLUMES)
set(OSPRAY_VOLUMES_CPP_SOURCES
geometry/Isosurfaces.cpp
texture/TextureVolume.cpp
volume/Volume.cpp
volume/VolumetricModel.cpp
volume/transferFunction/LinearTransferFunction.cpp
volume/transferFunction/TransferFunction.cpp
volume/transferFunction/registration.cpp
)
set(OSPRAY_VOLUMES_ISPC_SOURCES
geometry/Isosurfaces.ispc
render/ao/volumes.ispc
render/scivis/volumes.ispc
render/pathtracer/volumes/VolumeSampler.ispc
texture/TextureVolume.ispc
volume/Volume.ispc
volume/transferFunction/LinearTransferFunction.ispc
)
endif()
##############################################################
# Build ISPC device
##############################################################
@@ -271,54 +323,196 @@ install(FILES ${CMAKE_BINARY_DIR}/OSPConfig.h
## Build the module ##
include_directories_ispc(
${PROJECT_SOURCE_DIR}/ospray/include
${PROJECT_SOURCE_DIR}/ospray
${PROJECT_SOURCE_DIR}
${PROJECT_BINARY_DIR}
${RKCOMMON_INCLUDE_DIRS}
${EMBREE_INCLUDE_DIRS}
${OPENVKL_INCLUDE_DIRS}
)
add_library(ospray_module_cpu SHARED)
if (OSPRAY_MODULE_CPU)
include_directories_ispc(
${PROJECT_SOURCE_DIR}/ospray/include
${PROJECT_SOURCE_DIR}/ospray
${PROJECT_SOURCE_DIR}
${PROJECT_BINARY_DIR}
${RKCOMMON_INCLUDE_DIRS}
${EMBREE_INCLUDE_DIRS}
)
ispc_target_add_sources(ospray_module_cpu ${OSPRAY_ISPC_DEVICE_SOURCES})
add_definitions_ispc(
-DOSPRAY_BEGIN_ISPC_NAMESPACE=
-DOSPRAY_END_ISPC_NAMESPACE=
-DSYCL_EXTERNAL=
-D__noinline=
)
ospray_install_library(ospray_module_cpu lib)
if (OSPRAY_ENABLE_VOLUMES)
include_directories_ispc(
${OPENVKL_INCLUDE_DIRS}
)
target_link_libraries(ospray_module_cpu
PUBLIC
ospray
rkcommon::rkcommon
$<BUILD_INTERFACE:ispcrt::ispcrt>
$<BUILD_INTERFACE:embree>
$<BUILD_INTERFACE:openvkl::openvkl>
)
add_definitions_ispc(
-DOSPRAY_ENABLE_VOLUMES
)
endif()
target_include_directories(ospray_module_cpu
add_library(ospray_module_cpu SHARED)
ispc_target_add_sources(ospray_module_cpu
${OSPRAY_CPP_SOURCES}
${OSPRAY_VOLUMES_CPP_SOURCES}
${OSPRAY_ISPC_SOURCES}
${OSPRAY_VOLUMES_ISPC_SOURCES}
geometry/Subdivision.cpp
geometry/Subdivision.ispc)
ospray_install_library(ospray_module_cpu lib)
target_link_libraries(ospray_module_cpu
PUBLIC
$<BUILD_INTERFACE:${CMAKE_CURRENT_SOURCE_DIR}>
$<BUILD_INTERFACE:${CMAKE_SOURCE_DIR}/ospray/api>
$<BUILD_INTERFACE:${CMAKE_SOURCE_DIR}/ospray/common>
$<BUILD_INTERFACE:${CMAKE_SOURCE_DIR}/ospray>
$<BUILD_INTERFACE:${CMAKE_CURRENT_BINARY_DIR}>
$<BUILD_INTERFACE:${PROJECT_BINARY_DIR}>
# NOTE(jda) - the following includes are missing despite PUBLIC linking
$<BUILD_INTERFACE:${ISPCRT_INCLUDE_DIR}>
$<BUILD_INTERFACE:${OPENVKL_INCLUDE_DIRS}>
$<BUILD_INTERFACE:${EMBREE_INCLUDE_DIRS}>
#######################################################################
$<INSTALL_INTERFACE:${CMAKE_INSTALL_INCLUDEDIR}/ospray/SDK>
$<INSTALL_INTERFACE:${CMAKE_INSTALL_INCLUDEDIR}>
)
ospray
rkcommon::rkcommon
$<BUILD_INTERFACE:ispcrt::ispcrt>
$<BUILD_INTERFACE:embree>
)
target_compile_definitions(ospray_module_cpu
PUBLIC
TILE_SIZE=${OSPRAY_TILE_SIZE}
MAX_TILE_SIZE=${OSPRAY_MAX_STACK_TILE_SIZE}
$<$<BOOL:${OSPRAY_DEBUG_NO_TASKING}>:OSPRAY_SERIAL_RENDERING>
)
target_include_directories(ospray_module_cpu
PUBLIC
$<BUILD_INTERFACE:${CMAKE_CURRENT_SOURCE_DIR}>
$<BUILD_INTERFACE:${CMAKE_SOURCE_DIR}/ospray/api>
$<BUILD_INTERFACE:${CMAKE_SOURCE_DIR}/ospray/common>
$<BUILD_INTERFACE:${CMAKE_SOURCE_DIR}/ospray>
$<BUILD_INTERFACE:${CMAKE_CURRENT_BINARY_DIR}>
$<BUILD_INTERFACE:${PROJECT_BINARY_DIR}>
# NOTE(jda) - the following includes are missing despite PUBLIC linking
$<BUILD_INTERFACE:${ISPCRT_INCLUDE_DIR}>
$<BUILD_INTERFACE:${OPENVKL_INCLUDE_DIRS}>
$<BUILD_INTERFACE:${EMBREE_INCLUDE_DIRS}>
#######################################################################
$<INSTALL_INTERFACE:${CMAKE_INSTALL_INCLUDEDIR}/ospray/SDK>
$<INSTALL_INTERFACE:${CMAKE_INSTALL_INCLUDEDIR}>
)
target_compile_definitions(ospray_module_cpu
PUBLIC
TILE_SIZE=${OSPRAY_TILE_SIZE}
MAX_TILE_SIZE=${OSPRAY_MAX_STACK_TILE_SIZE}
$<$<BOOL:${OSPRAY_DEBUG_NO_TASKING}>:OSPRAY_SERIAL_RENDERING>
SYCL_EXTERNAL=
__noinline=
)
if (OSPRAY_ENABLE_VOLUMES)
target_link_libraries(ospray_module_cpu
PUBLIC
$<BUILD_INTERFACE:openvkl::openvkl>
)
target_include_directories(ospray_module_cpu
PUBLIC
$<BUILD_INTERFACE:${OPENVKL_INCLUDE_DIRS}>
)
target_compile_definitions(ospray_module_cpu
PUBLIC
-DOSPRAY_ENABLE_VOLUMES
)
endif()
endif()
if (OSPRAY_MODULE_GPU)
# Have CMake's build infrastructure treat the ISPC files as C++ for DPC++
set_source_files_properties(${OSPRAY_ISPC_SOURCES} PROPERTIES
LANGUAGE CXX)
add_library(ospray_module_gpu SHARED
${OSPRAY_CPP_SOURCES} ${OSPRAY_ISPC_SOURCES})
target_link_libraries(ospray_module_gpu
PUBLIC
ospray
rkcommon::rkcommon_sycl
$<BUILD_INTERFACE:ispcrt::ispcrt>
$<BUILD_INTERFACE:embree>
# TODO: needs OpenVKL GPU port
#$<BUILD_INTERFACE:openvkl::openvkl>
)
target_include_directories(ospray_module_gpu
PUBLIC
$<BUILD_INTERFACE:${CMAKE_CURRENT_SOURCE_DIR}>
$<BUILD_INTERFACE:${CMAKE_SOURCE_DIR}/ospray/api>
$<BUILD_INTERFACE:${CMAKE_SOURCE_DIR}/ospray/common>
$<BUILD_INTERFACE:${CMAKE_SOURCE_DIR}/ospray>
$<BUILD_INTERFACE:${CMAKE_CURRENT_BINARY_DIR}>
$<BUILD_INTERFACE:${PROJECT_BINARY_DIR}>
# NOTE(jda) - the following includes are missing despite PUBLIC linking
$<BUILD_INTERFACE:${ISPCRT_INCLUDE_DIR}>
# TODO: needs OpenVKL GPU port
#$<BUILD_INTERFACE:${OPENVKL_INCLUDE_DIRS}>
$<BUILD_INTERFACE:${EMBREE_INCLUDE_DIRS}>
#######################################################################
$<INSTALL_INTERFACE:${CMAKE_INSTALL_INCLUDEDIR}/ospray/SDK>
$<INSTALL_INTERFACE:${CMAKE_INSTALL_INCLUDEDIR}>
)
target_compile_definitions(ospray_module_gpu
PUBLIC
TILE_SIZE=${OSPRAY_TILE_SIZE}
MAX_TILE_SIZE=${OSPRAY_MAX_STACK_TILE_SIZE}
$<$<BOOL:${OSPRAY_DEBUG_NO_TASKING}>:OSPRAY_SERIAL_RENDERING>
)
target_compile_options(ospray_module_gpu
PRIVATE
-x c++
)
target_compile_options(ospray_module_gpu PRIVATE
-fsycl
${OSPRAY_CXX_FLAGS_SYCL}
-fsycl-targets=${OSPRAY_SYCL_TARGET}
)
target_link_options(ospray_module_gpu PRIVATE
-fsycl
-fsycl-targets=${OSPRAY_SYCL_TARGET}
-Xsycl-target-backend=${OSPRAY_SYCL_TARGET} "-device ${OSPRAY_DPCPP_AOT_DEVICES} -revision_id ${OSPRAY_DPCPP_AOT_DEVICE_REVISION} ${OSPRAY_OCL_OPTIONS_STR} -options \"${OSPRAY_OCL_OTHER_OPTIONS_STR} -igc_opts='${OSPRAY_IGC_OPTIONS_STR}'\""
)
target_compile_definitions(ospray_module_gpu
PUBLIC
OSPRAY_TARGET_DPCPP
__SYCL_USE_NON_VARIADIC_SPIRV_OCL_PRINTF__
OSPRAY_BEGIN_ISPC_NAMESPACE=namespace\ ispc{
OSPRAY_END_ISPC_NAMESPACE=}
uniform=
varying=
unmasked=
export=
make_vec2i=vec2i
make_vec3i=vec3i
make_vec4i=vec4i
make_vec2ui=vec2ui
make_vec3ui=vec3ui
make_vec4ui=vec4ui
make_vec2f=vec2f
make_vec3f=vec3f
make_vec4f=vec4f
make_box1f=box1f
make_box2f=box2f
make_box3f=box3f
make_box3fa=box3fa
make_box4f=box4f
make_LinearSpace3f=LinearSpace3f
make_LinearSpace3f_varying_identity=LinearSpace3f
rtcIntersectV=rtcIntersect1
rtcIntersectExV=rtcIntersectEx1
rtcOccludedV=rtcOccluded1
rtcOccludedExV=rtcOccludedEx1
rtcInterpolateV1=rtcInterpolate1
programCount=1
programIndex=0
signbits=sycl::signbit
)
ospray_install_library(ospray_module_gpu lib)
endif()
# -------------------------------------------------------
# Install SDK headers
@@ -349,6 +543,7 @@ endforeach()
add_library(ospray_sdk INTERFACE)
target_link_libraries(ospray_sdk
INTERFACE
ospray_module_cpu rkcommon::rkcommon
ospray
rkcommon::rkcommon
)
ospray_install_target(ospray_sdk devel)
modules/cpu/ISPCDevice.cpp
+69 -3
View File
@@ -1,6 +1,12 @@
// Copyright 2009 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
#ifdef OSPRAY_TARGET_DPCPP
#include <level_zero/ze_api.h>
// ze_api and sycl level zero backend must be in this order
#include <sycl/ext/oneapi/backend/level_zero.hpp>
#endif
// ospray
#include "ISPCDevice.h"
#include "camera/Camera.h"
@@ -19,8 +25,10 @@
#include "render/Renderer.h"
#include "texture/Texture.h"
#include "texture/Texture2D.h"
#ifdef OSPRAY_ENABLE_VOLUMES
#include "volume/VolumetricModel.h"
#include "volume/transferFunction/TransferFunction.h"
#endif
// stl
#include <algorithm>
@@ -29,7 +37,9 @@
#include "rkcommon/tasking/tasking_system_init.h"
#include "rkcommon/utility/CodeTimer.h"
#ifndef OSPRAY_TARGET_DPCPP
#include "ISPCDevice_ispc.h"
#endif
namespace ospray {
namespace api {
@@ -84,9 +94,11 @@ static std::map<OSPDataType, std::function<SetParamFcn>> setParamFcns = {
declare_param_setter_object(Material *),
declare_param_setter_object(Renderer *),
declare_param_setter_object(Texture *),
#ifdef OSPRAY_ENABLE_VOLUMES
declare_param_setter_object(TransferFunction *),
declare_param_setter_object(Volume *),
declare_param_setter_object(VolumetricModel *),
#endif
declare_param_setter_object(World *),
declare_param_setter_string(const char *),
declare_param_setter(char *),
@@ -147,8 +159,36 @@ static std::map<OSPDataType, std::function<SetParamFcn>> setParamFcns = {
ISPCDevice::ISPCDevice()
: loadBalancer(std::make_shared<LocalTiledLoadBalancer>()),
ispcrtDevice(ISPCRT_DEVICE_TYPE_CPU)
{}
#ifdef OSPRAY_TARGET_DPCPP
ispcrtDevice(ISPCRT_DEVICE_TYPE_GPU),
#else
ispcrtDevice(ISPCRT_DEVICE_TYPE_CPU),
#endif
ispcrtQueue(ispcrtDevice)
{
#ifdef OSPRAY_TARGET_DPCPP
syclPlatform = sycl::ext::oneapi::level_zero::make_platform(
reinterpret_cast<pi_native_handle>(ispcrtDevice.nativePlatformHandle()));
syclDevice = sycl::ext::oneapi::level_zero::make_device(syclPlatform,
reinterpret_cast<pi_native_handle>(ispcrtDevice.nativeDeviceHandle()));
std::cout << "Using GPU device from ISPCRT: "
<< syclDevice.get_info<sycl::info::device::name>() << std::endl;
syclContext = sycl::ext::oneapi::level_zero::make_context(
std::vector<sycl::device>{syclDevice},
reinterpret_cast<pi_native_handle>(ispcrtDevice.nativeContextHandle()),
true);
syclQueue = sycl::ext::oneapi::level_zero::make_queue(syclContext,
syclDevice,
reinterpret_cast<pi_native_handle>(ispcrtQueue.nativeTaskQueueHandle()),
true);
// TODO needed?
setenv("SYCL_PI_LEVEL_ZERO_USM_ALLOCATOR", "1;0;shared:64K,0,2M", 1);
loadBalancer->setQueue(&syclQueue);
#endif
} // namespace api
ISPCDevice::~ISPCDevice()
{
@@ -160,9 +200,11 @@ ISPCDevice::~ISPCDevice()
// silently move on, sometimes a pthread mutex lock fails in Embree
}
#ifdef OSPRAY_ENABLE_VOLUMES
if (vklDevice) {
vklReleaseDevice(vklDevice);
}
#endif
}
static void embreeErrorFunc(void *, const RTCError code, const char *str)
@@ -173,6 +215,7 @@ static void embreeErrorFunc(void *, const RTCError code, const char *str)
handleError(e, "Embree internal error '" + std::string(str) + "'");
}
#ifdef OSPRAY_ENABLE_VOLUMES
static void vklErrorFunc(void *, const VKLError code, const char *str)
{
postStatusMsg() << "#osp: Open VKL internal error " << code << " : " << str;
@@ -180,6 +223,7 @@ static void vklErrorFunc(void *, const VKLError code, const char *str)
(code > VKL_UNSUPPORTED_CPU) ? OSP_UNKNOWN_ERROR : (OSPError)code;
handleError(e, "Open VKL internal error '" + std::string(str) + "'");
}
#endif
void ISPCDevice::commit()
{
@@ -188,7 +232,12 @@ void ISPCDevice::commit()
tasking::initTaskingSystem(numThreads, true);
if (!embreeDevice) {
#ifdef OSPRAY_TARGET_DPCPP
embreeDevice = rtcNewSYCLDevice(
&syclContext, &syclQueue, generateEmbreeDeviceCfg(*this).c_str());
#else
embreeDevice = rtcNewDevice(generateEmbreeDeviceCfg(*this).c_str());
#endif
rtcSetDeviceErrorFunction(embreeDevice, embreeErrorFunc, nullptr);
RTCError erc = rtcGetDeviceError(embreeDevice);
if (erc != RTC_ERROR_NONE) {
@@ -198,6 +247,7 @@ void ISPCDevice::commit()
}
}
#ifdef OSPRAY_ENABLE_VOLUMES
if (!vklDevice) {
vklLoadModule("cpu_device");
@@ -230,7 +280,9 @@ void ISPCDevice::commit()
vklCommitDevice(vklDevice);
}
#endif
#ifndef OSPRAY_TARGET_DPCPP
// Output device info string
const char *isaNames[] = {"unknown",
"SSE2",
@@ -243,6 +295,7 @@ void ISPCDevice::commit()
postStatusMsg(OSP_LOG_INFO)
<< "Using ISPC device with " << isaNames[ispc::ISPCDevice_isa()]
<< " instruction set...";
#endif
}
///////////////////////////////////////////////////////////////////////////
@@ -299,7 +352,11 @@ OSPGeometry ISPCDevice::newGeometry(const char *type)
OSPVolume ISPCDevice::newVolume(const char *type)
{
#ifdef OSPRAY_ENABLE_VOLUMES
return (OSPVolume) new Volume(*this, type);
#else
return (OSPVolume) nullptr;
#endif
}
OSPGeometricModel ISPCDevice::newGeometricModel(OSPGeometry _geom)
@@ -311,9 +368,13 @@ OSPGeometricModel ISPCDevice::newGeometricModel(OSPGeometry _geom)
OSPVolumetricModel ISPCDevice::newVolumetricModel(OSPVolume _volume)
{
#ifdef OSPRAY_ENABLE_VOLUMES
auto *volume = (Volume *)_volume;
auto *model = new VolumetricModel(*this, volume);
return (OSPVolumetricModel)model;
#else
return (OSPVolumetricModel) nullptr;
#endif
}
///////////////////////////////////////////////////////////////////////////
@@ -328,7 +389,11 @@ OSPMaterial ISPCDevice::newMaterial(
OSPTransferFunction ISPCDevice::newTransferFunction(const char *type)
{
#ifdef OSPRAY_ENABLE_VOLUMES
return (OSPTransferFunction)TransferFunction::createInstance(type, *this);
#else
return (OSPTransferFunction) nullptr;
#endif
}
OSPTexture ISPCDevice::newTexture(const char *type)
@@ -430,7 +495,8 @@ OSPFrameBuffer ISPCDevice::frameBufferCreate(
OSPImageOperation ISPCDevice::newImageOp(const char *type)
{
return (OSPImageOperation)ImageOp::createInstance(type);
ospray::ImageOp *ret = ImageOp::createInstance(type);
return (OSPImageOperation)ret;
}
const void *ISPCDevice::frameBufferMap(
modules/cpu/ISPCDevice.h
+34 -1
View File
@@ -3,17 +3,30 @@
#pragma once
#ifdef OSPRAY_TARGET_DPCPP
#include <CL/sycl.hpp>
#endif
// ospray
#include "Device.h"
// ispcrt
#include "ispcrt.hpp"
// embree
#include "embree3/rtcore.h"
#include "embree4/rtcore.h"
#ifdef OSPRAY_ENABLE_VOLUMES
// openvkl
#include "openvkl/openvkl.h"
#endif
/*! \file ISPCDevice.h Implements the "local" device for local rendering */
#ifdef OSPRAY_TARGET_DPCPP
namespace ispc {
int ISPCDevice_programCount();
int ISPCDevice_isa();
} // namespace ispc
#endif
namespace ospray {
struct LocalTiledLoadBalancer;
@@ -134,20 +147,40 @@ struct OSPRAY_SDK_INTERFACE ISPCDevice : public Device
return embreeDevice;
}
#ifdef OSPRAY_ENABLE_VOLUMES
VKLDevice getVklDevice()
{
return vklDevice;
}
#endif
ispcrt::Device &getIspcrtDevice()
{
return ispcrtDevice;
}
#ifdef OSPRAY_TARGET_DPCPP
sycl::queue &getSyclQueue()
{
return syclQueue;
}
#endif
private:
ispcrt::Device ispcrtDevice;
ispcrt::TaskQueue ispcrtQueue;
RTCDevice embreeDevice = nullptr;
#ifdef OSPRAY_ENABLE_VOLUMES
VKLDevice vklDevice = nullptr;
#endif
#ifdef OSPRAY_TARGET_DPCPP
sycl::platform syclPlatform;
sycl::device syclDevice;
sycl::context syclContext;
sycl::queue syclQueue;
#endif
};
} // namespace api
modules/cpu/ISPCDevice.ispc
+19
View File
@@ -1,6 +1,23 @@
// Copyright 2009 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
OSPRAY_BEGIN_ISPC_NAMESPACE
#ifdef OSPRAY_TARGET_DPCPP
int ISPCDevice_programCount()
{
// TODO: This should query the GPU width for ISPC,
// But for DPC++ we're in a scalar context
return 1;
}
int ISPCDevice_isa()
{
// TODO: GPU ISA/width detection
return 0;
}
#else
export uniform int ISPCDevice_programCount()
{
return programCount;
@@ -26,3 +43,5 @@ export uniform int ISPCDevice_isa()
return 0;
#endif
}
#endif
OSPRAY_END_ISPC_NAMESPACE
modules/cpu/ModuleInit.cpp
+16 -3
View File
@@ -15,17 +15,28 @@
#include "lights/registration.h"
#include "render/registration.h"
#include "texture/registration.h"
#ifdef OSPRAY_ENABLE_VOLUMES
#include "volume/transferFunction/registration.h"
#endif
using namespace ospray;
extern "C" OSPError OSPRAY_DLLEXPORT ospray_module_init_cpu(
int16_t versionMajor, int16_t versionMinor, int16_t /*versionPatch*/)
extern "C" OSPError OSPRAY_DLLEXPORT
#ifdef OSPRAY_TARGET_DPCPP
ospray_module_init_gpu
#else
ospray_module_init_cpu
#endif
(int16_t versionMajor, int16_t versionMinor, int16_t /*versionPatch*/)
{
auto status = moduleVersionCheck(versionMajor, versionMinor);
if (status == OSP_NO_ERROR) {
api::Device::registerType<api::ISPCDevice>("cpu");
#ifdef OSPRAY_TARGET_DPCPP
api::Device::registerType<ospray::api::ISPCDevice>("gpu");
#else
api::Device::registerType<ospray::api::ISPCDevice>("cpu");
#endif
registerAllCameras();
registerAllImageOps();
@@ -34,7 +45,9 @@ extern "C" OSPError OSPRAY_DLLEXPORT ospray_module_init_cpu(
registerAllMaterials();
registerAllRenderers();
registerAllTextures();
#ifdef OSPRAY_ENABLE_VOLUMES
registerAllTransferFunctions();
#endif
}
return status;
modules/cpu/camera/Camera.ih
+14 -1
View File
@@ -5,11 +5,22 @@
#pragma once
#include "common/Ray.ih"
#ifdef OSPRAY_TARGET_DPCPP
#include "rkcommon/math/box.h"
using namespace rkcommon;
using namespace rkcommon::math;
#include "common/rkcommonDPCPPWrappers.h"
#else
#include "rkcommon/math/box.ih"
#endif
#include "common/Ray.ih"
// c++ shared
#include "CameraShared.h"
OSPRAY_BEGIN_ISPC_NAMESPACE
// Specifies the input parameters (time, screen, and lens
// samples) required for 'initRay' to generate a primary ray
struct CameraSample
@@ -34,3 +45,5 @@ inline float Camera_shutterTime(
const float pos = self->rollingShutterHorizontal ? screen.x : screen.y;
return clamp(time + lerp(pos, self->shutter.lower, self->shutter.upper));
}
OSPRAY_END_ISPC_NAMESPACE
modules/cpu/camera/CameraShared.h
+10 -2
View File
@@ -3,23 +3,30 @@
#pragma once
#include "CameraType.ih"
#ifdef __cplusplus
#include "common/StructShared.h"
namespace ispc {
#endif // __cplusplus
#if defined(__cplusplus) && !defined(OSPRAY_TARGET_DPCPP)
typedef void *Camera_initRay;
#else
struct Camera;
struct CameraSample;
struct Ray;
// Fct pointer type for 'virtual' method that sets a pixel,
// generated ray.dir must be normalized to ensure ray.t is world-space distance
typedef void (*Camera_initRay)(const Camera *uniform,
varying Ray &ray,
const varying CameraSample &sample);
#endif // __cplusplus
#endif
struct Camera
{
CameraType type;
Camera_initRay initRay; // the 'virtual' initRay() method
bool motionBlur; // for the camera itself only, not in general
@@ -34,7 +41,8 @@ struct Camera
#ifdef __cplusplus
Camera()
: initRay(nullptr),
: type(CAMERA_TYPE_UNKNOWN),
initRay(nullptr),
motionBlur(false),
nearClip(1e-6f),
subImage(0.f),
modules/cpu/camera/CameraType.ih
+36
View File
@@ -0,0 +1,36 @@
// Copyright 2022 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
#pragma once
#ifdef OSPRAY_TARGET_DPCPP
#include <CL/sycl.hpp>
#include "rkcommon/platform.h"
#endif
#ifdef __cplusplus
namespace ispc {
#endif // __cplusplus
struct Camera;
struct CameraSample;
struct Ray;
enum CameraType
{
CAMERA_TYPE_PERSPECTIVE,
CAMERA_TYPE_ORTHOGRAPHIC,
CAMERA_TYPE_PANORAMIC,
CAMERA_TYPE_UNKNOWN,
};
#if defined(ISPC) || defined(OSPRAY_TARGET_DPCPP)
SYCL_EXTERNAL __noinline void Camera_dispatch_initRay(const Camera *uniform self,
varying Ray &ray,
const varying CameraSample &sample);
#endif
#ifdef __cplusplus
}
#endif
modules/cpu/camera/CameraType.ispc
+39
View File
@@ -0,0 +1,39 @@
// Copyright 2022 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
#include "Camera.ih"
#include "CameraType.ih"
#include "OrthographicCamera.ih"
#include "PanoramicCamera.ih"
#include "PerspectiveCamera.ih"
OSPRAY_BEGIN_ISPC_NAMESPACE
SYCL_EXTERNAL __noinline void Camera_dispatch_initRay(const Camera *uniform self,
varying Ray &ray,
const varying CameraSample &sample)
{
switch (self->type) {
case CAMERA_TYPE_PERSPECTIVE:
PerspectiveCamera_initRay(self, ray, sample);
break;
case CAMERA_TYPE_ORTHOGRAPHIC:
OrthographicCamera_initRay(self, ray, sample);
break;
case CAMERA_TYPE_PANORAMIC:
PanoramicCamera_initRay(self, ray, sample);
break;
default:
// TODO: No function pointer for DPCPP version because it'll hang.
// Should we have an error here or still allow it to be called?
#ifndef OSPRAY_TARGET_DPCPP
self->initRay(self, ray, sample);
break;
#else
break;
#endif
}
}
OSPRAY_END_ISPC_NAMESPACE
modules/cpu/camera/OrthographicCamera.cpp
+6 -1
View File
@@ -2,15 +2,20 @@
// SPDX-License-Identifier: Apache-2.0
#include "OrthographicCamera.h"
#ifndef OSPRAY_TARGET_DPCPP
// ispc exports
#include "camera/OrthographicCamera_ispc.h"
#endif
namespace ospray {
OrthographicCamera::OrthographicCamera(api::ISPCDevice &device)
: AddStructShared(device.getIspcrtDevice(), device)
{
getSh()->super.initRay = ispc::OrthographicCamera_initRay_addr();
#ifndef OSPRAY_TARGET_DPCPP
getSh()->super.initRay = reinterpret_cast<ispc::Camera_initRay>(
ispc::OrthographicCamera_initRay_addr());
#endif
}
std::string OrthographicCamera::toString() const
modules/cpu/camera/OrthographicCamera.ih
+14
View File
@@ -0,0 +1,14 @@
// Copyright 2022 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
#include "camera/Camera.ih"
OSPRAY_BEGIN_ISPC_NAMESPACE
SYCL_EXTERNAL void OrthographicCamera_initRay(const Camera *uniform _self,
varying Ray &ray,
const varying CameraSample &sample);
OSPRAY_END_ISPC_NAMESPACE
modules/cpu/camera/OrthographicCamera.ispc
+10 -2
View File
@@ -6,7 +6,9 @@
// c++ shared
#include "OrthographicCameraShared.h"
void OrthographicCamera_initRay(const Camera *uniform _self,
OSPRAY_BEGIN_ISPC_NAMESPACE
SYCL_EXTERNAL void OrthographicCamera_initRay(const Camera *uniform _self,
varying Ray &ray,
const varying CameraSample &sample)
{
@@ -43,7 +45,7 @@ void OrthographicCamera_initRay(const Camera *uniform _self,
export void *uniform OrthographicCamera_initRay_addr()
{
return OrthographicCamera_initRay;
return (void *uniform)OrthographicCamera_initRay;
}
export void OrthographicCamera_projectBox(
@@ -55,7 +57,11 @@ export void OrthographicCamera_projectBox(
const uniform vec3f dvn = self->dv_up / dot(self->dv_up, self->dv_up);
vec3f projectedPt = make_vec3f(-1.f, -1.f, 1e20f);
#ifdef OSPRAY_TARGET_DPCPP
for (uint32 i = 0; i < 8; ++i) {
#else
foreach (i = 0 ... 8) {
#endif
// Get the point we should be projecting
vec3f p;
switch (i) {
@@ -117,3 +123,5 @@ export void OrthographicCamera_projectBox(
projection.upper.z = reduce_max(projectedPt.z);
}
}
OSPRAY_END_ISPC_NAMESPACE
modules/cpu/camera/OrthographicCameraShared.h
+11 -1
View File
@@ -23,8 +23,18 @@ struct OrthographicCamera
// up direction of camera
#ifdef __cplusplus
OrthographicCamera() : dir(0.f), org(0.f), du_size(0.f), dv_up(0.f) {}
OrthographicCamera() : dir(0.f), org(0.f), du_size(0.f), dv_up(0.f)
{
super.type = CAMERA_TYPE_ORTHOGRAPHIC;
}
};
#ifdef OSPRAY_TARGET_DPCPP
void *OrthographicCamera_initRay_addr();
void OrthographicCamera_projectBox(
void *_self, const box3f &box, box3f &projection);
#endif
} // namespace ispc
#else
};
modules/cpu/camera/PanoramicCamera.cpp
+6 -1
View File
@@ -3,14 +3,19 @@
#include "PanoramicCamera.h"
// ispc exports
#ifndef OSPRAY_TARGET_DPCPP
#include "camera/PanoramicCamera_ispc.h"
#endif
namespace ospray {
PanoramicCamera::PanoramicCamera(api::ISPCDevice &device)
: AddStructShared(device.getIspcrtDevice(), device)
{
getSh()->super.initRay = ispc::PanoramicCamera_initRay_addr();
#ifndef OSPRAY_TARGET_DPCPP
getSh()->super.initRay = reinterpret_cast<ispc::Camera_initRay>(
ispc::PanoramicCamera_initRay_addr());
#endif
}
std::string PanoramicCamera::toString() const
modules/cpu/camera/PanoramicCamera.ih
+13
View File
@@ -0,0 +1,13 @@
// Copyright 2022 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
#include "camera/Camera.ih"
OSPRAY_BEGIN_ISPC_NAMESPACE
SYCL_EXTERNAL void PanoramicCamera_initRay(const Camera *uniform _self,
varying Ray &ray,
const varying CameraSample &sample);
OSPRAY_END_ISPC_NAMESPACE
modules/cpu/camera/PanoramicCamera.ispc
+11 -4
View File
@@ -5,10 +5,15 @@
#include "common/MotionTransform.ih"
#include "math/sampling.ih"
#include "ospray/OSPEnums.h"
#ifdef OSPRAY_TARGET_DPCPP
#include "common/rkcommonDPCPPWrappers.h"
#endif
// c++ shared
#include "PanoramicCameraShared.h"
void PanoramicCamera_initRay(const Camera *uniform _self,
OSPRAY_BEGIN_ISPC_NAMESPACE
SYCL_EXTERNAL void PanoramicCamera_initRay(const Camera *uniform _self,
varying Ray &ray,
const varying CameraSample &sample)
{
@@ -40,8 +45,8 @@ void PanoramicCamera_initRay(const Camera *uniform _self,
screen = Camera_subRegion(_self, screen);
const float phi = two_pi * screen.x;
const float theta = M_PI * screen.y;
const float phi = (float)two_pi * screen.x;
const float theta = (float)pi * screen.y;
float sinTheta, cosTheta;
sincos(theta, &sinTheta, &cosTheta);
@@ -76,5 +81,7 @@ void PanoramicCamera_initRay(const Camera *uniform _self,
export void *uniform PanoramicCamera_initRay_addr()
{
return PanoramicCamera_initRay;
return (void *uniform)PanoramicCamera_initRay;
}
OSPRAY_END_ISPC_NAMESPACE
modules/cpu/camera/PanoramicCameraShared.h
+7 -1
View File
@@ -21,8 +21,14 @@ struct PanoramicCamera
#ifdef __cplusplus
PanoramicCamera()
: org(0.f), frame(one), stereoMode(OSP_STEREO_NONE), ipd_offset(0.f)
{}
{
super.type = CAMERA_TYPE_PANORAMIC;
}
};
#ifdef OSPRAY_TARGET_DPCPP
void *PanoramicCamera_initRay_addr();
#endif
} // namespace ispc
#else
};
modules/cpu/camera/PerspectiveCamera.cpp
+38 -3
View File
@@ -2,15 +2,20 @@
// SPDX-License-Identifier: Apache-2.0
#include "PerspectiveCamera.h"
#ifndef OSPRAY_TARGET_DPCPP
// ispc exports
#include "camera/PerspectiveCamera_ispc.h"
#endif
namespace ospray {
PerspectiveCamera::PerspectiveCamera(api::ISPCDevice &device)
: AddStructShared(device.getIspcrtDevice(), device)
{
getSh()->super.initRay = ispc::PerspectiveCamera_initRay_addr();
#ifndef OSPRAY_TARGET_DPCPP
getSh()->super.initRay = reinterpret_cast<ispc::Camera_initRay>(
ispc::PerspectiveCamera_initRay_addr());
#endif
}
std::string PerspectiveCamera::toString() const
@@ -57,13 +62,19 @@ void PerspectiveCamera::commit()
getSh()->imgPlaneSize = imgPlaneSize;
if (getSh()->super.motionBlur) {
getSh()->super.initRay = ispc::PerspectiveCamera_initRayMB_addr();
#ifndef OSPRAY_TARGET_DPCPP
getSh()->super.initRay = reinterpret_cast<ispc::Camera_initRay>(
ispc::PerspectiveCamera_initRayMB_addr());
#endif
getSh()->du_size = vec3f(imgPlaneSize.x, imgPlaneSize.y, architectural);
getSh()->dv_up = up;
getSh()->ipd_offset =
vec3f(0.5f * interpupillaryDistance, focusDistance, 0.0f);
} else {
getSh()->super.initRay = ispc::PerspectiveCamera_initRay_addr();
#ifndef OSPRAY_TARGET_DPCPP
getSh()->super.initRay = reinterpret_cast<ispc::Camera_initRay>(
ispc::PerspectiveCamera_initRay_addr());
#endif
getSh()->du_size = normalize(cross(getSh()->dir_00, up));
if (architectural) // orient film to be parallel to 'up'
getSh()->dv_up = normalize(up);
@@ -103,6 +114,30 @@ void PerspectiveCamera::commit()
}
}
#ifdef OSPRAY_TARGET_DPCPP
/*
void PerspectiveCamera::setGPUFunctionPtrs(sycl::queue &syclQueue)
{
auto *sSh = getSh();
if (sSh->super.motionBlur) {
throw std::runtime_error(
"TODO Motion blur not supported in SYCL mode see
rtcGetGeometryTransform");
}
auto event = syclQueue.submit([&](sycl::handler &cgh) {
cgh.parallel_for(1, [=](cl::sycl::id<1>) RTC_SYCL_KERNEL {
// if (sSh->super.motionBlur) {
// sSh->super.initRay = ispc::PerspectiveCamera_initRayMB;
//} else {
sSh->super.initRay = ispc::PerspectiveCamera_initRay;
//}
});
});
event.wait();
}
*/
#endif
box3f PerspectiveCamera::projectBox(const box3f &b) const
{
if (stereoMode != OSP_STEREO_NONE) {
modules/cpu/camera/PerspectiveCamera.h
+4
View File
@@ -19,6 +19,10 @@ struct OSPRAY_SDK_INTERFACE PerspectiveCamera
virtual std::string toString() const override;
virtual void commit() override;
#ifdef OSPRAY_TARGET_DPCPP
// virtual void setGPUFunctionPtrs(sycl::queue &syclQueue) override;
#endif
box3f projectBox(const box3f &b) const override;
// Data members //
modules/cpu/camera/PerspectiveCamera.ih
+13
View File
@@ -0,0 +1,13 @@
// Copyright 2022 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
#include "camera/Camera.ih"
OSPRAY_BEGIN_ISPC_NAMESPACE
SYCL_EXTERNAL void PerspectiveCamera_initRay(const Camera *uniform _self,
varying Ray &ray,
const varying CameraSample &sample);
OSPRAY_END_ISPC_NAMESPACE
modules/cpu/camera/PerspectiveCamera.ispc
+22 -4
View File
@@ -8,10 +8,13 @@
// c++ shared
#include "PerspectiveCameraShared.h"
void PerspectiveCamera_initRay(const Camera *uniform _self,
OSPRAY_BEGIN_ISPC_NAMESPACE
SYCL_EXTERNAL void PerspectiveCamera_initRay(const Camera *uniform _self,
varying Ray &ray,
const varying CameraSample &sample)
{
// TODO: We need ot call initRayMB if we have motionBlur
const PerspectiveCamera *uniform self =
(const PerspectiveCamera *uniform)_self;
@@ -20,7 +23,11 @@ void PerspectiveCamera_initRay(const Camera *uniform _self,
const uniform bool sbs = self->stereoMode == OSP_STEREO_SIDE_BY_SIDE;
varying float *uniform split = sbs ? &screen.x : &screen.y;
#ifdef OSPRAY_TARGET_DPCPP
if (sbs || self->stereoMode == OSP_STEREO_TOP_BOTTOM) {
#else
if (or (sbs, self->stereoMode == OSP_STEREO_TOP_BOTTOM)) {
#endif
*split *= 2.f;
if (*split < 1.f) {
org = org - self->ipd_offset;
@@ -47,7 +54,7 @@ void PerspectiveCamera_initRay(const Camera *uniform _self,
setRay(ray, org, normalize(dir), self->super.nearClip, inf, time);
}
void PerspectiveCamera_initRayMB(const Camera *uniform _self,
SYCL_EXTERNAL void PerspectiveCamera_initRayMB(const Camera *uniform _self,
varying Ray &ray,
const varying CameraSample &sample)
{
@@ -59,7 +66,11 @@ void PerspectiveCamera_initRayMB(const Camera *uniform _self,
const uniform bool sbs = self->stereoMode == OSP_STEREO_SIDE_BY_SIDE;
varying float *uniform split = sbs ? &screen.x : &screen.y;
#ifdef OSPRAY_TARGET_DPCPP
if (sbs || self->stereoMode == OSP_STEREO_TOP_BOTTOM) {
#else
if (or (sbs, self->stereoMode == OSP_STEREO_TOP_BOTTOM)) {
#endif
*split *= 2.f;
if (*split < 1.f)
ipd_offset1 = -ipd_offset1;
@@ -69,6 +80,7 @@ void PerspectiveCamera_initRayMB(const Camera *uniform _self,
screen = Camera_subRegion(_self, screen);
const float time = Camera_shutterTime(_self, screen, sample.time);
// TODO: rtcGetGeometryTransform not callable on device?
const affine3f xfm = getInterpolatedTransform(self->super.geom, time);
vec3f org = xfmPoint(xfm, self->org);
@@ -124,12 +136,12 @@ void PerspectiveCamera_initRayMB(const Camera *uniform _self,
export void *uniform PerspectiveCamera_initRay_addr()
{
return PerspectiveCamera_initRay;
return (void *uniform)PerspectiveCamera_initRay;
}
export void *uniform PerspectiveCamera_initRayMB_addr()
{
return PerspectiveCamera_initRayMB;
return (void *uniform)PerspectiveCamera_initRayMB;
}
export void PerspectiveCamera_projectBox(
@@ -143,7 +155,11 @@ export void PerspectiveCamera_projectBox(
const uniform vec3f dvn = normalize(self->dv_up) / self->imgPlaneSize.y;
vec3f projectedPt = make_vec3f(-1.f, -1.f, 1e20f);
#ifdef OSPRAY_TARGET_DPCPP
for (uint32 i = 0; i < 8; ++i) {
#else
foreach (i = 0 ... 8) {
#endif
// Get the point we should be projecting
vec3f p;
switch (i) {
@@ -221,3 +237,5 @@ export void PerspectiveCamera_projectBox(
projection.upper.y = 1.f;
}
}
OSPRAY_END_ISPC_NAMESPACE
modules/cpu/camera/PerspectiveCameraShared.h
+9 -1
View File
@@ -43,8 +43,16 @@ struct PerspectiveCamera
scaledAperture(0.f),
aspect(1.f),
stereoMode(OSP_STEREO_NONE)
{}
{
super.type = CAMERA_TYPE_PERSPECTIVE;
}
};
#ifdef OSPRAY_TARGET_DPCPP
void PerspectiveCamera_projectBox(
void *_self, const box3f &box, box3f &projection);
#endif
} // namespace ispc
#else
};
modules/cpu/common/BufferShared.h
+57 -1
View File
@@ -3,6 +3,8 @@
#pragma once
#include <cstring>
#include <memory>
#include "ispcrt.hpp"
namespace ospray {
@@ -31,6 +33,18 @@ struct BufferShared : public ispcrt::Array<T, ispcrt::AllocType::Shared>
BufferShared(ispcrt::Device &device, size_t size);
BufferShared(ispcrt::Device &device, const std::vector<T> &v);
BufferShared(ispcrt::Device &device, const T *data, size_t size);
// TODO WILL: Probably best to put these methods/operators in ISPCRT
T *data();
T *begin();
T *end();
const T *cbegin() const;
const T *cend() const;
T &operator[](const size_t i);
const T &operator[](const size_t i) const;
};
template <typename T>
@@ -58,9 +72,51 @@ BufferShared<T>::BufferShared(
std::memcpy(sharedPtr(), data, sizeof(T) * size);
}
template <typename T>
T *BufferShared<T>::data()
{
return begin();
}
template <typename T>
T *BufferShared<T>::begin()
{
return sharedPtr();
}
template <typename T>
T *BufferShared<T>::end()
{
return begin() + ispcrt::Array<T, ispcrt::AllocType::Shared>::size();
}
template <typename T>
const T *BufferShared<T>::cbegin() const
{
return sharedPtr();
}
template <typename T>
const T *BufferShared<T>::cend() const
{
return cbegin() + ispcrt::Array<T, ispcrt::AllocType::Shared>::size();
}
template <typename T>
T &BufferShared<T>::operator[](const size_t i)
{
return *(sharedPtr() + i);
}
template <typename T>
const T &BufferShared<T>::operator[](const size_t i) const
{
return *(sharedPtr() + i);
}
template <typename T, typename... Args>
inline std::unique_ptr<BufferShared<T>> make_buffer_shared_unique(
Args &&... args)
Args &&...args)
{
return std::unique_ptr<BufferShared<T>>(
new BufferShared<T>(std::forward<Args>(args)...));
modules/cpu/common/Clipping.ih
+14 -1
View File
@@ -4,8 +4,19 @@
#pragma once
// ospray
#include "../math/AffineSpace.ih"
#ifdef OSPRAY_TARGET_DPCPP
#include "rkcommon/math/AffineSpace.h"
#include "rkcommon/math/box.h"
using namespace rkcommon;
using namespace rkcommon::math;
#else
#include "rkcommon/math/AffineSpace.ih"
#include "rkcommon/math/box.ih"
#endif
#include "common/OSPCommon.ih"
OSPRAY_BEGIN_ISPC_NAMESPACE
// Array of all ray intervals across ray
#define CLIPPING_INTERVALS_MAX_COUNT 64
@@ -50,3 +61,5 @@ inline bool overlapped(const range1f &a, const range1f &b)
{
return (a.upper > b.lower) && (a.lower < b.upper);
}
OSPRAY_END_ISPC_NAMESPACE
modules/cpu/common/DGEnum.h
+1 -1
View File
@@ -22,4 +22,4 @@ typedef enum
#ifdef __cplusplus
} // namespace ospray
#endif // __cplusplus
#endif // __cplusplus
modules/cpu/common/Data.cpp
+11 -1
View File
@@ -21,7 +21,17 @@ Data::Data(api::ISPCDevice &device,
{
if (sharedData == nullptr)
throw std::runtime_error("OSPData: shared buffer is NULL");
addr = (char *)sharedData;
// addr = (char *)sharedData;
// TODO: Here we need to check if it's a GPU pointer or not and do the
// alloc/copy Shared data on GPU needs a way to share the USM context from
// OSPRay. What about the ANARI mapping stuff? This is a hack right now to
// copy all data into USM so that we can run the ospExamples
shared = false;
view = BufferSharedCreate(
device.getIspcrtDevice().handle(), size() * sizeOf(type) + 16);
addr = (char *)ispcrtSharedPtr(view);
std::memcpy(addr, sharedData, size() * sizeOf(type));
init();
if (isObjectType(type)) {
modules/cpu/common/Data.h
+31
View File
@@ -8,6 +8,9 @@
#include "StructShared.h"
#include "ispcrt.hpp"
#ifdef OSPRAY_TARGET_DPCPP
#include "Data.ih"
#else
// including "Data_ispc.h" breaks app code using SDK headers
#ifndef __ISPC_STRUCT_Data1D__
#define __ISPC_STRUCT_Data1D__
@@ -21,6 +24,7 @@ struct Data1D
};
} // namespace ispc
#endif
#endif
namespace ospray {
@@ -53,6 +57,9 @@ struct OSPRAY_SDK_INTERFACE Data : public ISPCDeviceObject
template <typename T, int DIM = 1>
const DataT<T, DIM> &as() const;
template <typename T, int DIM = 1>
DataT<T, DIM> &as();
template <typename T, int DIM>
typename std::enable_if<std::is_pointer<T>::value, bool>::type is() const;
@@ -60,6 +67,10 @@ struct OSPRAY_SDK_INTERFACE Data : public ISPCDeviceObject
typename std::enable_if<!std::is_pointer<T>::value, bool>::type is() const;
protected:
// The actual buffer storing the data
// TODO: What to do with OSPRay shared data that is not in USM?
// we need to detect & copy, we also need an API for apps to pass us USM
// and this USM has to be in the same L0 context...
ISPCRTMemoryView view{nullptr};
char *addr{nullptr};
bool shared;
@@ -245,6 +256,12 @@ const ispc::Data1D *ispc(Ref<const DataT<T, 1>> &dataRef)
return dataRef ? &dataRef->ispc : &Data::emptyData1D;
}
template <typename T>
ispc::Data1D *ispc(Ref<DataT<T, 1>> &dataRef)
{
return dataRef ? &dataRef->ispc : &Data::emptyData1D;
}
inline size_t Data::size() const
{
return numItems.x * numItems.y * numItems.z;
@@ -298,6 +315,20 @@ inline const DataT<T, DIM> &Data::as() const
}
}
template <typename T, int DIM>
inline DataT<T, DIM> &Data::as()
{
if (is<T, DIM>())
return (DataT<T, DIM> &)*this;
else {
std::stringstream ss;
ss << "Incompatible type or dimension for DataT; requested type[dim]: "
<< stringFor(OSPTypeFor<T>::value) << "[" << DIM
<< "], actual: " << stringFor(type) << "[" << dimensions << "].";
throw std::runtime_error(ss.str());
}
}
template <typename T, int DIM>
inline const Ref<const DataT<T, DIM>> ISPCDeviceObject::getParamDataT(
const char *name, bool required, bool promoteScalar)
modules/cpu/common/Data.ih
+22 -1
View File
@@ -3,8 +3,18 @@
#pragma once
#include "../math/LinearSpace.ih"
#ifdef OSPRAY_TARGET_DPCPP
#include "common/OSPCommon.ih"
#include "rkcommon/math/LinearSpace.h"
#include "rkcommon/math/box.h"
using namespace rkcommon;
using namespace rkcommon::math;
#else
#include "rkcommon/math/LinearSpace.ih"
#include "rkcommon/math/box.ih"
#endif
OSPRAY_BEGIN_ISPC_NAMESPACE
struct Data1D
{
@@ -32,6 +42,14 @@ inline uniform bool valid(const uniform Data1D data)
#define BITS 20
// ok to pass-by-value, will be inlined and optimized
#ifdef OSPRAY_TARGET_DPCPP
// DPC++ execution in a thread is scalar
#define __define_get(T) \
inline uniform T get_##T(const uniform Data1D data, const uniform int index) \
{ \
return *((const T *)(data.addr + data.byteStride * index)); \
}
#else
#define __define_get(T) \
inline T get_##T(const uniform Data1D data, const varying int index) \
{ \
@@ -55,6 +73,7 @@ inline uniform bool valid(const uniform Data1D data)
{ \
return *((const T *)(data.addr + data.byteStride * index)); \
}
#endif
__define_get(int32);
__define_get(vec2i);
@@ -81,3 +100,5 @@ __define_get(linear3f);
#undef __define_get
#undef BITS
OSPRAY_END_ISPC_NAMESPACE
modules/cpu/common/DifferentialGeometry.ih
+13 -2
View File
@@ -3,8 +3,18 @@
#pragma once
#include "DGEnum.h"
#ifdef OSPRAY_TARGET_DPCPP
#include "rkcommon/math/vec.h"
using namespace rkcommon;
using namespace rkcommon::math;
#else
#include "rkcommon/math/vec.ih"
#endif
#include "DGEnum.h"
#include "common/OSPCommon.ih"
OSPRAY_BEGIN_ISPC_NAMESPACE
struct Material;
struct Renderer;
@@ -39,7 +49,7 @@ struct DifferentialGeometry
};
// assumed precision of intersection routines
#define ulpEpsilon 0x1.fp-21
#define ulpEpsilon 0x1.0fp-21
inline float calcEpsilon(const vec3f &P, const vec3f &dir, float t)
{
@@ -53,3 +63,4 @@ inline float calcEpsilon(const vec3f &P, float dist)
const vec3f dummy_dir = make_vec3f(1.0f);
return calcEpsilon(P, dummy_dir, dist);
}
OSPRAY_END_ISPC_NAMESPACE
modules/cpu/common/FilterIntersect.ih
+27 -11
View File
@@ -6,16 +6,23 @@
#include "Intersect.ih"
#include "Ray.ih"
OSPRAY_BEGIN_ISPC_NAMESPACE
inline bool filterCall(const RTCIntersectFunctionNArguments *uniform args,
const RTCHit &rtcHitTemp)
{
// prepare execution mask
#ifdef OSPRAY_TARGET_DPCPP
// We can't #define __mask because it conflicts with some stuff in C++
int imask = -1;
#else
varying int imask;
varying bool mask = __mask;
unmasked
{
imask = mask ? -1 : 0;
}
#endif
// call filter intersection callback
uniform RTCFilterFunctionNArguments filterArgs;
@@ -28,7 +35,7 @@ inline bool filterCall(const RTCIntersectFunctionNArguments *uniform args,
rtcFilterIntersection(args, &filterArgs);
// return true if hit has been accepted
return (imask == -1);
return imask == -1;
}
inline bool filterIntersectionSingle(
@@ -39,7 +46,11 @@ inline bool filterIntersectionSingle(
{
// if there is a hit
varying Ray *uniform ray = (varying Ray * uniform) args->rayhit;
#ifdef OSPRAY_TARGET_DPCPP
if (hit.hit && hit.t > ray->t0 && hit.t <= ray->t) {
#else
if (and(hit.hit, and(hit.t > ray->t0, hit.t <= ray->t))) {
#endif
// set a new hit value
const float tOld = ray->t;
ray->t = hit.t;
@@ -60,19 +71,22 @@ inline bool filterIntersectionSingle(
if (alwaysReject) {
// the hit is always rejected
ray->t = tOld;
}
else
{
} else {
if (accept) {
// hit has been accepted
if (isOcclusionTest) {
ray->t = neg_inf;
}
else
{
varying RTCHit *uniform rtcHit =
(varying RTCHit * uniform) & ray->Ng.x;
*rtcHit = rtcHitTemp;
} else {
ray->Ng.x = hit.N.x;
ray->Ng.y = hit.N.y;
ray->Ng.z = hit.N.z;
ray->u = hit.u;
ray->v = 0.f;
ray->primID = args->primID;
ray->geomID = args->geomID;
ray->instID = args->context->instID[0];
}
return true;
} else {
@@ -93,6 +107,8 @@ inline bool filterIntersectionBoth(
const uniform bool isOcclusionTest)
{
if (!filterIntersectionSingle(args, isect.entry, isOcclusionTest, false)) {
filterIntersectionSingle(args, isect.exit, isOcclusionTest, false);
return filterIntersectionSingle(args, isect.exit, isOcclusionTest, false);
}
return true;
}
OSPRAY_END_ISPC_NAMESPACE
modules/cpu/common/Group.cpp
+31 -3
View File
@@ -6,7 +6,9 @@
#include "ISPCDevice.h"
#include "geometry/GeometricModel.h"
#include "lights/Light.h"
#ifdef OSPRAY_ENABLE_VOLUMES
#include "volume/VolumetricModel.h"
#endif
namespace ospray {
@@ -44,11 +46,15 @@ Group::Group(api::ISPCDevice &device)
Group::~Group()
{
freeAndNullifyEmbreeScene(sceneGeometries);
#ifdef OSPRAY_ENABLE_VOLUMES
freeAndNullifyEmbreeScene(sceneVolumes);
#endif
freeAndNullifyEmbreeScene(sceneClippers);
getSh()->geometricModels = nullptr;
#ifdef OSPRAY_ENABLE_VOLUMES
getSh()->volumetricModels = nullptr;
#endif
getSh()->clipModels = nullptr;
}
@@ -60,20 +66,26 @@ std::string Group::toString() const
void Group::commit()
{
geometricModels = getParamDataT<GeometricModel *>("geometry");
#ifdef OSPRAY_ENABLE_VOLUMES
volumetricModels = getParamDataT<VolumetricModel *>("volume");
#endif
clipModels = getParamDataT<GeometricModel *>("clippingGeometry");
lights = getParamDataT<Light *>("light");
size_t numGeometries = geometricModels ? geometricModels->size() : 0;
#ifdef OSPRAY_ENABLE_VOLUMES
size_t numVolumes = volumetricModels ? volumetricModels->size() : 0;
#endif
size_t numClippers = clipModels ? clipModels->size() : 0;
size_t numLights = lights ? lights->size() : 0;
postStatusMsg(OSP_LOG_DEBUG)
<< "=======================================================\n"
<< "Finalizing instance, which has " << numGeometries << " geometries, "
<< numVolumes << " volumes, " << numClippers
<< " clipping geometries and " << numLights << " lights";
#ifdef OSPRAY_ENABLE_VOLUMES
<< numVolumes << " volumes, "
#endif
<< numClippers << " clipping geometries and " << numLights << " lights";
int sceneFlags = RTC_SCENE_FLAG_NONE;
RTCBuildQuality buildQuality = RTC_BUILD_QUALITY_HIGH;
@@ -87,14 +99,21 @@ void Group::commit()
(getParam<bool>("robustMode", false) ? RTC_SCENE_FLAG_ROBUST : 0);
freeAndNullifyEmbreeScene(sceneGeometries);
#ifdef OSPRAY_ENABLE_VOLUMES
freeAndNullifyEmbreeScene(sceneVolumes);
#endif
freeAndNullifyEmbreeScene(sceneClippers);
geometricModelsArray = nullptr;
#ifdef OSPRAY_ENABLE_VOLUMES
volumetricModelsArray = nullptr;
#endif
clipModelsArray = nullptr;
getSh()->geometricModels = nullptr;
#ifdef OSPRAY_ENABLE_VOLUMES
getSh()->volumetricModels = nullptr;
#endif
getSh()->clipModels = nullptr;
RTCDevice embreeDevice = getISPCDevice().getEmbreeDevice();
@@ -115,6 +134,7 @@ void Group::commit()
rtcCommitScene(sceneGeometries);
}
#ifdef OSPRAY_ENABLE_VOLUMES
if (numVolumes > 0) {
sceneVolumes = rtcNewScene(embreeDevice);
createEmbreeScene(
@@ -127,6 +147,7 @@ void Group::commit()
rtcCommitScene(sceneVolumes);
}
#endif
if (numClippers > 0) {
sceneClippers = rtcNewScene(embreeDevice);
@@ -149,13 +170,18 @@ void Group::commit()
}
getSh()->numGeometricModels = numGeometries;
#ifdef OSPRAY_ENABLE_VOLUMES
getSh()->numVolumetricModels = numVolumes;
#endif
getSh()->numClipModels = numClippers;
// Create empty scene for lights-only group,
// it is needed to have rtcGeometry created in Instance object
// which in turn is needed for motion blur matrices interpolation
if ((numLights > 0) && (numGeometries == 0) && (numVolumes == 0)
if ((numLights > 0) && (numGeometries == 0)
#ifdef OSPRAY_ENABLE_VOLUMES
&& (numVolumes == 0)
#endif
&& (numClippers == 0)) {
sceneGeometries = rtcNewScene(embreeDevice);
rtcCommitScene(sceneGeometries);
@@ -172,10 +198,12 @@ box3f Group::getBounds() const
sceneBounds.extend(box3f(vec3f(&bounds.lower[0]), vec3f(&bounds.upper[0])));
}
#ifdef OSPRAY_ENABLE_VOLUMES
if (sceneVolumes) {
rtcGetSceneBounds(sceneVolumes, (RTCBounds *)&bounds);
sceneBounds.extend(box3f(vec3f(&bounds.lower[0]), vec3f(&bounds.upper[0])));
}
#endif
return sceneBounds;
}
modules/cpu/common/Group.h
+9 -1
View File
@@ -9,14 +9,16 @@
// stl
#include <vector>
// embree
#include "embree3/rtcore.h"
#include "embree4/rtcore.h"
// ispc shared
#include "GroupShared.h"
namespace ospray {
struct GeometricModel;
#ifdef OSPRAY_ENABLE_VOLUMES
struct VolumetricModel;
#endif
struct Light;
struct OSPRAY_SDK_INTERFACE Group
@@ -33,19 +35,25 @@ struct OSPRAY_SDK_INTERFACE Group
// Data members //
Ref<const DataT<GeometricModel *>> geometricModels;
#ifdef OSPRAY_ENABLE_VOLUMES
Ref<const DataT<VolumetricModel *>> volumetricModels;
#endif
Ref<const DataT<GeometricModel *>> clipModels;
int numInvertedClippers{0};
Ref<const DataT<Light *>> lights;
RTCScene sceneGeometries{nullptr};
#ifdef OSPRAY_ENABLE_VOLUMES
RTCScene sceneVolumes{nullptr};
#endif
RTCScene sceneClippers{nullptr};
private:
std::unique_ptr<BufferShared<ispc::GeometricModel *>> geometricModelsArray;
#ifdef OSPRAY_ENABLE_VOLUMES
std::unique_ptr<BufferShared<ispc::VolumetricModel *>> volumetricModelsArray;
#endif
std::unique_ptr<BufferShared<ispc::GeometricModel *>> clipModelsArray;
};
modules/cpu/common/GroupShared.h
+8
View File
@@ -9,25 +9,33 @@ namespace ispc {
#endif // __cplusplus
struct GeometricModel;
#ifdef OSPRAY_ENABLE_VOLUMES
struct VolumetricModel;
#endif
struct Group
{
GeometricModel **geometricModels;
int32 numGeometricModels;
#ifdef OSPRAY_ENABLE_VOLUMES
VolumetricModel **volumetricModels;
int32 numVolumetricModels;
#endif
GeometricModel **clipModels;
int32 numClipModels;
RTCScene testScene;
#ifdef __cplusplus
Group()
: geometricModels(nullptr),
numGeometricModels(0),
#ifdef OSPRAY_ENABLE_VOLUMES
volumetricModels(nullptr),
numVolumetricModels(0),
#endif
clipModels(nullptr),
numClipModels(0)
{}
modules/cpu/common/Instance.h
+1 -1
View File
@@ -7,7 +7,7 @@
#include "Group.h"
#include "MotionTransform.h"
// embree
#include "embree3/rtcore.h"
#include "embree4/rtcore.h"
// ispc shared
#include "InstanceShared.h"
modules/cpu/common/Instance.ih
+26 -2
View File
@@ -12,19 +12,26 @@
#include "GroupShared.h"
#include "InstanceShared.h"
OSPRAY_BEGIN_ISPC_NAMESPACE
struct Renderer;
inline uniform AffineSpace3f Instance_getTransform(
const Instance *uniform self, const uniform float time)
{
return getInterpolatedTransform(self->geom, time);
// TODO: rtcGetGeometryTransform is missing the SYCL_EXTERNAL attribute
// Is there a different function we should call?
return self->xfm;
// return getInterpolatedTransform(self->geom, time);
}
#ifndef OSPRAY_TARGET_DPCPP
inline AffineSpace3f Instance_getTransform(
const Instance *uniform self, const float time)
{
return getInterpolatedTransform(self->geom, time);
}
#endif
inline void Instance_postIntersect(const Instance *uniform self,
const Renderer *uniform renderer,
@@ -35,8 +42,15 @@ inline void Instance_postIntersect(const Instance *uniform self,
{
GeometricModel **uniform models =
clip ? self->group->clipModels : self->group->geometricModels;
foreach_unique (geomID in ray.geomID)
#ifdef OSPRAY_TARGET_DPCPP
{
// DPC++ execution in a thread is scalar
const uint32 geomID = ray.geomID;
#else
foreach_unique (geomID in ray.geomID) {
#endif
GeometricModel_postIntersect(*(models + geomID), renderer, dg, ray, flags);
}
dg.instID =
(self->userID == RTC_INVALID_GEOMETRY_ID) ? ray.instID : self->userID;
@@ -44,10 +58,17 @@ inline void Instance_postIntersect(const Instance *uniform self,
uniform AffineSpace3f uxfm = self->xfm;
uniform AffineSpace3f rcp_uxfm = self->rcp_xfm;
#if 0
if ((flags & DG_MOTIONBLUR) == DG_MOTIONBLUR && self->motionBlur) {
uniform float utime;
uniform int cnt = 0;
#ifdef OSPRAY_TARGET_DPCPP
{
// DPC++ execution in a thread is scalar
const float t = ray.time;
#else
foreach_unique (t in ray.time) {
#endif
utime = t;
cnt++;
}
@@ -73,6 +94,7 @@ inline void Instance_postIntersect(const Instance *uniform self,
rcp_uxfm = rcp(uxfm);
}
}
#endif
dg.lP = xfmPoint(rcp_uxfm, dg.P);
dg.Ns = xfmVector(transposed(rcp_uxfm.l), dg.Ns);
@@ -86,3 +108,5 @@ inline void Instance_postIntersect(const Instance *uniform self,
dg.dPdt = xfmVector(uxfm, dg.dPdt);
}
}
OSPRAY_END_ISPC_NAMESPACE
modules/cpu/common/Intersect.ih
+12 -2
View File
@@ -3,10 +3,19 @@
#pragma once
#ifdef OSPRAY_TARGET_DPCPP
#include "common/rkcommonDPCPPWrappers.h"
#include "rkcommon/math/box.h"
using namespace rkcommon;
using namespace rkcommon::math;
#else
#include "rkcommon/math/box.ih"
#endif
// Ray intersection structures //////////////////////////////////////////////
OSPRAY_BEGIN_ISPC_NAMESPACE
struct Hit
{
bool hit;
@@ -33,7 +42,7 @@ inline Intersections intersectSphere(const vec3f &rayOrg,
isect.entry.hit = false;
isect.exit.hit = false;
isect.entry.t = inf;
isect.exit.t = -inf;
isect.exit.t = -(float)inf;
const vec3f d = rayDir;
const float rd2 = 1.0f / dot(d, d); // 1/a
@@ -85,7 +94,7 @@ inline Intersections intersectCylinder(const vec3f &rayOrg,
isect.entry.hit = false;
isect.exit.hit = false;
isect.entry.t = inf;
isect.exit.t = -inf;
isect.exit.t = -(float)inf;
const uniform vec3f cZ = v1 - v0;
const vec3f q = rayOrg - v0;
@@ -233,3 +242,4 @@ inline Hit intersectPlane(
return hit;
}
OSPRAY_END_ISPC_NAMESPACE
modules/cpu/common/Managed.cpp
+12
View File
@@ -3,7 +3,9 @@
#include "Managed.h"
#include "Data.h"
#ifndef OSPRAY_TARGET_DPCPP
#include "common/OSPCommon_ispc.h"
#endif
namespace ospray {
@@ -21,6 +23,16 @@ ManagedObject::~ManagedObject()
void ManagedObject::commit() {}
#ifdef OSPRAY_TARGET_DPCPP
/*
void ManagedObject::setGPUFunctionPtrs(sycl::queue &)
{
std::cout << "MOST LIKELY TODO: " << toString()
<< " does not override setGPUFunctionPtrs\n";
}
*/
#endif
std::string ManagedObject::toString() const
{
return "ospray::ManagedObject";
modules/cpu/common/Managed.h
+8
View File
@@ -3,6 +3,10 @@
#pragma once
#ifdef OSPRAY_TARGET_DPCPP
#include <CL/sycl.hpp>
#endif
#include "rkcommon/utility/Any.h"
#include "rkcommon/utility/Optional.h"
#include "rkcommon/utility/ParameterizedObject.h"
@@ -29,6 +33,10 @@ struct OSPRAY_SDK_INTERFACE ManagedObject : public memory::RefCount,
virtual void commit();
#ifdef OSPRAY_TARGET_DPCPP
// virtual void setGPUFunctionPtrs(sycl::queue &syclQueue);
#endif
virtual std::string toString() const;
template <typename T>
modules/cpu/common/MotionTransform.h
+1 -1
View File
@@ -6,7 +6,7 @@
// ospray stuff
#include "./Group.h"
// embree
#include "embree3/rtcore.h"
#include "embree4/rtcore.h"
namespace ospray {
modules/cpu/common/MotionTransform.ih
+18 -3
View File
@@ -3,17 +3,30 @@
#pragma once
#include "embree3/rtcore.isph"
#include "math/AffineSpace.ih"
#ifdef OSPRAY_TARGET_DPCPP
#include "rkcommon/math/AffineSpace.h"
using namespace rkcommon;
using namespace rkcommon::math;
#include "embree4/rtcore.h"
#else
#include "embree4/rtcore.isph"
#include "rkcommon/math/AffineSpace.ih"
#endif
OSPRAY_BEGIN_ISPC_NAMESPACE
inline uniform AffineSpace3f getInterpolatedTransform(
const RTCGeometry geom, const uniform float time)
{
uniform AffineSpace3f xfm;
rtcGetGeometryTransform(geom, time, RTC_FORMAT_FLOAT3X4_COLUMN_MAJOR, &xfm);
// TODO: This function is missing the SYCL_EXTERNAL attribute?
// rtcGetGeometryTransform(geom, time, RTC_FORMAT_FLOAT3X4_COLUMN_MAJOR,
// &xfm);
return xfm;
}
#ifndef OSPRAY_TARGET_DPCPP
inline AffineSpace3f getInterpolatedTransform(
const RTCGeometry geom, const float time)
{
@@ -31,3 +44,5 @@ inline AffineSpace3f getInterpolatedTransform(
*/
return xfm;
}
#endif
OSPRAY_END_ISPC_NAMESPACE
modules/cpu/common/OSPCommon.ih
+43 -1
View File
@@ -3,16 +3,51 @@
#pragma once
#include "embree3/rtcore.isph"
#ifdef OSPRAY_TARGET_DPCPP
#include <CL/sycl.hpp>
#include "embree4/rtcore.h"
#else
#include "embree4/rtcore.isph"
#endif
#include "OSPConfig.h"
#ifdef OSPRAY_TARGET_DPCPP
#include <cstdint>
#include "rkcommon/math/math.h"
typedef uint64_t uint64;
typedef uint32_t uint32;
typedef uint16_t uint16;
typedef uint8_t uint8;
typedef int64_t int64;
typedef int32_t int32;
typedef int16_t int16;
typedef int8_t int8;
#else
#ifndef ISPC_UINT_IS_DEFINED
typedef unsigned int64 uint64;
typedef unsigned int32 uint32;
typedef unsigned int16 uint16;
typedef unsigned int8 uint8;
#endif
#endif
#ifdef OSPRAY_TARGET_DPCPP
#include "rkcommon/platform.h"
#define LOG(x)
//#define PRINT(x)
//#define PRINT3(x)
#define PRINTU(x)
#define PRINT3U(x)
#else
#define LOG(x)
@@ -29,10 +64,16 @@ typedef unsigned int8 uint8;
extract(get(x, 0), count_trailing_zeros(lanemask())), \
extract(get(x, 1), count_trailing_zeros(lanemask())), \
extract(get(x, 2), count_trailing_zeros(lanemask())))
#endif
OSPRAY_BEGIN_ISPC_NAMESPACE
/*! a C++-callable 'delete' of ISPC-side allocated memory of uniform objects */
export void delete_uniform(void *uniform uptr);
OSPRAY_END_ISPC_NAMESPACE
#ifndef OSPRAY_TARGET_DPCPP
/*! 64-bit malloc. allows for alloc'ing memory larger than 64 bits */
extern "C" void *uniform malloc64(uniform uint64 size);
extern "C" void free64(void *uniform ptr);
@@ -41,3 +82,4 @@ extern "C" void free64(void *uniform ptr);
extern "C" void *uniform pushTLS(uniform uint64 size);
extern "C" void *uniform reallocTLS(void *uniform ptr, uniform uint64 size);
extern "C" void popTLS(void *uniform ptr);
#endif
modules/cpu/common/OSPCommon.ispc
+4
View File
@@ -3,7 +3,11 @@
#include "OSPCommon.ih"
OSPRAY_BEGIN_ISPC_NAMESPACE
export void delete_uniform(void *uniform uptr)
{
delete uptr;
}
OSPRAY_END_ISPC_NAMESPACE
modules/cpu/common/Ray.ih
+24 -8
View File
@@ -4,10 +4,24 @@
#pragma once
// ospray
#include "../math/AffineSpace.ih"
#ifdef OSPRAY_TARGET_DPCPP
#include "common/rkcommonDPCPPWrappers.h"
#include "rkcommon/math/AffineSpace.h"
#include "rkcommon/math/box.h"
#include "rkcommon/math/vec.h"
#include "embree4/rtcore.h"
using namespace rkcommon;
using namespace rkcommon::math;
#else
#include "rkcommon/math/AffineSpace.ih"
#include "rkcommon/math/vec.ih"
// embree
#include "embree3/rtcore.isph"
#include "embree4/rtcore.isph"
#endif
OSPRAY_BEGIN_ISPC_NAMESPACE
// NOTE(jda) - this MUST match Embree's RayHit structure layout!
struct Ray
@@ -21,9 +35,9 @@ struct Ray
float t; /*!< end of valid ray interval, or distance to hit point after
'intersect' */
int32 mask; //!< Used to mask out objects during traversal
int32 rayID;
int32 flags;
uint32 mask; //!< Used to mask out objects during traversal
uint32 rayID;
uint32 flags;
/* hit data */
vec3f Ng; /*! geometry normal*/
@@ -31,9 +45,9 @@ struct Ray
float u; //!< Barycentric u coordinate of hit
float v; //!< Barycentric v coordinate of hit
int primID; //!< primitive ID
int geomID; //!< geometry ID
int instID; //!< instance ID
uint32 primID; //!< primitive ID
uint32 geomID; //!< geometry ID
uint32 instID; //!< instance ID
};
struct EmbreeIntersectionContext
@@ -102,3 +116,5 @@ inline void transformRay(Ray &ray, AffineSpace3f xfm)
ray.dir = xfmVector(xfm, ray.dir);
ray.Ng = xfmVector(transposed(xfm.l), ray.Ng);
}
OSPRAY_END_ISPC_NAMESPACE
modules/cpu/common/VolumeIntervals.ih
+16 -2
View File
@@ -1,12 +1,23 @@
// Copyright 2009 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
#ifdef OSPRAY_ENABLE_VOLUMES
#pragma once
#include "../math/AffineSpace.ih"
#include "Instance.ih"
#include "OSPCommon.ih"
#ifdef OSPRAY_TARGET_DPCPP
// Note: Keep this here for now since we will need these includes later
#include "common/rkcommonDPCPPWrappers.h"
#include "rkcommon/math/AffineSpace.h"
#include "rkcommon/math/box.h"
#else
#include "rkcommon/math/AffineSpace.ih"
#include "rkcommon/math/box.ih"
#endif
OSPRAY_BEGIN_ISPC_NAMESPACE
struct VolumetricModel;
@@ -26,7 +37,7 @@ inline void initVolumeInterval(VolumeInterval &hit)
{
hit.volumetricModel = NULL;
hit.interval.lower = inf;
hit.interval.upper = -inf;
hit.interval.upper = -hit.interval.lower;
}
inline bool hasInterval(const VolumeInterval &vi)
@@ -53,3 +64,6 @@ inline void freeVolumeIntervals(VolumeIntervals &intervals)
popTLS(intervals.intervals);
intervals.intervals = NULL;
}
OSPRAY_END_ISPC_NAMESPACE
#endif
modules/cpu/common/World.cpp
+20 -3
View File
@@ -46,9 +46,13 @@ static void freeAndNullifyEmbreeScene(RTCScene &scene)
World::~World()
{
// TODO WILL :Debugging other issues and just using the group's scene for now,
// so this would double-free the geom scene
// Release Embree scenes
freeAndNullifyEmbreeScene(getSh()->embreeSceneHandleGeometries);
#ifdef OSPRAY_ENABLE_VOLUMES
freeAndNullifyEmbreeScene(getSh()->embreeSceneHandleVolumes);
#endif
freeAndNullifyEmbreeScene(getSh()->embreeSceneHandleClippers);
}
@@ -66,11 +70,15 @@ std::string World::toString() const
void World::commit()
{
RTCScene &esGeom = getSh()->embreeSceneHandleGeometries;
#ifdef OSPRAY_ENABLE_VOLUMES
RTCScene &esVol = getSh()->embreeSceneHandleVolumes;
#endif
RTCScene &esClip = getSh()->embreeSceneHandleClippers;
freeAndNullifyEmbreeScene(esGeom);
#ifdef OSPRAY_ENABLE_VOLUMES
freeAndNullifyEmbreeScene(esVol);
#endif
freeAndNullifyEmbreeScene(esClip);
scivisData = nullptr;
@@ -117,15 +125,20 @@ void World::commit()
unsigned int id = 0;
for (auto &&inst : *instances) {
getSh()->instances[id] = inst->getSh();
if (inst->group->sceneGeometries)
if (inst->group->sceneGeometries) {
addGeometryInstance(
esGeom, inst->group->sceneGeometries, inst, embreeDevice, id);
if (inst->group->sceneVolumes)
}
#ifdef OSPRAY_ENABLE_VOLUMES
if (inst->group->sceneVolumes) {
addGeometryInstance(
esVol, inst->group->sceneVolumes, inst, embreeDevice, id);
if (inst->group->sceneClippers)
}
#endif
if (inst->group->sceneClippers) {
addGeometryInstance(
esClip, inst->group->sceneClippers, inst, embreeDevice, id);
}
id++;
}
}
@@ -135,11 +148,13 @@ void World::commit()
rtcSetSceneBuildQuality(esGeom, buildQuality);
rtcCommitScene(esGeom);
}
#ifdef OSPRAY_ENABLE_VOLUMES
if (esVol) {
rtcSetSceneFlags(esVol, static_cast<RTCSceneFlags>(sceneFlags));
rtcSetSceneBuildQuality(esVol, buildQuality);
rtcCommitScene(esVol);
}
#endif
if (esClip) {
rtcSetSceneFlags(esClip,
static_cast<RTCSceneFlags>(
@@ -160,10 +175,12 @@ box3f World::getBounds() const
sceneBounds.extend(box3f(vec3f(&bounds.lower[0]), vec3f(&bounds.upper[0])));
}
#ifdef OSPRAY_ENABLE_VOLUMES
if (getSh()->embreeSceneHandleVolumes) {
rtcGetSceneBounds(getSh()->embreeSceneHandleVolumes, (RTCBounds *)&bounds);
sceneBounds.extend(box3f(vec3f(&bounds.lower[0]), vec3f(&bounds.upper[0])));
}
#endif
return sceneBounds;
}
modules/cpu/common/World.h
+1 -1
View File
@@ -9,7 +9,7 @@
// stl
#include <vector>
// embree
#include "embree3/rtcore.h"
#include "embree4/rtcore.h"
// ispc shared
#include "WorldShared.h"
modules/cpu/common/World.ih
+97 -26
View File
@@ -3,15 +3,29 @@
#pragma once
#include "embree4/rtcore_config.h"
// ospray stuff
#include "common/Clipping.ih"
#include "common/DifferentialGeometry.ih"
#include "common/Instance.ih"
#include "common/Ray.ih"
#include "common/VolumeIntervals.ih"
#include "geometry/GeometryType.ih"
#include "volume/Volume.ih"
// c++ shared
#include "WorldShared.h"
#ifdef OSPRAY_TARGET_DPCPP
using namespace ospray;
#endif
OSPRAY_BEGIN_ISPC_NAMESPACE
struct Renderer;
// TODO: Test/Remove: will making this SYCL_EXTERNAL help
// with the bad compile times?
inline void traceGeometryRay(const World *uniform world, varying Ray &ray)
{
// Skip if no geometries scene
@@ -20,11 +34,18 @@ inline void traceGeometryRay(const World *uniform world, varying Ray &ray)
uniform EmbreeIntersectionContext context;
rtcInitIntersectContext(&context.ectx);
rtcIntersectV(world->embreeSceneHandleGeometries,
uniform RTCIntersectArguments intersectArgs;
rtcInitIntersectArguments(&intersectArgs);
intersectArgs.intersect = (RTCIntersectFunctionN)Geometry_dispatch_intersect;
rtcIntersectExV(world->embreeSceneHandleGeometries,
&context.ectx,
(varying RTCRayHit * uniform) & ray);
(varying RTCRayHit * uniform) & ray,
&intersectArgs);
}
#ifdef OSPRAY_ENABLE_VOLUMES
inline void traceVolumeRay(
const World *uniform world, varying Ray &ray, VolumeInterval &interval)
{
@@ -38,9 +59,15 @@ inline void traceVolumeRay(
uniform EmbreeIntersectionContext context;
rtcInitIntersectContext(&context.ectx);
context.userPtr = &intervals;
rtcIntersectV(world->embreeSceneHandleVolumes,
uniform RTCIntersectArguments intersectArgs;
rtcInitIntersectArguments(&intersectArgs);
intersectArgs.intersect = (RTCIntersectFunctionN)Volume_intersect_kernel;
rtcIntersectExV(world->embreeSceneHandleVolumes,
&context.ectx,
(varying RTCRayHit * uniform) & ray);
(varying RTCRayHit * uniform) & ray,
&intersectArgs);
if (intervals.numVolumeIntervals > 0 && hasInterval(intervals.intervals[0])) {
interval.instance = *(world->instances + intervals.intervals[0].instID);
@@ -61,9 +88,15 @@ inline void traceVolumeRay(
uniform EmbreeIntersectionContext context;
rtcInitIntersectContext(&context.ectx);
context.userPtr = &intervals;
rtcIntersectV(world->embreeSceneHandleVolumes,
uniform RTCIntersectArguments intersectArgs;
rtcInitIntersectArguments(&intersectArgs);
intersectArgs.intersect = (RTCIntersectFunctionN)Volume_intersect_kernel;
rtcIntersectExV(world->embreeSceneHandleVolumes,
&context.ectx,
(varying RTCRayHit * uniform) & ray);
(varying RTCRayHit * uniform) & ray,
&intersectArgs);
if (intervals.numVolumeIntervals > 0) {
for (uniform int i = 0; i < intervals.numVolumeIntervals; ++i) {
@@ -75,6 +108,7 @@ inline void traceVolumeRay(
}
}
}
#endif
// Intersection context structure used for clipping geometries
struct UserIntersectionContextForClipping
@@ -86,20 +120,21 @@ struct UserIntersectionContextForClipping
varying ClippingHit hits[CLIPPING_HITS_MAX_COUNT];
};
unmasked void clippingIntersectionFilterV(
const struct RTCFilterFunctionNArguments *uniform args);
unmasked void clippingIntersectionFilterV(const void *uniform _args);
inline void traceClippingRay(const World *uniform world,
varying Ray &ray,
varying RayIntervals &rayIntervals)
{
// A scene with clipping geometries has to exist
if (!world->embreeSceneHandleClippers) {
rayIntervals.intervals[0] = make_box1f(ray.t0, ray.t);
rayIntervals.count = 1;
return;
}
// if (!world->embreeSceneHandleClippers) {
rayIntervals.intervals[0] = make_box1f(ray.t0, ray.t);
rayIntervals.count = 1;
return;
//}
// TODO:Clipping disabled for now
#if 0
// Create and initialize intersection context
UserIntersectionContextForClipping context;
rtcInitIntersectContext(&context.ectx);
@@ -149,13 +184,21 @@ inline void traceClippingRay(const World *uniform world,
}
// Start interval if previous one is closed and we are not in clipping area
#ifdef OSPRAY_TARGET_DPCPP
if (clippingDepth == 0 && intervalCompleted) {
#else
if (and(clippingDepth == 0, intervalCompleted)) {
#endif
rayIntervals.intervals[intervalId].lower = absT;
intervalCompleted = false;
}
// End interval if current one is not closed and we enter clipping area
#ifdef OSPRAY_TARGET_DPCPP
if (clippingDepth > 0 && !intervalCompleted) {
#else
if (and(clippingDepth > 0, !intervalCompleted)) {
#endif
rayIntervals.intervals[intervalId].upper = absT;
intervalCompleted = true;
intervalId++;
@@ -170,6 +213,7 @@ inline void traceClippingRay(const World *uniform world,
// Save number of ray intervals
rayIntervals.count = intervalId;
#endif
}
inline void traceGeometryRayIntervals(const World *uniform world,
@@ -187,7 +231,11 @@ inline void traceGeometryRayIntervals(const World *uniform world,
ray.t = rayIntervals.intervals[i].upper;
// Skip intervals outside of the ray range
#ifdef OSPRAY_TARGET_DPCPP
if (t < ray.t0 || ray.t < t0)
#else
if (or (t < ray.t0, ray.t < t0))
#endif
continue;
// Clip interval to the ray range
@@ -218,11 +266,11 @@ inline void traceRay(const World *uniform world, varying Ray &ray)
}
// Trace ray in clipping geometries scene, fill array with ray intervals
varying RayIntervals rayIntervals;
traceClippingRay(world, ray, rayIntervals);
// varying RayIntervals rayIntervals;
// traceClippingRay(world, ray, rayIntervals);
// Trace ray intervals
traceGeometryRayIntervals(world, ray, rayIntervals);
// traceGeometryRayIntervals(world, ray, rayIntervals);
}
inline bool isOccludedNoClipping(const World *uniform world, varying Ray &ray)
@@ -233,9 +281,21 @@ inline bool isOccludedNoClipping(const World *uniform world, varying Ray &ray)
uniform EmbreeIntersectionContext context;
rtcInitIntersectContext(&context.ectx);
rtcOccludedV(world->embreeSceneHandleGeometries,
uniform RTCIntersectArguments intersectArgs;
rtcInitIntersectArguments(&intersectArgs);
#ifdef ISPC
// Note: no union type in ISPC, intersect/occluded are a union in the
// RTCIntersectArguments struct.
intersectArgs.intersect = Geometry_dispatch_occluded;
#else
intersectArgs.occluded = (RTCOccludedFunctionN)Geometry_dispatch_occluded;
#endif
rtcOccludedExV(world->embreeSceneHandleGeometries,
&context.ectx,
(varying RTCRay * uniform) & ray);
(varying RTCRay * uniform) & ray,
&intersectArgs);
return ray.t < ray.t0;
}
@@ -262,19 +322,23 @@ inline bool areIntervalsOccluded(const World *uniform world,
inline bool isOccluded(const World *uniform world, varying Ray &ray)
{
// Fast path if no clipping geometry
if (!world->embreeSceneHandleClippers) {
// if (!world->embreeSceneHandleClippers) {
if (true) {
return isOccludedNoClipping(world, ray);
}
// Allocate array for ray intervals
varying RayIntervals rayIntervals;
rayIntervals.count = 0;
// TODO: use new Embree4 callback setup for clipping
/*
// Allocate array for ray intervals
varying RayIntervals rayIntervals;
rayIntervals.count = 0;
// Trace ray in clipping geometries scene, fill array with ray intervals
traceClippingRay(world, ray, rayIntervals);
// Trace ray in clipping geometries scene, fill array with ray intervals
traceClippingRay(world, ray, rayIntervals);
// Is there any occluder within given ray intervals
return areIntervalsOccluded(world, ray, rayIntervals);
// Is there any occluder within given ray intervals
return areIntervalsOccluded(world, ray, rayIntervals);
*/
}
/*! Perform post-intersect computations, i.e. fill the members of
@@ -307,7 +371,13 @@ inline void postIntersect(const World *uniform world,
dg.P = ray.org + ray.t * ray.dir;
dg.epsilon = 0.f; // per default no geometry-type specific epsilon
#ifdef OSPRAY_TARGET_DPCPP
{
// DPC++ execution in a thread is scalar
const uint32 instID = ray.instID;
#else
foreach_unique (instID in ray.instID) {
#endif
if (instID != RTC_INVALID_GEOMETRY_ID) {
Instance *uniform instance = *(world->instances + instID);
Instance_postIntersect(instance, renderer, dg, ray, flags);
@@ -351,3 +421,4 @@ inline void postIntersect(const World *uniform world,
#undef DG_NG_NORMALIZE
#undef DG_NS_NORMALIZE
}
OSPRAY_END_ISPC_NAMESPACE
modules/cpu/common/World.ispc
+16 -3
View File
@@ -1,15 +1,20 @@
// Copyright 2009 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
#include "fb/Tile.ih"
#include "World.ih"
#include "math/Distribution1D.ih"
#include "render/pathtracer/GeometryLight.ih"
// c++ shared
#include "lights/LightShared.h"
unmasked void clippingIntersectionFilterV(
const struct RTCFilterFunctionNArguments *uniform args)
OSPRAY_BEGIN_ISPC_NAMESPACE
unmasked void clippingIntersectionFilterV(const void *uniform _args)
{
const struct RTCFilterFunctionNArguments *uniform args =
(const struct RTCFilterFunctionNArguments *uniform)_args;
// Set execution mask
varying int *uniform valid = (varying int *uniform)args->valid;
if (*valid == 0)
@@ -43,8 +48,14 @@ unmasked void clippingIntersectionFilterV(
ray->primID = hit->primID;
ray->geomID = hit->geomID;
// We need to call postIntersect on clipping geometry
// We need to call postIntersect on clipping geometry
#ifdef OSPRAY_TARGET_DPCPP
{
// DPC++ execution in a thread is scalar
const uint32 instID = hit->instID[0];
#else
foreach_unique (instID in hit->instID[0]) {
#endif
// Call postIntersect to get shading normal
Instance *uniform instance = *(context->world->instances + instID);
Instance_postIntersect(instance, NULL, dg, *ray, DG_NG | DG_NS, true);
@@ -98,3 +109,5 @@ unmasked void clippingIntersectionFilterV(
context->hits[context->hitsCount] = cHit;
context->hitsCount++;
}
OSPRAY_END_ISPC_NAMESPACE
modules/cpu/common/WorldShared.h
+10
View File
@@ -6,6 +6,12 @@
#include "render/pathtracer/PathTracerDataShared.h"
#include "render/scivis/SciVisDataShared.h"
#ifdef __cplusplus
#include "embree4/rtcore.h"
#else
#include "embree4/rtcore.isph"
#endif
#ifdef __cplusplus
#include "common/StructShared.h"
namespace ispc {
@@ -19,7 +25,9 @@ struct World
int32 numInvertedClippers;
RTCScene embreeSceneHandleGeometries;
#ifdef OSPRAY_ENABLE_VOLUMES
RTCScene embreeSceneHandleVolumes;
#endif
RTCScene embreeSceneHandleClippers;
SciVisData *scivisData;
@@ -30,7 +38,9 @@ struct World
: instances(nullptr),
numInvertedClippers(0),
embreeSceneHandleGeometries(nullptr),
#ifdef OSPRAY_ENABLE_VOLUMES
embreeSceneHandleVolumes(nullptr),
#endif
embreeSceneHandleClippers(nullptr),
scivisData(nullptr),
pathtracerData(nullptr)
modules/cpu/common/rkcommonDPCPPWrappers.h
+21
View File
@@ -0,0 +1,21 @@
// Copyright 2022 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
#pragma once
#ifdef OSPRAY_TARGET_DPCPP
#include <CL/sycl.hpp>
#include <type_traits>
#include "common/OSPCommon.ih"
#include "rkcommon/math/AffineSpace.h"
#include "rkcommon/math/LinearSpace.h"
#include "rkcommon/math/math.h"
#include "rkcommon/math/rkmath.h"
#include "rkcommon/math/vec.h"
#include "rkcommon/utility/random.h"
using namespace rkcommon::utility;
using namespace rkcommon::math;
#endif
modules/cpu/fb/FrameBuffer.cpp
+13
View File
@@ -2,7 +2,9 @@
// SPDX-License-Identifier: Apache-2.0
#include "FrameBuffer.h"
#ifndef OSPRAY_TARGET_DPCPP
#include "ISPCDevice_ispc.h"
#endif
#include "OSPConfig.h"
namespace {
@@ -50,6 +52,10 @@ FrameBuffer::FrameBuffer(api::ISPCDevice &device,
getSh()->channels = channels;
getSh()->colorBufferFormat = _colorBufferFormat;
#ifdef OSPRAY_TARGET_DPCPP
// Note: using 2x2, 4x4, etc doesn't change perf much
vec2i renderTaskSize(1);
#else
#if OSPRAY_RENDER_TASK_SIZE == -1
// Compute render task size based on the simd width to get as "square" as
// possible a task size that has simdWidth pixels
@@ -63,6 +69,7 @@ FrameBuffer::FrameBuffer(api::ISPCDevice &device,
// Note: we could also allow changing this at runtime if we want to add this
// to the API
vec2i renderTaskSize(OSPRAY_RENDER_TASK_SIZE);
#endif
#endif
getSh()->renderTaskSize = renderTaskSize;
}
@@ -95,6 +102,7 @@ float FrameBuffer::getVariance() const
void FrameBuffer::beginFrame()
{
cancelRender = false;
// TODO: Maybe better as a kernel to avoid USM thrash to host
getSh()->cancelRender = 0;
getSh()->numPixelsRendered = 0;
getSh()->frameID++;
@@ -130,8 +138,13 @@ void FrameBuffer::waitForEvent(OSPSyncEvent event) const
float FrameBuffer::getCurrentProgress() const
{
#ifdef OSPRAY_TARGET_DPCPP
// TODO: Continually polling this will cause a lot of USM thrashing
return 0.f;
#else
return static_cast<float>(getSh()->numPixelsRendered)
/ getNumPixels().long_product();
#endif
}
void FrameBuffer::cancelFrame()
modules/cpu/fb/FrameBuffer.ih
+8 -1
View File
@@ -6,18 +6,23 @@
#include "FrameBufferShared.h"
#include "common/OSPCommon.ih"
#ifndef OSPRAY_TARGET_DPCPP
extern "C" uint32 *uniform getThreadNumPixelsRendered();
extern "C" int *uniform getThreadLastFrameID();
#endif
OSPRAY_BEGIN_ISPC_NAMESPACE
// Run the chain of pixel ops (if any) to produce the final color to be written
// into the framebuffer
vec4f FrameBuffer_runPixelOps(FrameBuffer *uniform self,
SYCL_EXTERNAL vec4f FrameBuffer_runPixelOps(FrameBuffer *uniform self,
const vec4f &color,
const float depth,
const vec3f &normal,
const vec3f &albedo);
#ifndef OSPRAY_TARGET_DPCPP
inline void FrameBuffer_updateProgress(
FrameBuffer *uniform self, const uniform uint32 numPixelsRendered)
{
@@ -42,3 +47,5 @@ inline void FrameBuffer_updateProgress(
*pixelsRendered = 0;
}
}
#endif
OSPRAY_END_ISPC_NAMESPACE
modules/cpu/fb/FrameBuffer.ispc
+4 -1
View File
@@ -3,7 +3,9 @@
#include "fb/FrameBuffer.ih"
vec4f FrameBuffer_runPixelOps(FrameBuffer *uniform self,
OSPRAY_BEGIN_ISPC_NAMESPACE
SYCL_EXTERNAL vec4f FrameBuffer_runPixelOps(FrameBuffer *uniform self,
const vec4f &color,
const float depth,
const vec3f &normal,
@@ -16,3 +18,4 @@ vec4f FrameBuffer_runPixelOps(FrameBuffer *uniform self,
}
return outColor;
}
OSPRAY_END_ISPC_NAMESPACE
modules/cpu/fb/FrameBufferShared.h
+13 -4
View File
@@ -3,18 +3,25 @@
#pragma once
#include "fb/FrameBufferType.ih"
#include "fb/ImageOpShared.h"
#include "ospray/OSPEnums.h"
#if !defined(__cplusplus) || defined(OSPRAY_TARGET_DPCPP)
#include "fb/RenderTaskDesc.ih"
#include "render/ScreenSample.ih"
#endif
#ifdef __cplusplus
#include "common/StructShared.h"
namespace ispc {
#endif // __cplusplus
#if defined(__cplusplus) && !defined(OSPRAY_TARGET_DPCPP)
typedef void *FrameBuffer_accumulateSampleFct;
typedef void *FrameBuffer_getRenderTaskDescFct;
typedef void *FrameBuffer_completeTaskFct;
#else
#include "fb/RenderTaskDesc.ih"
#include "render/ScreenSample.ih"
struct FrameBuffer;
@@ -27,7 +34,7 @@ typedef uniform RenderTaskDesc (*FrameBuffer_getRenderTaskDescFct)(
typedef void (*FrameBuffer_completeTaskFct)(
FrameBuffer *uniform fb, const uniform RenderTaskDesc &taskDesc);
#endif // __cplusplus
#endif
/* The ISPC-side FrameBuffer allows tasks to write directly to the framebuffer
* memory from ISPC. Given the set of task IDs to be rendered, a renderer must:
@@ -42,6 +49,7 @@ typedef void (*FrameBuffer_completeTaskFct)(
*/
struct FrameBuffer
{
FrameBufferType type;
/* Get the task description for a given render task ID. The task description
* stores the region that should be rendered for the task and its accumID
*/
@@ -86,7 +94,8 @@ struct FrameBuffer
#ifdef __cplusplus
FrameBuffer()
: getRenderTaskDesc(nullptr),
: type(FRAMEBUFFER_TYPE_UNKNOWN),
getRenderTaskDesc(nullptr),
accumulateSample(nullptr),
completeTask(nullptr),
size(0),
modules/cpu/fb/FrameBufferType.ih
+44
View File
@@ -0,0 +1,44 @@
// Copyright 2022 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
#pragma once
#ifdef OSPRAY_TARGET_DPCPP
#include <CL/sycl.hpp>
#include "common/rkcommonDPCPPWrappers.h"
#include "rkcommon/platform.h"
#endif
#ifdef __cplusplus
namespace ispc {
#endif // __cplusplus
struct FrameBuffer;
struct RenderTaskDesc;
struct ScreenSample;
enum FrameBufferType
{
FRAMEBUFFER_TYPE_LOCAL,
FRAMEBUFFER_TYPE_SPARSE,
FRAMEBUFFER_TYPE_UNKNOWN,
};
#if defined(ISPC) || defined(OSPRAY_TARGET_DPCPP)
SYCL_EXTERNAL __noinline uniform RenderTaskDesc
FrameBuffer_dispatch_getRenderTaskDesc(
FrameBuffer *uniform fb, const uniform uint32 taskID);
SYCL_EXTERNAL __noinline void FrameBuffer_dispatch_accumulateSample(
FrameBuffer *uniform fb,
const varying ScreenSample &screenSample,
uniform RenderTaskDesc &taskDesc);
SYCL_EXTERNAL __noinline void FrameBuffer_dispatch_completeTask(
FrameBuffer *uniform fb, const uniform RenderTaskDesc &taskDesc);
#endif
#ifdef __cplusplus
}
#endif
modules/cpu/fb/FrameBufferType.ispc
+76
View File
@@ -0,0 +1,76 @@
// Copyright 2022 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
#include "fb/FrameBufferShared.h"
#include "fb/FrameBufferType.ih"
#include "fb/RenderTaskDesc.ih"
#include "render/ScreenSample.ih"
#include "fb/LocalFB.ih"
#include "fb/SparseFB.ih"
OSPRAY_BEGIN_ISPC_NAMESPACE
SYCL_EXTERNAL __noinline uniform RenderTaskDesc
FrameBuffer_dispatch_getRenderTaskDesc(
FrameBuffer *uniform fb, const uniform uint32 taskID)
{
switch (fb->type) {
case FRAMEBUFFER_TYPE_LOCAL:
return LocalFB_getRenderTaskDesc(fb, taskID);
case FRAMEBUFFER_TYPE_SPARSE:
return SparseFB_getRenderTaskDesc(fb, taskID);
default:
#ifndef OSPRAY_TARGET_DPCPP
return fb->getRenderTaskDesc(fb, taskID);
#else
break;
#endif
}
// TODO: Should be an error here
uniform RenderTaskDesc rt;
rt.region.lower = make_vec2i(0, 0);
rt.region.upper = rt.region.lower;
return rt;
}
SYCL_EXTERNAL __noinline void FrameBuffer_dispatch_accumulateSample(
FrameBuffer *uniform fb,
const varying ScreenSample &screenSample,
uniform RenderTaskDesc &taskDesc)
{
switch (fb->type) {
case FRAMEBUFFER_TYPE_LOCAL:
LocalFB_accumulateSample(fb, screenSample, taskDesc);
break;
case FRAMEBUFFER_TYPE_SPARSE:
SparseFB_accumulateSample(fb, screenSample, taskDesc);
break;
default:
#ifndef OSPRAY_TARGET_DPCPP
fb->accumulateSample(fb, screenSample, taskDesc);
#endif
break;
}
}
SYCL_EXTERNAL __noinline void FrameBuffer_dispatch_completeTask(
FrameBuffer *uniform fb, const uniform RenderTaskDesc &taskDesc)
{
switch (fb->type) {
case FRAMEBUFFER_TYPE_LOCAL:
LocalFB_completeTask(fb, taskDesc);
break;
case FRAMEBUFFER_TYPE_SPARSE:
SparseFB_completeTask(fb, taskDesc);
break;
default:
#ifndef OSPRAY_TARGET_DPCPP
fb->completeTask(fb, taskDesc);
#endif
break;
}
}
OSPRAY_END_ISPC_NAMESPACE
modules/cpu/fb/FrameBufferView.ih
+11 -1
View File
@@ -3,10 +3,19 @@
#pragma once
#include "math/AffineSpace.ih"
#ifdef OSPRAY_TARGET_DPCPP
#include "rkcommon/math/AffineSpace.h"
using namespace rkcommon;
using namespace rkcommon::math;
#else
#include "rkcommon/math/AffineSpace.ih"
#endif
// c++ shared
#include "FrameBufferShared.h"
OSPRAY_BEGIN_ISPC_NAMESPACE
struct FrameBufferView
{
// The total dimensions of the global framebuffer
@@ -28,3 +37,4 @@ struct FrameBufferView
// accumulated albedo, one RGBF32 per pixel
vec3f *albedoBuffer;
};
OSPRAY_END_ISPC_NAMESPACE
modules/cpu/fb/ImageOpShared.h
+6 -1
View File
@@ -5,10 +5,15 @@
#ifdef __cplusplus
#include "common/StructShared.h"
#include "rkcommon/math/vec.h"
namespace ispc {
typedef void *LivePixelOp_processPixel;
#else
#include "rkcommon/math/vec.ih"
#endif
#if defined(__cplusplus) && !defined(OSPRAY_TARGET_DPCPP)
typedef void *LivePixelOp_processPixel;
#else
struct LivePixelOp;
typedef vec4f (*LivePixelOp_processPixel)(const LivePixelOp *uniform self,
modules/cpu/fb/LocalFB.cpp
+92 -46
View File
@@ -5,9 +5,12 @@
#include <algorithm>
#include <iostream>
#include <iterator>
#include <numeric>
#include "ImageOp.h"
#include "SparseFB.h"
#ifndef OSPRAY_TARGET_DPCPP
#include "fb/LocalFB_ispc.h"
#endif
#include "render/util.h"
#include "rkcommon/common.h"
#include "rkcommon/tasking/parallel_for.h"
@@ -21,76 +24,97 @@ LocalFrameBuffer::LocalFrameBuffer(api::ISPCDevice &device,
: AddStructShared(
device.getIspcrtDevice(), device, _size, _colorBufferFormat, channels),
numRenderTasks(divRoundUp(size, getRenderTaskSize())),
taskErrorRegion(hasVarianceBuffer ? getNumRenderTasks() : vec2i(0))
taskErrorRegion(device.getIspcrtDevice(),
hasVarianceBuffer ? getNumRenderTasks() : vec2i(0))
{
const size_t pixelBytes = sizeOf(_colorBufferFormat);
const size_t numPixels = _size.long_product();
colorBuffer.resize(pixelBytes * numPixels);
if (getColorBufferFormat() != OSP_FB_NONE) {
colorBuffer = make_buffer_shared_unique<uint8_t>(
device.getIspcrtDevice(), pixelBytes * numPixels);
}
if (hasDepthBuffer)
depthBuffer.resize(numPixels);
depthBuffer =
make_buffer_shared_unique<float>(device.getIspcrtDevice(), numPixels);
if (hasAccumBuffer) {
accumBuffer.resize(numPixels);
accumBuffer =
make_buffer_shared_unique<vec4f>(device.getIspcrtDevice(), numPixels);
taskAccumID.resize(getTotalRenderTasks(), 0);
taskAccumID = make_buffer_shared_unique<int32_t>(
device.getIspcrtDevice(), getTotalRenderTasks());
std::memset(taskAccumID->data(), 0, taskAccumID->size() * sizeof(int32_t));
}
if (hasVarianceBuffer)
varianceBuffer.resize(numPixels);
varianceBuffer =
make_buffer_shared_unique<vec4f>(device.getIspcrtDevice(), numPixels);
if (hasNormalBuffer)
normalBuffer.resize(numPixels);
normalBuffer =
make_buffer_shared_unique<vec3f>(device.getIspcrtDevice(), numPixels);
if (hasAlbedoBuffer)
albedoBuffer.resize(numPixels);
albedoBuffer =
make_buffer_shared_unique<vec3f>(device.getIspcrtDevice(), numPixels);
if (hasPrimitiveIDBuffer)
primitiveIDBuffer.resize(numPixels);
primitiveIDBuffer = make_buffer_shared_unique<uint32_t>(
device.getIspcrtDevice(), numPixels);
if (hasObjectIDBuffer)
objectIDBuffer.resize(numPixels);
objectIDBuffer = make_buffer_shared_unique<uint32_t>(
device.getIspcrtDevice(), numPixels);
if (hasInstanceIDBuffer)
instanceIDBuffer.resize(numPixels);
instanceIDBuffer = make_buffer_shared_unique<uint32_t>(
device.getIspcrtDevice(), numPixels);
// TODO: Better way to pass the task IDs that doesn't require just storing
// them all? Maybe as blocks/tiles similar to when we just had tiles? Will
// make task ID lookup more expensive for sparse case though
renderTaskIDs.reserve(getTotalRenderTasks());
for (uint32_t i = 0; i < getTotalRenderTasks(); ++i) {
renderTaskIDs.push_back(i);
}
renderTaskIDs = make_buffer_shared_unique<uint32_t>(
device.getIspcrtDevice(), getTotalRenderTasks());
std::iota(renderTaskIDs->begin(), renderTaskIDs->end(), 0);
// TODO: Could use TBB parallel sort here if it's exposed through the rkcommon
// tasking system
std::sort(renderTaskIDs.begin(),
renderTaskIDs.end(),
std::sort(renderTaskIDs->begin(),
renderTaskIDs->end(),
[&](const uint32_t &a, const uint32_t &b) {
const vec2i p_a = getTaskStartPos(a);
const vec2i p_b = getTaskStartPos(b);
return interleaveZOrder(p_a.x, p_a.y) < interleaveZOrder(p_b.x, p_b.y);
});
#ifndef OSPRAY_TARGET_DPCPP
getSh()->super.accumulateSample =
ispc::LocalFrameBuffer_accumulateSample_addr();
reinterpret_cast<ispc::FrameBuffer_accumulateSampleFct>(
ispc::LocalFrameBuffer_accumulateSample_addr());
getSh()->super.getRenderTaskDesc =
ispc::LocalFrameBuffer_getRenderTaskDesc_addr();
getSh()->super.completeTask = ispc::LocalFrameBuffer_completeTask_addr();
reinterpret_cast<ispc::FrameBuffer_getRenderTaskDescFct>(
ispc::LocalFrameBuffer_getRenderTaskDesc_addr());
getSh()->super.completeTask =
reinterpret_cast<ispc::FrameBuffer_completeTaskFct>(
ispc::LocalFrameBuffer_completeTask_addr());
#endif
getSh()->colorBuffer = getDataSafe(colorBuffer);
getSh()->depthBuffer = getDataSafe(depthBuffer);
getSh()->accumBuffer = getDataSafe(accumBuffer);
getSh()->varianceBuffer = getDataSafe(varianceBuffer);
getSh()->normalBuffer = getDataSafe(normalBuffer);
getSh()->albedoBuffer = getDataSafe(albedoBuffer);
getSh()->taskAccumID = getDataSafe(taskAccumID);
getSh()->colorBuffer = colorBuffer ? colorBuffer->data() : nullptr;
getSh()->depthBuffer = depthBuffer ? depthBuffer->data() : nullptr;
getSh()->accumBuffer = accumBuffer ? accumBuffer->data() : nullptr;
getSh()->varianceBuffer = varianceBuffer ? varianceBuffer->data() : nullptr;
getSh()->normalBuffer = normalBuffer ? normalBuffer->data() : nullptr;
getSh()->albedoBuffer = albedoBuffer ? albedoBuffer->data() : nullptr;
getSh()->taskAccumID = taskAccumID ? taskAccumID->data() : nullptr;
getSh()->taskRegionError = taskErrorRegion.errorBuffer();
getSh()->numRenderTasks = numRenderTasks;
getSh()->primitiveIDBuffer = getDataSafe(primitiveIDBuffer);
getSh()->objectIDBuffer = getDataSafe(objectIDBuffer);
getSh()->instanceIDBuffer = getDataSafe(instanceIDBuffer);
getSh()->primitiveIDBuffer =
primitiveIDBuffer ? primitiveIDBuffer->data() : nullptr;
getSh()->objectIDBuffer = objectIDBuffer ? objectIDBuffer->data() : nullptr;
getSh()->instanceIDBuffer =
instanceIDBuffer ? instanceIDBuffer->data() : nullptr;
}
void LocalFrameBuffer::commit()
@@ -101,10 +125,10 @@ void LocalFrameBuffer::commit()
if (imageOpData) {
FrameBufferView fbv(this,
getColorBufferFormat(),
getDataSafe(colorBuffer),
getDataSafe(depthBuffer),
getDataSafe(normalBuffer),
getDataSafe(albedoBuffer));
colorBuffer ? colorBuffer->data() : nullptr,
depthBuffer ? depthBuffer->data() : nullptr,
normalBuffer ? normalBuffer->data() : nullptr,
albedoBuffer ? albedoBuffer->data() : nullptr);
for (auto &&obj : *imageOpData)
imageOps.push_back(obj->attach(fbv));
@@ -112,6 +136,23 @@ void LocalFrameBuffer::commit()
prepareImageOps();
}
#ifdef OSPRAY_TARGET_DPCPP
/*
void LocalFrameBuffer::setGPUFunctionPtrs(sycl::queue &syclQueue)
{
auto *sSh = getSh();
auto event = syclQueue.submit([&](sycl::handler &cgh) {
cgh.parallel_for(1, [=](cl::sycl::id<1>) RTC_SYCL_KERNEL {
sSh->super.accumulateSample = ispc::LocalFB_accumulateSample;
sSh->super.getRenderTaskDesc = ispc::LocalFB_getRenderTaskDesc;
sSh->super.completeTask = ispc::LocalFB_completeTask;
});
});
event.wait();
}
*/
#endif
vec2i LocalFrameBuffer::getNumRenderTasks() const
{
return numRenderTasks;
@@ -124,7 +165,8 @@ uint32_t LocalFrameBuffer::getTotalRenderTasks() const
utility::ArrayView<uint32_t> LocalFrameBuffer::getRenderTaskIDs()
{
return utility::ArrayView<uint32_t>(renderTaskIDs);
return utility::ArrayView<uint32_t>(
renderTaskIDs->data(), renderTaskIDs->size());
}
std::string LocalFrameBuffer::toString() const
@@ -138,7 +180,9 @@ void LocalFrameBuffer::clear()
// it is only necessary to reset the accumID,
// LocalFrameBuffer_accumulateTile takes care of clearing the
// accumulating buffers
std::fill(taskAccumID.begin(), taskAccumID.end(), 0);
if (taskAccumID) {
std::fill(taskAccumID->begin(), taskAccumID->end(), 0);
}
// always also clear error buffer (if present)
if (hasVarianceBuffer) {
@@ -160,7 +204,7 @@ void LocalFrameBuffer::writeTiles(const containers::AlignedVector<Tile> &tiles)
if (hasAlbedoBuffer) {
ispc::LocalFrameBuffer_writeAuxTile(getSh(),
tile,
(ispc::vec3f *)albedoBuffer.data(),
(ispc::vec3f *)albedoBuffer->data(),
tile->ar,
tile->ag,
tile->ab);
@@ -184,12 +228,12 @@ void LocalFrameBuffer::writeTiles(const containers::AlignedVector<Tile> &tiles)
if (hasNormalBuffer)
ispc::LocalFrameBuffer_writeAuxTile(getSh(),
tile,
(ispc::vec3f *)normalBuffer.data(),
(ispc::vec3f *)normalBuffer->data(),
tile->nx,
tile->ny,
tile->nz);
if (!colorBuffer.empty()) {
if (colorBuffer) {
switch (getColorBufferFormat()) {
case OSP_FB_RGBA8:
ispc::LocalFrameBuffer_writeTile_RGBA8(getSh(), tile);
@@ -277,29 +321,31 @@ void LocalFrameBuffer::endFrame(
const void *LocalFrameBuffer::mapBuffer(OSPFrameBufferChannel channel)
{
// TODO: Mapping the USM back to the app like this will cause a lot of USM
// thrashing between host/device
const void *buf = nullptr;
switch (channel) {
case OSP_FB_COLOR:
buf = getDataSafe(colorBuffer);
buf = colorBuffer ? colorBuffer->data() : nullptr;
break;
case OSP_FB_DEPTH:
buf = getDataSafe(depthBuffer);
buf = depthBuffer ? depthBuffer->data() : nullptr;
break;
case OSP_FB_NORMAL:
buf = getDataSafe(normalBuffer);
buf = normalBuffer ? normalBuffer->data() : nullptr;
break;
case OSP_FB_ALBEDO:
buf = getDataSafe(albedoBuffer);
buf = albedoBuffer ? albedoBuffer->data() : nullptr;
break;
case OSP_FB_ID_PRIMITIVE:
buf = getDataSafe(primitiveIDBuffer);
buf = primitiveIDBuffer ? primitiveIDBuffer->data() : nullptr;
break;
case OSP_FB_ID_OBJECT:
buf = getDataSafe(objectIDBuffer);
buf = objectIDBuffer ? objectIDBuffer->data() : nullptr;
break;
case OSP_FB_ID_INSTANCE:
buf = getDataSafe(instanceIDBuffer);
buf = instanceIDBuffer ? instanceIDBuffer->data() : nullptr;
break;
default:
break;
modules/cpu/fb/LocalFB.h
+25 -11
View File
@@ -3,7 +3,12 @@
#pragma once
#ifdef OSPRAY_TARGET_DPCPP
#include <CL/sycl.hpp>
#endif
// ospray
#include "common/BufferShared.h"
#include "fb/FrameBuffer.h"
#include "fb/TaskError.h"
// rkcommon
@@ -28,6 +33,10 @@ struct OSPRAY_SDK_INTERFACE LocalFrameBuffer
virtual void commit() override;
#ifdef OSPRAY_TARGET_DPCPP
// void setGPUFunctionPtrs(sycl::queue &syclQueue) override;
#endif
// Return the number of render tasks in the x and y direction
// This is the kernel launch dims to render the image
virtual vec2i getNumRenderTasks() const override;
@@ -71,22 +80,26 @@ struct OSPRAY_SDK_INTERFACE LocalFrameBuffer
// NOTE: All per-pixel data is only allocated if the corresponding channel
// flag was passed on construction
// TODO/NOTE: we use unique ptr here because the default Array constructor
// allocates a single element "array". We probably want to change this in
// ispcrt to have behavior more like std::vector/other containers
// format depends on FrameBuffer::colorBufferFormat
std::vector<uint8_t> colorBuffer;
std::unique_ptr<BufferShared<uint8_t>> colorBuffer;
// one float per pixel
containers::AlignedVector<float> depthBuffer;
std::unique_ptr<BufferShared<float>> depthBuffer;
// one RGBA per pixel
containers::AlignedVector<vec4f> accumBuffer;
std::unique_ptr<BufferShared<vec4f>> accumBuffer;
// one RGBA per pixel, accumulates every other sample, for variance estimation
containers::AlignedVector<vec4f> varianceBuffer;
std::unique_ptr<BufferShared<vec4f>> varianceBuffer;
// accumulated world-space normal per pixel
containers::AlignedVector<vec3f> normalBuffer;
std::unique_ptr<BufferShared<vec3f>> normalBuffer;
// accumulated, one RGB per pixel
containers::AlignedVector<vec3f> albedoBuffer;
std::unique_ptr<BufferShared<vec3f>> albedoBuffer;
// primitive ID, object ID, and instance ID
containers::AlignedVector<uint32> primitiveIDBuffer;
containers::AlignedVector<uint32> objectIDBuffer;
containers::AlignedVector<uint32> instanceIDBuffer;
std::unique_ptr<BufferShared<uint32_t>> primitiveIDBuffer;
std::unique_ptr<BufferShared<uint32_t>> objectIDBuffer;
std::unique_ptr<BufferShared<uint32_t>> instanceIDBuffer;
protected:
vec2i getTaskStartPos(const uint32_t taskID) const;
@@ -94,10 +107,11 @@ struct OSPRAY_SDK_INTERFACE LocalFrameBuffer
//// Data ////
vec2i numRenderTasks;
std::vector<uint32_t> renderTaskIDs;
std::unique_ptr<BufferShared<uint32_t>> renderTaskIDs;
// holds accumID per render task, for adaptive accumulation
containers::AlignedVector<int32> taskAccumID;
std::unique_ptr<BufferShared<int32_t>> taskAccumID;
// holds error per tile and adaptive regions
TaskError taskErrorRegion;
};
modules/cpu/fb/LocalFB.ih
+26
View File
@@ -0,0 +1,26 @@
// Copyright 2009-2022 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
#include "fb/FrameBuffer.ih"
#include "fb/Tile.ih"
// c++ shared
#include "LocalFBShared.h"
#ifdef OSPRAY_TARGET_DPCPP
#include "common/rkcommonDPCPPWrappers.h"
#endif
OSPRAY_BEGIN_ISPC_NAMESPACE
SYCL_EXTERNAL void LocalFB_accumulateSample(FrameBuffer *uniform _fb,
const varying ScreenSample &screenSample,
uniform RenderTaskDesc &taskDesc);
SYCL_EXTERNAL uniform RenderTaskDesc LocalFB_getRenderTaskDesc(
FrameBuffer *uniform _fb, const uniform uint32 taskID);
SYCL_EXTERNAL void LocalFB_completeTask(
FrameBuffer *uniform _fb, const uniform RenderTaskDesc &taskDesc);
OSPRAY_END_ISPC_NAMESPACE
modules/cpu/fb/LocalFB.ispc
+21 -7
View File
@@ -6,6 +6,12 @@
// c++ shared
#include "LocalFBShared.h"
#ifdef OSPRAY_TARGET_DPCPP
#include "common/rkcommonDPCPPWrappers.h"
#endif
OSPRAY_BEGIN_ISPC_NAMESPACE
// write tile into the given frame buffer's color buffer,
// this buffer _must_ exist when this fct is called, and it
// _must_ have format 'name'
@@ -30,6 +36,7 @@
for (uint32 iix = tile->region.lower.x + programIndex; \
iix < tile->region.upper.x; \
iix += programCount, chunkID++) { \
type cvtCol; \
unmasked \
{ \
varying vec4f col = make_vec4f(varyTile->r[chunkID], \
@@ -37,7 +44,7 @@
varyTile->b[chunkID], \
varyTile->a[chunkID]); \
\
const type cvtCol = cvt(col); \
cvtCol = cvt(col); \
} \
color[iix] = cvtCol; \
} \
@@ -162,7 +169,8 @@ inline void LocalFB_accumulateAux(uniform vec3f *uniform auxBuffer,
* enabled) and writes aux channels (depth, normal, albedo) that the framebuffer
* has and writes the accumulated color in to the framebuffer for display.
*/
void LocalFB_accumulateSample(FrameBuffer *uniform _fb,
SYCL_EXTERNAL void LocalFB_accumulateSample(
FrameBuffer *uniform _fb,
const varying ScreenSample &screenSample,
uniform RenderTaskDesc &taskDesc)
{
@@ -218,8 +226,11 @@ void LocalFB_accumulateSample(FrameBuffer *uniform _fb,
// Run any pixel ops for this pixel
// TODO: Best as inline in this kernel? Or a separate kernel?
// TODO also needs sycl external
/*
finalColor =
FrameBuffer_runPixelOps(&fb->super, finalColor, depth, normal, albedo);
*/
// Write the final color out to the color buffer in the color format
// requested by the app
@@ -240,10 +251,12 @@ void LocalFB_accumulateSample(FrameBuffer *uniform _fb,
}
}
#ifndef OSPRAY_TARGET_DPCPP
FrameBuffer_updateProgress(&fb->super, popcnt(lanemask()));
#endif
}
uniform RenderTaskDesc LocalFB_getRenderTaskDesc(
SYCL_EXTERNAL uniform RenderTaskDesc LocalFB_getRenderTaskDesc(
FrameBuffer *uniform _fb, const uniform uint32 taskID)
{
LocalFB *uniform fb = (LocalFB * uniform) _fb;
@@ -260,7 +273,7 @@ uniform RenderTaskDesc LocalFB_getRenderTaskDesc(
return desc;
}
void LocalFB_completeTask(
SYCL_EXTERNAL void LocalFB_completeTask(
FrameBuffer *uniform _fb, const uniform RenderTaskDesc &taskDesc)
{
LocalFB *uniform fb = (LocalFB * uniform) _fb;
@@ -377,15 +390,16 @@ export void LocalFrameBuffer_writeAuxTile(void *uniform _fb,
export void *uniform LocalFrameBuffer_accumulateSample_addr()
{
return LocalFB_accumulateSample;
return (void *uniform)LocalFB_accumulateSample;
}
export void *uniform LocalFrameBuffer_getRenderTaskDesc_addr()
{
return LocalFB_getRenderTaskDesc;
return (void *uniform)LocalFB_getRenderTaskDesc;
}
export void *uniform LocalFrameBuffer_completeTask_addr()
{
return LocalFB_completeTask;
return (void *uniform)LocalFB_completeTask;
}
OSPRAY_END_ISPC_NAMESPACE
modules/cpu/fb/LocalFBShared.h
+30 -1
View File
@@ -4,6 +4,7 @@
#pragma once
#include "FrameBufferShared.h"
#include "embree4/rtcore_config.h"
#ifdef __cplusplus
namespace ispc {
@@ -44,8 +45,36 @@ struct LocalFB
primitiveIDBuffer(nullptr),
objectIDBuffer(nullptr),
instanceIDBuffer(nullptr)
{}
{
super.type = FRAMEBUFFER_TYPE_LOCAL;
}
};
#ifdef OSPRAY_TARGET_DPCPP
void LocalFrameBuffer_writeTile_RGBA8(void *_fb, const void *_tile);
void LocalFrameBuffer_writeTile_SRGBA(void *_fb, const void *_tile);
void LocalFrameBuffer_writeTile_RGBA32F(void *_fb, const void *_tile);
void LocalFrameBuffer_writeDepthTile(void *_fb, const void *uniform _tile);
void LocalFrameBuffer_writeAuxTile(void *_fb,
const void *_tile,
vec3f *aux,
const void *_ax,
const void *_ay,
const void *_az);
void LocalFrameBuffer_writeIDTile(void *uniform _fb,
const void *uniform _tile,
uniform uint32 *uniform dst,
const void *uniform src);
SYCL_EXTERNAL void LocalFB_accumulateSample(FrameBuffer *uniform _fb,
const varying ScreenSample &screenSample,
uniform RenderTaskDesc &taskDesc);
SYCL_EXTERNAL uniform RenderTaskDesc LocalFB_getRenderTaskDesc(
FrameBuffer *uniform _fb, const uniform uint32 taskID);
SYCL_EXTERNAL void LocalFB_completeTask(
FrameBuffer *uniform _fb, const uniform RenderTaskDesc &taskDesc);
#endif
} // namespace ispc
#else
};
modules/cpu/fb/RenderTaskDesc.ih
+9
View File
@@ -3,7 +3,15 @@
#pragma once
#ifdef OSPRAY_TARGET_DPCPP
#include "rkcommon/math/box.h"
using namespace rkcommon;
using namespace rkcommon::math;
#else
#include "rkcommon/math/box.ih"
#endif
OSPRAY_BEGIN_ISPC_NAMESPACE
struct RenderTaskDesc
{
@@ -12,3 +20,4 @@ struct RenderTaskDesc
int accumID;
float error;
};
OSPRAY_END_ISPC_NAMESPACE
modules/cpu/fb/SparseFB.cpp
+11 -3
View File
@@ -5,7 +5,9 @@
#include <rkcommon/tasking/parallel_for.h>
#include <cstdlib>
#include "ImageOp.h"
#ifndef OSPRAY_TARGET_DPCPP
#include "fb/SparseFB_ispc.h"
#endif
#include "render/util.h"
#include "rkcommon/common.h"
@@ -228,11 +230,17 @@ void SparseFrameBuffer::setTiles(const std::vector<uint32_t> &_tileIDs)
clear();
#ifndef OSPRAY_TARGET_DPCPP
getSh()->super.accumulateSample =
ispc::SparseFrameBuffer_accumulateSample_addr();
reinterpret_cast<ispc::FrameBuffer_accumulateSampleFct>(
ispc::SparseFrameBuffer_accumulateSample_addr());
getSh()->super.getRenderTaskDesc =
ispc::SparseFrameBuffer_getRenderTaskDesc_addr();
getSh()->super.completeTask = ispc::SparseFrameBuffer_completeTask_addr();
reinterpret_cast<ispc::FrameBuffer_getRenderTaskDescFct>(
ispc::SparseFrameBuffer_getRenderTaskDesc_addr());
getSh()->super.completeTask =
reinterpret_cast<ispc::FrameBuffer_completeTaskFct>(
ispc::SparseFrameBuffer_completeTask_addr());
#endif
getSh()->numRenderTasks = numRenderTasks;
getSh()->totalTiles = totalTiles;
modules/cpu/fb/SparseFB.ih
+13
View File
@@ -6,6 +6,18 @@
#include "SparseFBShared.h"
#include "fb/Tile.ih"
OSPRAY_BEGIN_ISPC_NAMESPACE
SYCL_EXTERNAL uniform RenderTaskDesc SparseFB_getRenderTaskDesc(
FrameBuffer *uniform _fb, const uniform uint32 taskID);
SYCL_EXTERNAL void SparseFB_accumulateSample(FrameBuffer *uniform _fb,
const varying ScreenSample &screenSample,
uniform RenderTaskDesc &taskDesc);
SYCL_EXTERNAL void SparseFB_completeTask(
FrameBuffer *uniform _fb, const uniform RenderTaskDesc &taskDesc);
inline uniform uint32 SparseFB_getTileIndexForTask(
const SparseFB *uniform fb, const uniform uint32 taskID)
{
@@ -16,3 +28,4 @@ inline uniform uint32 SparseFB_getTileIndexForTask(
// tileIdx -> index in the SparseFB's list of tiles
return taskID / nTasksPerTile;
}
OSPRAY_END_ISPC_NAMESPACE
modules/cpu/fb/SparseFB.ispc
+8 -3
View File
@@ -9,6 +9,8 @@
#include "render/Renderer.ih"
#include "render/util.ih"
OSPRAY_BEGIN_ISPC_NAMESPACE
/* Return the position of the screen sample within its tile
*/
inline vec2i SparseFB_getTilePixelPos(SparseFB *uniform fb,
@@ -141,7 +143,9 @@ void SparseFB_accumulateSample(FrameBuffer *uniform _fb,
SparseFB_accumulateTileSample(
fb, screenSample, tileIdx, sampleTilePos, taskDesc);
#ifndef OSPRAY_TARGET_DPCPP
FrameBuffer_updateProgress(&fb->super, popcnt(lanemask()));
#endif
}
uniform RenderTaskDesc SparseFB_getRenderTaskDesc(
@@ -207,15 +211,16 @@ void SparseFB_completeTask(
export void *uniform SparseFrameBuffer_accumulateSample_addr()
{
return SparseFB_accumulateSample;
return (void *uniform)SparseFB_accumulateSample;
}
export void *uniform SparseFrameBuffer_getRenderTaskDesc_addr()
{
return SparseFB_getRenderTaskDesc;
return (void *uniform)SparseFB_getRenderTaskDesc;
}
export void *uniform SparseFrameBuffer_completeTask_addr()
{
return SparseFB_completeTask;
return (void *uniform)SparseFB_completeTask;
}
OSPRAY_END_ISPC_NAMESPACE
modules/cpu/fb/SparseFBShared.h
+3 -1
View File
@@ -41,7 +41,9 @@ struct SparseFB
accumulationBuffer(nullptr),
varianceBuffer(nullptr),
taskRegionError(nullptr)
{}
{
super.type = FRAMEBUFFER_TYPE_SPARSE;
}
};
} // namespace ispc
#else
modules/cpu/fb/TaskError.cpp
+24 -19
View File
@@ -6,17 +6,25 @@
namespace ospray {
TaskError::TaskError(const vec2i &_numTasks)
: numTasks(_numTasks), taskErrorBuffer(numTasks.long_product())
TaskError::TaskError(ispcrt::Device &device, const vec2i &_numTasks)
: numTasks(_numTasks)
{
// maximum number of regions: all regions are of size 3 are split in half
errorRegion.reserve(divRoundUp(taskErrorBuffer.size() * 2, size_t(3)));
clear();
if (numTasks.long_product() > 0) {
taskErrorBuffer =
make_buffer_shared_unique<float>(device, numTasks.long_product());
// maximum number of regions: all regions are of size 3 are split in
// half
errorRegion.reserve(divRoundUp(taskErrorBuffer->size() * 2, size_t(3)));
clear();
}
}
void TaskError::clear()
{
std::fill(taskErrorBuffer.begin(), taskErrorBuffer.end(), inf);
if (!taskErrorBuffer) {
return;
}
std::fill(taskErrorBuffer->begin(), taskErrorBuffer->end(), inf);
errorRegion.clear();
// initially create one region covering the complete tile/image
@@ -25,30 +33,30 @@ void TaskError::clear()
float TaskError::operator[](const int id) const
{
if (taskErrorBuffer.empty()) {
if (!taskErrorBuffer) {
return inf;
}
return taskErrorBuffer[id];
return (*taskErrorBuffer)[id];
}
void TaskError::update(const vec2i &task, const float err)
{
if (!taskErrorBuffer.empty()) {
taskErrorBuffer[task.y * numTasks.x + task.x] = err;
if (taskErrorBuffer) {
(*taskErrorBuffer)[task.y * numTasks.x + task.x] = err;
}
}
float TaskError::refine(const float errorThreshold)
{
if (taskErrorBuffer.empty()) {
if (!taskErrorBuffer) {
return inf;
}
float maxErr = 0.f;
float sumActErr = 0.f;
int activeTasks = 0;
for (const auto &err : taskErrorBuffer) {
for (const auto &err : *taskErrorBuffer) {
maxErr = std::max(maxErr, err);
if (err > errorThreshold) {
sumActErr += err;
@@ -66,8 +74,8 @@ float TaskError::refine(const float errorThreshold)
for (int y = region.lower.y; y < region.upper.y; y++)
for (int x = region.lower.x; x < region.upper.x; x++) {
int idx = y * numTasks.x + x;
err += taskErrorBuffer[idx];
maxErr = std::max(maxErr, taskErrorBuffer[idx]);
err += (*taskErrorBuffer)[idx];
maxErr = std::max(maxErr, (*taskErrorBuffer)[idx]);
}
if (maxErr > errorThreshold) {
// set all tasks of this region to >errorThreshold to enforce their
@@ -76,7 +84,7 @@ float TaskError::refine(const float errorThreshold)
for (int y = region.lower.y; y < region.upper.y; y++)
for (int x = region.lower.x; x < region.upper.x; x++) {
int idx = y * numTasks.x + x;
taskErrorBuffer[idx] = std::max(taskErrorBuffer[idx], minErr);
(*taskErrorBuffer)[idx] = std::max((*taskErrorBuffer)[idx], minErr);
}
}
const vec2i size = region.size();
@@ -109,10 +117,7 @@ float TaskError::refine(const float errorThreshold)
float *TaskError::errorBuffer()
{
if (taskErrorBuffer.empty()) {
return nullptr;
}
return taskErrorBuffer.data();
return taskErrorBuffer ? taskErrorBuffer->data() : nullptr;
}
} // namespace ospray
modules/cpu/fb/TaskError.h
+4 -2
View File
@@ -4,7 +4,9 @@
#pragma once
#include <vector>
#include "common/BufferShared.h"
#include "common/OSPCommon.h"
#include "ispcrt.hpp"
#include "rkcommon/containers/AlignedVector.h"
namespace ospray {
@@ -15,7 +17,7 @@ namespace ospray {
class OSPRAY_SDK_INTERFACE TaskError
{
public:
TaskError(const vec2i &numTasks);
TaskError(ispcrt::Device &ispcrtDevice, const vec2i &numTasks);
// The default constructor will make an empty task error region
TaskError() = default;
@@ -35,7 +37,7 @@ class OSPRAY_SDK_INTERFACE TaskError
protected:
vec2i numTasks = vec2i(0);
// holds error per task
containers::AlignedVector<float> taskErrorBuffer;
std::unique_ptr<BufferShared<float>> taskErrorBuffer;
// image regions (in #tasks) that do not yet estimate the error on
// per-task basis
std::vector<box2i> errorRegion;
modules/cpu/fb/Tile.ih
+11 -1
View File
@@ -3,9 +3,18 @@
#pragma once
#ifdef OSPRAY_TARGET_DPCPP
#include "rkcommon/math/box.h"
using namespace rkcommon;
using namespace rkcommon::math;
#else
#include "rkcommon/math/box.ih"
#endif
#include "TileShared.h"
#include "common/OSPCommon.ih"
#include "rkcommon/math/box.ih"
OSPRAY_BEGIN_ISPC_NAMESPACE
inline void setRGBA(uniform Tile &tile,
const varying uint32 i,
@@ -50,3 +59,4 @@ inline void setNormalAlbedo(
tile.ag[i] = albedo.y;
tile.ab[i] = albedo.z;
}
OSPRAY_END_ISPC_NAMESPACE
modules/cpu/fb/TileShared.h
+5 -1
View File
@@ -23,7 +23,9 @@ static_assert(TILE_SIZE > 0 && (TILE_SIZE & (TILE_SIZE - 1)) == 0,
// Extra defines to hide some differences between ISPC/C++ to unify the
// declaration
#ifdef __cplusplus
#ifndef OSPRAY_TARGET_DPCPP
#define uniform
#endif
#else
#define OSPRAY_SDK_INTERFACE
#endif
@@ -86,7 +88,7 @@ struct OSPRAY_SDK_INTERFACE Tile
#endif
};
#ifndef __cplusplus
#if !defined(__cplusplus) || defined(OSPRAY_TARGET_DPCPP)
struct VaryingTile
{
uniform range2i region; // 4 ints
@@ -117,7 +119,9 @@ struct VaryingTile
#endif
#ifdef __cplusplus
#ifndef OSPRAY_TARGET_DPCPP
#undef uniform
#endif
#else
#undef OSPRAY_SDK_INTERFACE
#endif
modules/cpu/fb/frame_ops/SSAO.h
+26
View File
@@ -7,8 +7,34 @@
#include "rkcommon/tasking/parallel_for.h"
// std
#include <algorithm>
#ifdef OSPRAY_TARGET_DPCPP
#include "common/OSPCommon.ih"
namespace ispc {
void *LiveSSAOFrameOp_create();
void LiveSSAOFrameOp_set(void *uniform _self,
float nearClip,
float ssaoStrength,
void *_windowSize,
void *_pixelSize,
void *_cameraSpace,
void *_kernel,
void *_randomVecs);
void LiveSSAOFrameOp_getOcclusion(const void *_self,
void *_fb,
float *occlusionBuffer,
const float radius,
const float checkRadius,
int kernelSize,
int programID);
void LiveSSAOFrameOp_applyOcclusion(
void *_self, void *_fb, void *_color, float *occlusionBuffer);
int8 getProgramCount();
} // namespace ispc
#else
// ispc
#include "fb/frame_ops/SSAO_ispc.h"
#endif
namespace ospray {
modules/cpu/fb/frame_ops/SSAO.ispc
+13 -2
View File
@@ -1,11 +1,21 @@
// Copyright 2009 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
#ifdef OSPRAY_TARGET_DPCPP
#include "rkcommon/math/LinearSpace.h"
#include "rkcommon/math/vec.h"
using namespace rkcommon;
using namespace rkcommon::math;
#else
#include "rkcommon/math/LinearSpace.ih"
#include "rkcommon/math/vec.ih"
#endif
// ospray
#include "../FrameBufferView.ih"
#include "../Tile.ih"
#include "math/LinearSpace.ih"
#include "rkcommon/math/vec.ih"
OSPRAY_BEGIN_ISPC_NAMESPACE
struct LiveSSAOFrameOp
{
@@ -155,3 +165,4 @@ export int8 uniform getProgramCount()
{
return programCount;
}
OSPRAY_END_ISPC_NAMESPACE
modules/cpu/fb/pixel_ops/ToneMapper.cpp
+5 -1
View File
@@ -3,7 +3,9 @@
#include "ToneMapper.h"
#include "fb/FrameBuffer.h"
#ifndef OSPRAY_TARGET_DPCPP
#include "fb/pixel_ops/ToneMapper_ispc.h"
#endif
using namespace rkcommon;
@@ -69,7 +71,9 @@ LiveToneMapper::LiveToneMapper(FrameBufferView &_fbView,
: AddStructShared(
_fbView.originalFB->getISPCDevice().getIspcrtDevice(), _fbView)
{
getSh()->super.processPixel = ispc::LiveToneMapper_processPixel_addr();
getSh()->super.processPixel =
reinterpret_cast<ispc::LivePixelOp_processPixel>(
ispc::LiveToneMapper_processPixel_addr());
getSh()->exposure = exposure;
getSh()->a = a;
getSh()->b = b;
modules/cpu/fb/pixel_ops/ToneMapper.ispc
+15 -3
View File
@@ -1,9 +1,20 @@
// Copyright 2009 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
#include "fb/pixel_ops/ToneMapperShared.h"
#include "math/LinearSpace.ih"
#ifdef OSPRAY_TARGET_DPCPP
#include "rkcommon/math/LinearSpace.h"
#include "rkcommon/math/vec.h"
using namespace rkcommon;
using namespace rkcommon::math;
#include "common/rkcommonDPCPPWrappers.h"
#else
#include "rkcommon/math/LinearSpace.ih"
#include "rkcommon/math/vec.ih"
#endif
#include "fb/pixel_ops/ToneMapperShared.h"
OSPRAY_BEGIN_ISPC_NAMESPACE
// ACES input transform matrix = RRT_SAT_MAT * XYZ_2_AP1_MAT * D65_2_D60_CAT *
// REC709_2_XYZ_PRI_MAT
@@ -45,5 +56,6 @@ vec4f LiveToneMapper_processPixel(void *uniform _self,
export void *uniform LiveToneMapper_processPixel_addr()
{
return LiveToneMapper_processPixel;
return (void *uniform)LiveToneMapper_processPixel;
}
OSPRAY_END_ISPC_NAMESPACE
modules/cpu/fb/pixel_ops/ToneMapperShared.h
+4
View File
@@ -22,6 +22,10 @@ struct LiveToneMapper
bool acesColor;
};
#ifdef OSPRAY_TARGET_DPCPP
void *LiveToneMapper_processPixel_addr();
#endif
#ifdef __cplusplus
}
#endif
modules/cpu/geometry/Boxes.cpp
+10 -4
View File
@@ -4,15 +4,23 @@
// ospray
#include "Boxes.h"
#include "common/Data.h"
#ifndef OSPRAY_TARGET_DPCPP
// ispc-generated files
#include "geometry/Boxes_ispc.h"
#else
#include "geometry/Boxes.ih"
#endif
namespace ospray {
Boxes::Boxes(api::ISPCDevice &device)
: AddStructShared(device.getIspcrtDevice(), device)
{
getSh()->super.postIntersect = ispc::Boxes_postIntersect_addr();
#ifndef OSPRAY_TARGET_DPCPP
getSh()->super.postIntersect =
reinterpret_cast<ispc::Geometry_postIntersectFct>(
ispc::Boxes_postIntersect_addr());
#endif
}
std::string Boxes::toString() const
@@ -24,9 +32,7 @@ void Boxes::commit()
{
boxData = getParamDataT<box3f>("box", true);
createEmbreeUserGeometry((RTCBoundsFunction)&ispc::Boxes_bounds,
(RTCIntersectFunctionN)&ispc::Boxes_intersect,
(RTCOccludedFunctionN)&ispc::Boxes_occluded);
createEmbreeUserGeometry((RTCBoundsFunction)&ispc::Boxes_bounds);
getSh()->boxes = *ispc(boxData);
getSh()->super.numPrimitives = numPrimitives();
modules/cpu/geometry/Boxes.ih
+21
View File
@@ -0,0 +1,21 @@
// Copyright 2022 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
#include "geometry/Geometry.ih"
OSPRAY_BEGIN_ISPC_NAMESPACE
// Only used on the host
void Boxes_bounds(const RTCBoundsFunctionArguments *uniform args);
SYCL_EXTERNAL void Boxes_postIntersect(const Geometry *uniform _self,
varying DifferentialGeometry &dg,
const varying Ray &ray,
uniform int64 flags);
SYCL_EXTERNAL unmasked void Boxes_intersect_kernel(
const RTCIntersectFunctionNArguments *uniform args,
const uniform bool isOcclusionTest);
OSPRAY_END_ISPC_NAMESPACE
modules/cpu/geometry/Boxes.ispc
+26 -8
View File
@@ -1,17 +1,27 @@
// Copyright 2009 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
#ifdef OSPRAY_TARGET_DPCPP
#include "rkcommon/math/box.h"
#include "rkcommon/math/vec.h"
using namespace rkcommon;
using namespace rkcommon::math;
#else
#include "rkcommon/math/box.ih"
#include "rkcommon/math/vec.ih"
#endif
// ospray
#include "common/Data.ih"
#include "common/DifferentialGeometry.ih"
#include "common/FilterIntersect.ih"
#include "common/Intersect.ih"
#include "common/Ray.ih"
#include "rkcommon/math/box.ih"
#include "rkcommon/math/vec.ih"
// c++ shared
#include "BoxesShared.h"
OSPRAY_BEGIN_ISPC_NAMESPACE
export void Boxes_bounds(const RTCBoundsFunctionArguments *uniform args)
{
Boxes *uniform self = (Boxes * uniform) args->geometryUserPtr;
@@ -22,12 +32,15 @@ export void Boxes_bounds(const RTCBoundsFunctionArguments *uniform args)
*out = make_box3fa(get_box3f(self->boxes, primID));
}
void Boxes_intersect_kernel(const RTCIntersectFunctionNArguments *uniform args,
SYCL_EXTERNAL unmasked void Boxes_intersect_kernel(
const RTCIntersectFunctionNArguments *uniform args,
const uniform bool isOcclusionTest)
{
// make sure to set the mask
if (!args->valid[programIndex])
if (!args->valid[programIndex]) {
return;
}
args->valid[programIndex] = 0;
Boxes *uniform self = (Boxes * uniform) args->geometryUserPtr;
varying Ray *uniform ray = (varying Ray * uniform) args->rayhit;
@@ -35,10 +48,13 @@ void Boxes_intersect_kernel(const RTCIntersectFunctionNArguments *uniform args,
uniform int primID = args->primID;
uniform box3f box = get_box3f(self->boxes, primID);
const Intersections isect = intersectBox(ray->org, ray->dir, box);
// call intersection filtering callback and setup hit if accepted
filterIntersectionBoth(args, isect, isOcclusionTest);
if (filterIntersectionBoth(args, isect, isOcclusionTest)) {
args->valid[programIndex] = -1;
}
}
export void Boxes_intersect(
@@ -53,7 +69,7 @@ export void Boxes_occluded(
Boxes_intersect_kernel((RTCIntersectFunctionNArguments * uniform) args, true);
}
void Boxes_postIntersect(const Geometry *uniform geometry,
SYCL_EXTERNAL void Boxes_postIntersect(const Geometry *uniform geometry,
varying DifferentialGeometry &dg,
const varying Ray &ray,
uniform int64 flags)
@@ -63,5 +79,7 @@ void Boxes_postIntersect(const Geometry *uniform geometry,
export void *uniform Boxes_postIntersect_addr()
{
return Boxes_postIntersect;
}
return (void *uniform)Boxes_postIntersect;
}
OSPRAY_END_ISPC_NAMESPACE
modules/cpu/geometry/BoxesShared.h
+7 -1
View File
@@ -13,8 +13,14 @@ struct Boxes
{
Geometry super;
Data1D boxes;
};
#ifdef __cplusplus
Boxes()
{
super.type = GEOMETRY_TYPE_BOXES;
}
#endif
};
#ifdef __cplusplus
} // namespace ispc
#endif // __cplusplus
modules/cpu/geometry/Curves.cpp
+5 -1
View File
@@ -5,8 +5,10 @@
#include "Curves.h"
#include "common/DGEnum.h"
#ifndef OSPRAY_TARGET_DPCPP
// ispc-generated files
#include "geometry/Curves_ispc.h"
#endif
// std
#include <map>
@@ -48,7 +50,9 @@ static std::map<std::pair<OSPCurveType, OSPCurveBasis>, RTCGeometryType>
Curves::Curves(api::ISPCDevice &device)
: AddStructShared(device.getIspcrtDevice(), device)
{
getSh()->super.postIntersect = ispc::Curves_postIntersect_addr();
getSh()->super.postIntersect =
reinterpret_cast<ispc::Geometry_postIntersectFct>(
ispc::Curves_postIntersect_addr());
// TODO implement area sampling of OldCurves for geometry lights
}
modules/cpu/geometry/Curves.ispc
+10 -2
View File
@@ -6,6 +6,12 @@
// c++ shared
#include "CurvesShared.h"
#ifdef OSPRAY_TARGET_DPCPP
using namespace ospray;
#endif
OSPRAY_BEGIN_ISPC_NAMESPACE
static void Curves_postIntersect(const Geometry *uniform _self,
varying DifferentialGeometry &dg,
const varying Ray &ray,
@@ -42,5 +48,7 @@ static void Curves_postIntersect(const Geometry *uniform _self,
export void *uniform Curves_postIntersect_addr()
{
return Curves_postIntersect;
}
return (void *uniform)Curves_postIntersect;
}
OSPRAY_END_ISPC_NAMESPACE
modules/cpu/geometry/CurvesShared.h
+3
View File
@@ -18,6 +18,9 @@ struct Curves
#ifdef __cplusplus
Curves() : geom(nullptr), flagMask(-1) {}
};
#ifdef OSPRAY_TARGET_DPCPP
void *Curves_postIntersect_addr();
#endif
} // namespace ispc
#else
};
modules/cpu/geometry/GeometricModel.ih
+20 -2
View File
@@ -3,13 +3,29 @@
#pragma once
#ifdef OSPRAY_TARGET_DPCPP
#include "common/DGEnum.h"
#include "common/rkcommonDPCPPWrappers.h"
#include "rkcommon/math/AffineSpace.h"
using namespace rkcommon;
using namespace rkcommon::math;
using namespace ospray;
#else
#include "rkcommon/math/AffineSpace.ih"
#endif
#include "Geometry.ih"
#include "common/Data.ih"
#include "math/AffineSpace.ih"
#include "render/Renderer.ih"
// c++ shared
#include "GeometricModelShared.h"
#ifdef OSPRAY_TARGET_DPCPP
#include "geometry/MeshShared.h"
#endif
OSPRAY_BEGIN_ISPC_NAMESPACE
inline Material *GeometricModel_getMaterial(const GeometricModel *uniform self,
const Renderer *uniform renderer,
const int32 primID)
@@ -44,7 +60,7 @@ inline void GeometricModel_postIntersect(const GeometricModel *uniform self,
{
Geometry *uniform geom = self->geom;
geom->postIntersect(geom, dg, ray, flags);
Geometry_dispatch_postIntersect(geom, dg, ray, flags);
dg.areaPDF = self->areaPDF;
dg.objID =
@@ -65,3 +81,5 @@ inline void GeometricModel_postIntersect(const GeometricModel *uniform self,
if (flags & DG_NS && self->invertedNormals)
dg.Ns = neg(dg.Ns);
}
OSPRAY_END_ISPC_NAMESPACE
modules/cpu/geometry/GeometricModel.ispc
+16 -2
View File
@@ -1,6 +1,10 @@
// Copyright 2009 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
#ifdef OSPRAY_TARGET_DPCPP
#include "rkcommon/math/math.h"
#endif
// ospray
#include "GeometricModel.ih"
#include "math/Distribution1D.ih"
@@ -8,6 +12,8 @@
// c++ shared
#include "common/InstanceShared.h"
OSPRAY_BEGIN_ISPC_NAMESPACE
export uniform int GeometricModel_gatherEmissivePrimIDs(void *uniform _model,
void *uniform _renderer,
void *uniform _instance,
@@ -22,7 +28,12 @@ export uniform int GeometricModel_gatherEmissivePrimIDs(void *uniform _model,
// create the list of emissive primitive IDs
uniform int32 numEmissivePrims = 0;
#ifdef OSPRAY_TARGET_DPCPP
// DPC++ execution in a thread is scalar
for (int primID = 0; primID < geo->numPrimitives; ++primID) {
#else
foreach (primID = 0 ... geo->numPrimitives) {
#endif
Material *mat =
(Material *)GeometricModel_getMaterial(model, renderer, primID);
const vec3f emission = mat->emission;
@@ -36,11 +47,14 @@ export uniform int GeometricModel_gatherEmissivePrimIDs(void *uniform _model,
// create the sampling distribution
// TODO: use emissive power instead of just area
// TODO: motion blur can introduce scale, which means areas cannot accurately
// be precomputed
// TODO: motion blur can introduce scale, which means areas cannot
// accurately be precomputed
// TODO: Can stay as a ptr call if not moved to GPU
geo->getAreas(geo, primIDs, numEmissivePrims, instance->xfm, distribution);
pdf = model->areaPDF =
1.f / Distribution1D_create(numEmissivePrims, distribution);
return numEmissivePrims;
}
OSPRAY_END_ISPC_NAMESPACE

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