Draw: Cleanup the GLSL intersection code

Only store center and size. Skip resource finalize computation.

.arcconfig

@@ -0,0 +1,8 @@
+{
+	"project_id" : "Blender",
+	"conduit_uri" : "https://developer.blender.org/",
+	"phabricator.uri" : "https://developer.blender.org/",
+	"git.default-relative-commit" : "origin/master",
+	"arc.land.update.default" : "rebase",
+	"arc.land.onto.default" : "master"
+}

.clang-format

@@ -236,8 +236,6 @@ ForEachMacros:
   - LOOP_UNSELECTED_POINTS
   - LOOP_VISIBLE_KEYS
   - LOOP_VISIBLE_POINTS
-  - LIGHT_FOREACH_BEGIN_DIRECTIONAL
-  - LIGHT_FOREACH_BEGIN_LOCAL
   - LISTBASE_CIRCULAR_BACKWARD_BEGIN
   - LISTBASE_CIRCULAR_FORWARD_BEGIN
   - LISTBASE_FOREACH

.gitea/default_merge_message/REBASE_TEMPLATE.md

@@ -1,5 +0,0 @@
-${CommitTitle}
-
-${CommitBody}
-
-Pull Request #${PullRequestIndex}

.gitea/default_merge_message/SQUASH_TEMPLATE.md

@@ -1,3 +0,0 @@
-${PullRequestTitle}
-
-Pull Request #${PullRequestIndex}

.gitea/issue_template.yaml

@@ -0,0 +1,45 @@
+name: Bug Report
+about: File a bug report
+labels:
+  - bug
+ref: master
+body:
+  - type: markdown
+    attributes:
+      value: |
+                ### First time bug reporting?
+                Read [these tips](https://wiki.blender.org/wiki/Process/Bug_Reports) and watch this **[How to Report a Bug](https://www.youtube.com/watch?v=JTD0OJq_rF4)** video to make a complete, valid bug report. Remember to write your bug report in **English**.
+
+                ### What not to report here
+                For feature requests, feedback, questions or issues building Blender, see [communication channels](https://wiki.blender.org/wiki/Communication/Contact#User_Feedback_and_Requests).
+
+                ### Please verify
+                * Always test with the latest official release from [blender.org](https://www.blender.org/) and daily build from [builder.blender.org](https://builder.blender.org/).
+                * Please use `Help > Report a Bug` in Blender to automatically fill system information and exact Blender version.
+                * Test [previous Blender versions](https://download.blender.org/release/) to find the latest version that was working as expected.
+                * Find steps to redo the bug consistently without any non-official add-ons, and include a **small and simple .blend file** to demonstrate the bug.
+                * If there are multiple bugs, make multiple bug reports.
+                * Sometimes, driver or software upgrades cause problems. On Windows, try a clean install of the graphics drivers.
+
+                ### Help the developers
+                Bug fixing is important, the developers will handle a report swiftly. For that reason, we need your help to carefully provide instructions that others can follow quickly. You do your half of the work, then we do our half!
+
+                If a report is tagged with Needs Information from User and it has no reply after a week, we will assume the issue is gone and close the report.
+
+  - type: textarea
+    attributes:
+      label: "Description"
+      value: |
+               **System Information**
+               Operating system:
+               Graphics card:
+
+               **Blender Version**
+               Broken: (example: 2.80, edbf15d3c044, master, 2018-11-28, as found on the splash screen)
+               Worked: (newest version of Blender that worked as expected)
+
+               **Short description of error**
+
+               **Exact steps for others to reproduce the error**
+               Based on the default startup or an attached .blend file (as simple as possible).
+

.gitea/issue_template/bug.yaml

@@ -1,44 +0,0 @@
-name: Bug Report
-about: File a bug report
-labels:
-  - "type::Report"
-  - "status::Needs Triage"
-  - "priority::Normal"
-body:
-  - type: markdown
-    attributes:
-      value: |
-                ### Instructions
-                First time reporting? See [tips](https://wiki.blender.org/wiki/Process/Bug_Reports).
-
-                * Use **Help > Report a Bug** in Blender to fill system information and exact Blender version.
-                * Test [daily builds](https://builder.blender.org/) to verify if the issue is already fixed.
-                * Test [previous versions](https://download.blender.org/release/) to find an older working version.
-                * For feature requests, feedback, questions or build issues, see [communication channels](https://wiki.blender.org/wiki/Communication/Contact#User_Feedback_and_Requests).
-                * If there are multiple bugs, make multiple bug reports.
-
-  - type: textarea
-    id: body
-    attributes:
-      label: "Description"
-      hide_label: true
-      value: |
-               **System Information**
-               Operating system:
-               Graphics card:
-
-               **Blender Version**
-               Broken: (example: 2.80, edbf15d3c044, master, 2018-11-28, as found on the splash screen)
-               Worked: (newest version of Blender that worked as expected)
-
-               **Short description of error**
-
-               **Exact steps for others to reproduce the error**
-               Based on the default startup or an attached .blend file (as simple as possible).
-
-  - type: markdown
-    attributes:
-      value: |
-                ### Help the developers
-
-                Bug fixing is important, the developers will handle reports swiftly. For that reason, carefully provide exact steps and a **small and simple .blend file** to reproduce the problem. You do your half of the work, then we do our half!

.gitea/issue_template/config.yaml

@@ -1 +0,0 @@
-blank_issues_enabled: false

.gitea/issue_template/design.yaml

@@ -1,10 +0,0 @@
-name: Design
-about: Create a design task (for developers only)
-labels:
-  - "type::Design"
-body:
-  - type: textarea
-    id: body
-    attributes:
-      label: "Description"
-      hide_label: true

.gitea/issue_template/todo.yaml

@@ -1,10 +0,0 @@
-name: To Do
-about: Create a to do task (for developers only)
-labels:
-  - "type::To Do"
-body:
-  - type: textarea
-    id: body
-    attributes:
-      label: "Description"
-      hide_label: true

.gitea/pull_request_template.md

@@ -0,0 +1,4 @@
+---
+name: Pull Request
+about: Submit a pull request
+---

.gitea/pull_request_template.yaml

@@ -1,17 +0,0 @@
-name: Pull Request
-about: Contribute code to Blender
-body:
-  - type: markdown
-    attributes:
-      value: |
-        ### Instructions
-
-        Guides to [contributing code](https://wiki.blender.org/index.php/Dev:Doc/Process/Contributing_Code) and effective [code review](https://wiki.blender.org/index.php/Dev:Doc/Tools/Code_Review).
-
-        By submitting code here, you agree that the code is (compatible with) GNU GPL v2 or later.
-
-  - type: textarea
-    id: body
-    attributes:
-      label: "Description"
-      hide_label: true

.github/pull_request_template.md

@@ -1,4 +1,5 @@
-This repository is only used as a mirror. Blender development happens on projects.blender.org.
+This repository is only used as a mirror of git.blender.org. Blender development happens on
+https://developer.blender.org.
 
 To get started with contributing code, please see:
 https://wiki.blender.org/wiki/Process/Contributing_Code

.github/stale.yml

@@ -15,7 +15,8 @@ staleLabel: stale
 # Comment to post when closing a stale Issue or Pull Request.
 closeComment: >
   This issue has been automatically closed, because this repository is only
-  used as a mirror. Blender development happens on projects.blender.org.
+  used as a mirror of git.blender.org. Blender development happens on
+  developer.blender.org.
 
   To get started contributing code, please read:
   https://wiki.blender.org/wiki/Process/Contributing_Code

.gitmodules

@@ -1,20 +1,20 @@
 [submodule "release/scripts/addons"]
 	path = release/scripts/addons
 	url = ../blender-addons.git
-	branch = main
+	branch = master
 	ignore = all
 [submodule "release/scripts/addons_contrib"]
 	path = release/scripts/addons_contrib
 	url = ../blender-addons-contrib.git
-	branch = main
+	branch = master
 	ignore = all
 [submodule "release/datafiles/locale"]
 	path = release/datafiles/locale
 	url = ../blender-translations.git
-	branch = main
+	branch = master
 	ignore = all
 [submodule "source/tools"]
 	path = source/tools
 	url = ../blender-dev-tools.git
-	branch = main
+	branch = master
 	ignore = all

CMakeLists.txt

@@ -167,26 +167,14 @@ get_blender_version()
 option(WITH_BLENDER "Build blender (disable to build only the blender player)" ON)
 mark_as_advanced(WITH_BLENDER)
 
-if(WIN32)
-  option(WITH_BLENDER_THUMBNAILER "\
-Build \"BlendThumb.dll\" helper for Windows explorer integration to support extracting \
-thumbnails from `.blend` files."
-    ON
-  )
+if(APPLE)
+  # In future, can be used with `quicklookthumbnailing/qlthumbnailreply` to create file
+  # thumbnails for say Finder. Turn it off for now.
+  option(WITH_BLENDER_THUMBNAILER "Build \"blender-thumbnailer\" thumbnail extraction utility" OFF)
+elseif(WIN32)
+  option(WITH_BLENDER_THUMBNAILER "Build \"BlendThumb.dll\" helper for Windows explorer integration" ON)
 else()
-  set(_option_default ON)
-  if(APPLE)
-    # In future, can be used with `quicklookthumbnailing/qlthumbnailreply`
-    # to create file thumbnails for say Finder.
-    # Turn it off for now, even though it can build on APPLE, it's not likely to be useful.
-    set(_option_default OFF)
-  endif()
-  option(WITH_BLENDER_THUMBNAILER "\
-Build stand-alone \"blender-thumbnailer\" command-line thumbnail extraction utility, \
-intended for use by file-managers to extract PNG images from `.blend` files."
-    ${_option_default}
-  )
-  unset(_option_default)
+  option(WITH_BLENDER_THUMBNAILER "Build \"blender-thumbnailer\" thumbnail extraction utility" ON)
 endif()
 
 option(WITH_INTERNATIONAL "Enable I18N (International fonts and text)" ON)
@@ -226,19 +214,14 @@ option(WITH_BULLET        "Enable Bullet (Physics Engine)" ON)
 option(WITH_SYSTEM_BULLET "Use the systems bullet library (currently unsupported due to missing features in upstream!)" )
 mark_as_advanced(WITH_SYSTEM_BULLET)
 option(WITH_OPENCOLORIO   "Enable OpenColorIO color management" ON)
-
-set(_option_default ON)
 if(APPLE)
   # There's no OpenXR runtime in sight for macOS, neither is code well
   # tested there -> disable it by default.
-  set(_option_default OFF)
-endif()
-option(WITH_XR_OPENXR "Enable VR features through the OpenXR specification" ${_option_default})
-if(APPLE)
+  option(WITH_XR_OPENXR   "Enable VR features through the OpenXR specification" OFF)
   mark_as_advanced(WITH_XR_OPENXR)
+else()
+  option(WITH_XR_OPENXR   "Enable VR features through the OpenXR specification" ON)
 endif()
-unset(_option_default)
-
 option(WITH_GMP "Enable features depending on GMP (Exact Boolean)" ON)
 
 # Compositor
@@ -370,13 +353,12 @@ else()
   set(WITH_COREAUDIO OFF)
 endif()
 if(NOT WIN32)
-  set(_option_default ON)
   if(APPLE)
-    set(_option_default OFF)
+    option(WITH_JACK          "Enable JACK Support (http://www.jackaudio.org)" OFF)
+  else()
+    option(WITH_JACK          "Enable JACK Support (http://www.jackaudio.org)" ON)
   endif()
-  option(WITH_JACK "Enable JACK Support (http://www.jackaudio.org)" ${_option_default})
-  unset(_option_default)
-  option(WITH_JACK_DYNLOAD "Enable runtime dynamic JACK libraries loading" OFF)
+  option(WITH_JACK_DYNLOAD  "Enable runtime dynamic JACK libraries loading" OFF)
 else()
   set(WITH_JACK OFF)
 endif()
@@ -417,26 +399,6 @@ mark_as_advanced(WITH_SYSTEM_GLOG)
 # Freestyle
 option(WITH_FREESTYLE     "Enable Freestyle (advanced edges rendering)" ON)
 
-# Libraries.
-if(UNIX AND NOT APPLE)
-  # Optionally build without pre-compiled libraries.
-  # NOTE: this could be supported on all platforms however in practice UNIX is the only platform
-  # that has good support for detecting installed libraries.
-  option(WITH_LIBS_PRECOMPILED "\
-Detect and link against pre-compiled libraries (typically found under \"../lib/\"). \
-Disabling this option will use the system libraries although cached paths \
-that point to pre-compiled libraries will be left as-is."
-    ON
-  )
-  mark_as_advanced(WITH_LIBS_PRECOMPILED)
-
-  option(WITH_STATIC_LIBS "Try to link with static libraries, as much as possible, to make blender more portable across distributions" OFF)
-  if(WITH_STATIC_LIBS)
-    option(WITH_BOOST_ICU "Boost uses ICU library (required for linking with static Boost built with libicu)." OFF)
-    mark_as_advanced(WITH_BOOST_ICU)
-  endif()
-endif()
-
 # Misc
 if(WIN32 OR APPLE)
   option(WITH_INPUT_IME "Enable Input Method Editor (IME) for complex Asian character input" ON)
@@ -444,6 +406,11 @@ endif()
 option(WITH_INPUT_NDOF "Enable NDOF input devices (SpaceNavigator and friends)" ON)
 if(UNIX AND NOT APPLE)
   option(WITH_INSTALL_PORTABLE "Install redistributable runtime, otherwise install into CMAKE_INSTALL_PREFIX" ON)
+  option(WITH_STATIC_LIBS "Try to link with static libraries, as much as possible, to make blender more portable across distributions" OFF)
+  if(WITH_STATIC_LIBS)
+    option(WITH_BOOST_ICU "Boost uses ICU library (required for linking with static Boost built with libicu)." OFF)
+    mark_as_advanced(WITH_BOOST_ICU)
+  endif()
 endif()
 
 option(WITH_PYTHON_INSTALL       "Copy system python into the blender install folder" ON)
@@ -625,10 +592,8 @@ mark_as_advanced(
 
 # Vulkan
 option(WITH_VULKAN_BACKEND "Enable Vulkan as graphics backend (only for development)" OFF)
-option(WITH_VULKAN_GUARDEDALLOC "Use guardedalloc for host allocations done inside Vulkan (development option)" OFF)
 mark_as_advanced(
   WITH_VULKAN_BACKEND
-  WITH_VULKAN_GUARDEDALLOC
 )
 
 # Metal
@@ -1028,8 +993,6 @@ set(PLATFORM_LINKLIBS "")
 # - CMAKE_EXE_LINKER_FLAGS_DEBUG
 set(PLATFORM_LINKFLAGS "")
 set(PLATFORM_LINKFLAGS_DEBUG "")
-set(PLATFORM_LINKFLAGS_RELEASE "")
-set(PLATFORM_LINKFLAGS_EXECUTABLE "")
 
 if(NOT CMAKE_BUILD_TYPE MATCHES "Release")
   if(WITH_COMPILER_ASAN)
@@ -1243,6 +1206,13 @@ if(WITH_OPENGL)
   add_definitions(-DWITH_OPENGL)
 endif()
 
+#-----------------------------------------------------------------------------
+# Configure Vulkan.
+
+if(WITH_VULKAN_BACKEND)
+  list(APPEND BLENDER_GL_LIBRARIES ${VULKAN_LIBRARIES})
+endif()
+
 # -----------------------------------------------------------------------------
 # Configure Metal
 
@@ -1292,14 +1262,12 @@ endif()
 # -----------------------------------------------------------------------------
 # Configure Bullet
 
-if(WITH_BULLET)
-  if(WITH_SYSTEM_BULLET)
-    find_package(Bullet)
-    set_and_warn_library_found("Bullet" BULLET_FOUND WITH_BULLET)
-  else()
-    set(BULLET_INCLUDE_DIRS "${CMAKE_SOURCE_DIR}/extern/bullet2/src")
-    set(BULLET_LIBRARIES "extern_bullet")
-  endif()
+if(WITH_BULLET AND WITH_SYSTEM_BULLET)
+  find_package(Bullet)
+  set_and_warn_library_found("Bullet" BULLET_FOUND WITH_BULLET)
+else()
+  set(BULLET_INCLUDE_DIRS "${CMAKE_SOURCE_DIR}/extern/bullet2/src")
+  # set(BULLET_LIBRARIES "")
 endif()
 
 

GNUmakefile

@@ -71,13 +71,6 @@ Static Source Code Checking
    * check_mypy:            Checks all Python scripts using mypy,
                             see: source/tools/check_source/check_mypy_config.py scripts which are included.
 
-Documentation Checking
-
-   * check_wiki_file_structure:
-     Check the WIKI documentation for the source-tree's file structure
-     matches Blender's source-code.
-     See: https://wiki.blender.org/wiki/Source/File_Structure
-
 Spell Checkers
    This runs the spell checker from the developer tools repositor.
 
@@ -299,11 +292,7 @@ else
 	ifneq ("$(wildcard $(DEPS_BUILD_DIR)/build.ninja)","")
 		DEPS_BUILD_COMMAND:=ninja
 	else
-		ifeq ($(OS), Darwin)
-			DEPS_BUILD_COMMAND:=make -s
-		else
-			DEPS_BUILD_COMMAND:="$(BLENDER_DIR)/build_files/build_environment/linux/make_deps_wrapper.sh" -s
-		endif
+		DEPS_BUILD_COMMAND:=make -s
 	endif
 endif
 
@@ -402,7 +391,7 @@ endif
 
 deps: .FORCE
 	@echo
-	@echo Configuring dependencies in \"$(DEPS_BUILD_DIR)\", install to \"$(DEPS_INSTALL_DIR)\"
+	@echo Configuring dependencies in \"$(DEPS_BUILD_DIR)\"
 
 	@cmake -H"$(DEPS_SOURCE_DIR)" \
 	       -B"$(DEPS_BUILD_DIR)" \
@@ -492,10 +481,6 @@ check_smatch: .FORCE
 check_mypy: .FORCE
 	@$(PYTHON) "$(BLENDER_DIR)/source/tools/check_source/check_mypy.py"
 
-check_wiki_file_structure: .FORCE
-	@PYTHONIOENCODING=utf_8 $(PYTHON) \
-	    "$(BLENDER_DIR)/source/tools/check_wiki/check_wiki_file_structure.py"
-
 check_spelling_py: .FORCE
 	@cd "$(BUILD_DIR)" ; \
 	PYTHONIOENCODING=utf_8 $(PYTHON) \

README.md

@@ -24,7 +24,7 @@ Development
 -----------
 
 - [Build Instructions](https://wiki.blender.org/wiki/Building_Blender)
-- [Code Review & Bug Tracker](https://projects.blender.org)
+- [Code Review & Bug Tracker](https://developer.blender.org)
 - [Developer Forum](https://devtalk.blender.org)
 - [Developer Documentation](https://wiki.blender.org)
 

build_files/build_environment/cmake/dpcpp.cmake

@@ -2,7 +2,7 @@
 
 # LLVM does not switch over to cpp17 until llvm 16 and building ealier versions with
 # MSVC is leading to some crashes in ISPC. Switch back to their default on all platforms
-# for now.
+# for now. 
 string(REPLACE "-DCMAKE_CXX_STANDARD=17" " " DPCPP_CMAKE_FLAGS "${DEFAULT_CMAKE_FLAGS}")
 
 if(WIN32)

build_files/build_environment/cmake/epoxy.cmake

@@ -10,7 +10,7 @@ ExternalProject_Add(external_epoxy
   URL_HASH ${EPOXY_HASH_TYPE}=${EPOXY_HASH}
   PREFIX ${BUILD_DIR}/epoxy
   PATCH_COMMAND ${PATCH_CMD} -p 1 -N -d ${BUILD_DIR}/epoxy/src/external_epoxy/ < ${PATCH_DIR}/epoxy.diff
-  CONFIGURE_COMMAND ${CONFIGURE_ENV} && ${MESON} setup --prefix ${LIBDIR}/epoxy --default-library ${EPOXY_LIB_TYPE} --libdir lib ${BUILD_DIR}/epoxy/src/external_epoxy-build ${BUILD_DIR}/epoxy/src/external_epoxy -Dtests=false ${MESON_BUILD_TYPE}
+  CONFIGURE_COMMAND ${CONFIGURE_ENV} && ${MESON} setup --prefix ${LIBDIR}/epoxy --default-library ${EPOXY_LIB_TYPE} --libdir lib ${BUILD_DIR}/epoxy/src/external_epoxy-build ${BUILD_DIR}/epoxy/src/external_epoxy -Dtests=false
   BUILD_COMMAND ninja
   INSTALL_COMMAND ninja install
 )

build_files/build_environment/cmake/fribidi.cmake

@@ -9,7 +9,7 @@ ExternalProject_Add(external_fribidi
   URL_HASH ${FRIBIDI_HASH_TYPE}=${FRIBIDI_HASH}
   DOWNLOAD_DIR ${DOWNLOAD_DIR}
   PREFIX ${BUILD_DIR}/fribidi
-  CONFIGURE_COMMAND ${MESON} setup --prefix ${LIBDIR}/fribidi ${MESON_BUILD_TYPE} -Ddocs=false --default-library static --libdir lib ${BUILD_DIR}/fribidi/src/external_fribidi-build ${BUILD_DIR}/fribidi/src/external_fribidi
+  CONFIGURE_COMMAND ${MESON} setup --prefix ${LIBDIR}/fribidi -Ddocs=false --default-library static --libdir lib ${BUILD_DIR}/fribidi/src/external_fribidi-build ${BUILD_DIR}/fribidi/src/external_fribidi
   BUILD_COMMAND ninja
   INSTALL_COMMAND ninja install
   INSTALL_DIR ${LIBDIR}/fribidi

build_files/build_environment/cmake/gmp.cmake

@@ -22,7 +22,7 @@ elseif(UNIX AND NOT APPLE)
   )
 endif()
 
-# Boolean crashes with Arm assembly, see #103423.
+# Boolean crashes with Arm assembly, see T103423.
 if(BLENDER_PLATFORM_ARM)
   set(GMP_OPTIONS
     ${GMP_OPTIONS}

build_files/build_environment/cmake/harfbuzz.cmake

@@ -21,7 +21,6 @@ set(HARFBUZZ_EXTRA_OPTIONS
   # Only used for command line utilities,
   # disable as this would add an addition & unnecessary build-dependency.
   -Dcairo=disabled
-  ${MESON_BUILD_TYPE}
 )
 
 ExternalProject_Add(external_harfbuzz
@@ -60,10 +59,3 @@ if(BUILD_MODE STREQUAL Release AND WIN32)
     DEPENDEES install
   )
 endif()
-
-if(BUILD_MODE STREQUAL Debug AND WIN32)
-  ExternalProject_Add_Step(external_harfbuzz after_install
-    COMMAND ${CMAKE_COMMAND} -E copy ${LIBDIR}/harfbuzz/lib/libharfbuzz.a ${HARVEST_TARGET}/harfbuzz/lib/libharfbuzz_d.lib
-    DEPENDEES install
-  )
-endif()

build_files/build_environment/cmake/igc.cmake

@@ -40,8 +40,7 @@ ExternalProject_Add(external_igc_llvm
     ${PATCH_CMD} -p 1 -d ${IGC_LLVM_SOURCE_DIR} < ${IGC_OPENCL_CLANG_PATCH_DIR}/clang/0004-OpenCL-support-cl_ext_float_atomics.patch &&
     ${PATCH_CMD} -p 1 -d ${IGC_LLVM_SOURCE_DIR} < ${IGC_OPENCL_CLANG_PATCH_DIR}/clang/0005-OpenCL-Add-cl_khr_integer_dot_product.patch &&
     ${PATCH_CMD} -p 1 -d ${IGC_LLVM_SOURCE_DIR} < ${IGC_OPENCL_CLANG_PATCH_DIR}/llvm/0001-Memory-leak-fix-for-Managed-Static-Mutex.patch &&
-    ${PATCH_CMD} -p 1 -d ${IGC_LLVM_SOURCE_DIR} < ${IGC_OPENCL_CLANG_PATCH_DIR}/llvm/0002-Remove-repo-name-in-LLVM-IR.patch &&
-    ${PATCH_CMD} -p 1 -d ${IGC_LLVM_SOURCE_DIR} < ${IGC_OPENCL_CLANG_PATCH_DIR}/llvm/0003-Add-missing-include-limit-in-benchmark.patch
+    ${PATCH_CMD} -p 1 -d ${IGC_LLVM_SOURCE_DIR} < ${IGC_OPENCL_CLANG_PATCH_DIR}/llvm/0002-Remove-repo-name-in-LLVM-IR.patch
 )
 add_dependencies(
   external_igc_llvm
@@ -56,6 +55,9 @@ ExternalProject_Add(external_igc_spirv_translator
   CONFIGURE_COMMAND echo .
   BUILD_COMMAND echo .
   INSTALL_COMMAND echo .
+  PATCH_COMMAND ${PATCH_CMD} -p 1 -d ${IGC_SPIRV_TRANSLATOR_SOURCE_DIR} < ${IGC_OPENCL_CLANG_PATCH_DIR}/spirv/0001-update-SPIR-V-headers-for-SPV_INTEL_split_barrier.patch &&
+    ${PATCH_CMD} -p 1 -d ${IGC_SPIRV_TRANSLATOR_SOURCE_DIR} < ${IGC_OPENCL_CLANG_PATCH_DIR}/spirv/0002-Add-support-for-split-barriers-extension-SPV_INTEL_s.patch &&
+    ${PATCH_CMD} -p 1 -d ${IGC_SPIRV_TRANSLATOR_SOURCE_DIR} < ${IGC_OPENCL_CLANG_PATCH_DIR}/spirv/0003-Support-cl_bf16_conversions.patch
 )
 add_dependencies(
   external_igc_spirv_translator

build_files/build_environment/cmake/llvm.cmake

@@ -42,7 +42,7 @@ endif()
 
 # LLVM does not switch over to cpp17 until llvm 16 and building ealier versions with
 # MSVC is leading to some crashes in ISPC. Switch back to their default on all platforms
-# for now.
+# for now. 
 string(REPLACE "-DCMAKE_CXX_STANDARD=17" " " LLVM_CMAKE_FLAGS "${DEFAULT_CMAKE_FLAGS}")
 
 # short project name due to long filename issues on windows

build_files/build_environment/cmake/mesa.cmake

@@ -15,7 +15,7 @@ llvm-config = '${LIBDIR}/llvm/bin/llvm-config'"
 )
 
 set(MESA_EXTRA_FLAGS
-  ${MESON_BUILD_TYPE}
+  -Dbuildtype=release
   -Dc_args=${MESA_CFLAGS}
   -Dcpp_args=${MESA_CXXFLAGS}
   -Dc_link_args=${MESA_LDFLAGS}

build_files/build_environment/cmake/options.cmake

@@ -16,10 +16,8 @@ message("BuildMode = ${BUILD_MODE}")
 
 if(BUILD_MODE STREQUAL "Debug")
   set(LIBDIR ${CMAKE_CURRENT_BINARY_DIR}/Debug)
-  set(MESON_BUILD_TYPE -Dbuildtype=debug)
 else()
   set(LIBDIR ${CMAKE_CURRENT_BINARY_DIR}/Release)
-  set(MESON_BUILD_TYPE -Dbuildtype=release)
 endif()
 
 set(DOWNLOAD_DIR "${CMAKE_CURRENT_BINARY_DIR}/downloads" CACHE STRING "Path for downloaded files")

build_files/build_environment/cmake/python.cmake

@@ -88,19 +88,6 @@ else()
     export LDFLAGS=${PYTHON_LDFLAGS} &&
     export PKG_CONFIG_PATH=${LIBDIR}/ffi/lib/pkgconfig)
 
-  # NOTE: untested on APPLE so far.
-  if(NOT APPLE)
-    set(PYTHON_CONFIGURE_EXTRA_ARGS
-      ${PYTHON_CONFIGURE_EXTRA_ARGS}
-      # Used on most release Linux builds (Fedora for e.g.),
-      # increases build times noticeably with the benefit of a modest speedup at runtime.
-      --enable-optimizations
-      # While LTO is OK when building on the same system, it's incompatible across GCC versions,
-      # making it impractical for developers to build against, so keep it disabled.
-      # `--with-lto`
-    )
-  endif()
-
   ExternalProject_Add(external_python
     URL file://${PACKAGE_DIR}/${PYTHON_FILE}
     DOWNLOAD_DIR ${DOWNLOAD_DIR}

build_files/build_environment/cmake/ssl.cmake

@@ -10,9 +10,9 @@ if(WIN32)
     DOWNLOAD_DIR ${DOWNLOAD_DIR}
     URL_HASH ${SSL_HASH_TYPE}=${SSL_HASH}
     PREFIX ${BUILD_DIR}/ssl
-    CONFIGURE_COMMAND echo "."
-    BUILD_COMMAND echo "."
-    INSTALL_COMMAND echo "."
+    CONFIGURE_COMMAND echo "." 
+    BUILD_COMMAND echo "." 
+    INSTALL_COMMAND echo "." 
     INSTALL_DIR ${LIBDIR}/ssl
   )
 else()
@@ -46,4 +46,4 @@ else()
     INSTALL_COMMAND ${CONFIGURE_ENV} && cd ${BUILD_DIR}/ssl/src/external_ssl/ && make install
     INSTALL_DIR ${LIBDIR}/ssl
   )
-endif()
+endif()

build_files/build_environment/cmake/usd.cmake

@@ -29,7 +29,7 @@ elseif(UNIX)
   set(USD_PLATFORM_FLAGS
     -DPYTHON_INCLUDE_DIR=${LIBDIR}/python/include/python${PYTHON_SHORT_VERSION}/
     -DPYTHON_LIBRARY=${LIBDIR}/tbb/lib/${LIBPREFIX}${TBB_LIBRARY}${SHAREDLIBEXT}
-  )
+   )
 
   if(APPLE)
     set(USD_SHARED_LINKER_FLAGS "-Xlinker -undefined -Xlinker dynamic_lookup")

build_files/build_environment/cmake/versions.cmake

@@ -668,9 +668,9 @@ set(SPIRV_HEADERS_FILE SPIR-V-Headers-${SPIRV_HEADERS_VERSION}.tar.gz)
 # compiler, the versions used are taken from the following location
 # https://github.com/intel/intel-graphics-compiler/releases
 
-set(IGC_VERSION 1.0.13064.7)
+set(IGC_VERSION 1.0.12149.1)
 set(IGC_URI https://github.com/intel/intel-graphics-compiler/archive/refs/tags/igc-${IGC_VERSION}.tar.gz)
-set(IGC_HASH a929abd4cca2b293961ec0437ee4b3b2147bd3b2c8a3c423af78c0c359b2e5ae)
+set(IGC_HASH 44f67f24e3bc5130f9f062533abf8154782a9d0a992bc19b498639a8521ae836)
 set(IGC_HASH_TYPE SHA256)
 set(IGC_FILE igc-${IGC_VERSION}.tar.gz)
 
@@ -690,15 +690,15 @@ set(IGC_LLVM_FILE ${IGC_LLVM_VERSION}.tar.gz)
 #
 # WARNING WARNING WARNING
 
-set(IGC_OPENCL_CLANG_VERSION ee31812ea8b89d08c2918f045d11a19bd33525c5)
+set(IGC_OPENCL_CLANG_VERSION 363a5262d8c7cff3fb28f3bdb5d85c8d7e91c1bb)
 set(IGC_OPENCL_CLANG_URI https://github.com/intel/opencl-clang/archive/${IGC_OPENCL_CLANG_VERSION}.tar.gz)
-set(IGC_OPENCL_CLANG_HASH 1db6735bbcfaa31e8a9ba39f121d6bafa806ea8919e9f56782d6aaa67771ddda)
+set(IGC_OPENCL_CLANG_HASH aa8cf72bb239722ce8ce44f79413c6887ecc8ca18477dd520aa5c4809756da9a)
 set(IGC_OPENCL_CLANG_HASH_TYPE SHA256)
 set(IGC_OPENCL_CLANG_FILE opencl-clang-${IGC_OPENCL_CLANG_VERSION}.tar.gz)
 
-set(IGC_VCINTRINSICS_VERSION v0.11.0)
+set(IGC_VCINTRINSICS_VERSION v0.5.0)
 set(IGC_VCINTRINSICS_URI https://github.com/intel/vc-intrinsics/archive/refs/tags/${IGC_VCINTRINSICS_VERSION}.tar.gz)
-set(IGC_VCINTRINSICS_HASH e5acd5626ce7fa6d41ce154c50ac805eda734ee66af94ef28e680ac2ad81bb9f)
+set(IGC_VCINTRINSICS_HASH 70bb47c5e32173cf61514941e83ae7c7eb4485e6d2fca60cfa1f50d4f42c41f2)
 set(IGC_VCINTRINSICS_HASH_TYPE SHA256)
 set(IGC_VCINTRINSICS_FILE vc-intrinsics-${IGC_VCINTRINSICS_VERSION}.tar.gz)
 
@@ -714,9 +714,9 @@ set(IGC_SPIRV_TOOLS_HASH 6e19900e948944243024aedd0a201baf3854b377b9cc7a386553bc1
 set(IGC_SPIRV_TOOLS_HASH_TYPE SHA256)
 set(IGC_SPIRV_TOOLS_FILE SPIR-V-Tools-${IGC_SPIRV_TOOLS_VERSION}.tar.gz)
 
-set(IGC_SPIRV_TRANSLATOR_VERSION d739c01d65ec00dee64dedd40deed805216a7193)
+set(IGC_SPIRV_TRANSLATOR_VERSION a31ffaeef77e23d500b3ea3d35e0c42ff5648ad9)
 set(IGC_SPIRV_TRANSLATOR_URI https://github.com/KhronosGroup/SPIRV-LLVM-Translator/archive/${IGC_SPIRV_TRANSLATOR_VERSION}.tar.gz)
-set(IGC_SPIRV_TRANSLATOR_HASH ddc0cc9ccbe59dadeaf291012d59de142b2e9f2b124dbb634644d39daddaa13e)
+set(IGC_SPIRV_TRANSLATOR_HASH 9e26c96a45341b8f8af521bacea20e752623346340addd02af95d669f6e89252)
 set(IGC_SPIRV_TRANSLATOR_HASH_TYPE SHA256)
 set(IGC_SPIRV_TRANSLATOR_FILE SPIR-V-Translator-${IGC_SPIRV_TRANSLATOR_VERSION}.tar.gz)
 
@@ -724,15 +724,15 @@ set(IGC_SPIRV_TRANSLATOR_FILE SPIR-V-Translator-${IGC_SPIRV_TRANSLATOR_VERSION}.
 ### Intel Graphics Compiler DEPS END ###
 ########################################
 
-set(GMMLIB_VERSION intel-gmmlib-22.3.0)
+set(GMMLIB_VERSION intel-gmmlib-22.1.8)
 set(GMMLIB_URI https://github.com/intel/gmmlib/archive/refs/tags/${GMMLIB_VERSION}.tar.gz)
-set(GMMLIB_HASH c1f33e1519edfc527127baeb0436b783430dfd256c643130169a3a71dc86aff9)
+set(GMMLIB_HASH bf23e9a3742b4fb98c7666c9e9b29f3219e4b2fb4d831aaf4eed71f5e2d17368)
 set(GMMLIB_HASH_TYPE SHA256)
 set(GMMLIB_FILE ${GMMLIB_VERSION}.tar.gz)
 
-set(OCLOC_VERSION 22.49.25018.21)
+set(OCLOC_VERSION 22.38.24278)
 set(OCLOC_URI https://github.com/intel/compute-runtime/archive/refs/tags/${OCLOC_VERSION}.tar.gz)
-set(OCLOC_HASH 92362dae08b503a34e5d3820ed284198c452bcd5e7504d90eb69887b20492c06)
+set(OCLOC_HASH db0c542fccd651e6404b15a74d46027f1ce0eda8dc9e25a40cbb6c0faef257ee)
 set(OCLOC_HASH_TYPE SHA256)
 set(OCLOC_FILE ocloc-${OCLOC_VERSION}.tar.gz)
 

build_files/build_environment/cmake/wayland.cmake

@@ -13,7 +13,7 @@ ExternalProject_Add(external_wayland
   # NOTE: `-lm` is needed for `libxml2` which is a static library that uses `libm.so`,
   # without this, math symbols such as `floor` aren't found.
   CONFIGURE_COMMAND ${CMAKE_COMMAND} -E env PKG_CONFIG_PATH=${LIBDIR}/expat/lib/pkgconfig:${LIBDIR}/xml2/lib/pkgconfig:${LIBDIR}/ffi/lib/pkgconfig:$PKG_CONFIG_PATH
-                    ${MESON} --prefix ${LIBDIR}/wayland ${MESON_BUILD_TYPE} -Ddocumentation=false -Dtests=false -D "c_link_args=-L${LIBDIR}/ffi/lib -lm" . ../external_wayland
+                    ${MESON} --prefix ${LIBDIR}/wayland -Ddocumentation=false -Dtests=false -D "c_link_args=-L${LIBDIR}/ffi/lib -lm" . ../external_wayland
   BUILD_COMMAND ninja
   INSTALL_COMMAND ninja install
 )

build_files/build_environment/cmake/wayland_protocols.cmake

@@ -7,7 +7,7 @@ ExternalProject_Add(external_wayland_protocols
   PREFIX ${BUILD_DIR}/wayland-protocols
   # Use `-E` so the `PKG_CONFIG_PATH` can be defined to link against our own WAYLAND.
   CONFIGURE_COMMAND ${CMAKE_COMMAND} -E env PKG_CONFIG_PATH=${LIBDIR}/wayland/lib64/pkgconfig:$PKG_CONFIG_PATH
-                    ${MESON} --prefix ${LIBDIR}/wayland-protocols ${MESON_BUILD_TYPE} . ../external_wayland_protocols -Dtests=false
+                    ${MESON} --prefix ${LIBDIR}/wayland-protocols . ../external_wayland_protocols -Dtests=false
   BUILD_COMMAND ninja
   INSTALL_COMMAND ninja install
 )

build_files/build_environment/cmake/xml2.cmake

@@ -1,7 +1,7 @@
 # SPDX-License-Identifier: GPL-2.0-or-later
 
 if(WIN32)
-  set(XML2_EXTRA_ARGS
+  set(XML2_EXTRA_ARGS 
     -DLIBXML2_WITH_ZLIB=OFF
     -DLIBXML2_WITH_LZMA=OFF
     -DLIBXML2_WITH_PYTHON=OFF

build_files/build_environment/linux/make_deps_wrapper.sh

@@ -1,74 +0,0 @@
-#!/usr/bin/env bash
-# SPDX-License-Identifier: GPL-2.0-or-later
-
-# This script ensures:
-# - One dependency is built at a time.
-# - That dependency uses all available cores.
-#
-# Without this, simply calling `make -j$(nproc)` from the `${CMAKE_BUILD_DIR}/deps/`
-# directory will build many projects at once.
-#
-# This is undesirable for the following reasons:
-#
-# - The output from projects is mixed together,
-#   making it difficult to track down the cause of a build failure.
-#
-# - Larger dependencies such as LLVM can bottleneck the build process,
-#   making it necessary to cancel the build and manually run build commands in each directory.
-#
-# - Building many projects at once means canceling (Control-C) can lead to the build being in an undefined state.
-#   It's possible canceling happens as a patch is being applied or files are being copied.
-#   (steps that aren't part of the compilation process where it's typically safe to cancel).
-
-if [[ -z "$MY_MAKE_CALL_LEVEL" ]]; then
-  export MY_MAKE_CALL_LEVEL=0
-  export MY_MAKEFLAGS=$MAKEFLAGS
-
-  # Extract the jobs argument (`-jN`, `-j N`, `--jobs=N`).
-  add_next=0
-  for i in "$@"; do
-    case $i in
-      -j*)
-        export MY_JOBS_ARG=$i
-        if [ "$MY_JOBS_ARG" = "-j" ]; then
-          add_next=1
-        fi
-        ;;
-      --jobs=*)
-        shift # past argument=value
-        MY_JOBS_ARG=$i
-        ;;
-      *)
-        if (( add_next == 1 )); then
-          MY_JOBS_ARG="$MY_JOBS_ARG $i"
-          add_next=0
-        fi
-        ;;
-    esac
-  done
-  unset i add_next
-
-  if [[ -z "$MY_JOBS_ARG" ]]; then
-    MY_JOBS_ARG="-j$(nproc)"
-  fi
-  export MY_JOBS_ARG
-  # Support user defined `MAKEFLAGS`.
-  export MAKEFLAGS="$MY_MAKEFLAGS -j1"
-else
-  export MY_MAKE_CALL_LEVEL=$(( MY_MAKE_CALL_LEVEL + 1 ))
-  if (( MY_MAKE_CALL_LEVEL == 1 )); then
-    # Important to set jobs to 1, otherwise user defined jobs argument is used.
-    export MAKEFLAGS="$MY_MAKEFLAGS -j1"
-  elif (( MY_MAKE_CALL_LEVEL == 2 )); then
-    # This is the level used by each sub-project.
-    export MAKEFLAGS="$MY_MAKEFLAGS $MY_JOBS_ARG"
-  fi
-  # Else leave `MY_MAKEFLAGS` flags as-is, avoids setting a high number of jobs on recursive
-  # calls (which may easily run out of memory). Let the job-server handle the rest.
-fi
-
-# Useful for troubleshooting the wrapper.
-# echo "Call level: $MY_MAKE_CALL_LEVEL, args=$@".
-
-# Call actual make but ensure recursive calls run via this script.
-exec make MAKE="$0" "$@"

build_files/build_environment/patches/igc_opencl_clang.diff

@@ -1,7 +1,7 @@
 diff -Naur external_igc_opencl_clang.orig/CMakeLists.txt external_igc_opencl_clang/CMakeLists.txt
 --- external_igc_opencl_clang.orig/CMakeLists.txt	2022-03-16 05:51:10 -0600
 +++ external_igc_opencl_clang/CMakeLists.txt	2022-05-23 10:40:09 -0600
-@@ -147,22 +147,24 @@
+@@ -126,22 +126,24 @@
          )
      endif()
  

build_files/cmake/Modules/FindMoltenVK.cmake

@@ -19,13 +19,9 @@ ENDIF()
 
 SET(_moltenvk_SEARCH_DIRS
   ${MOLTENVK_ROOT_DIR}
+  ${LIBDIR}/vulkan/MoltenVK
 )
 
-# FIXME: These finder modules typically don't use LIBDIR,
-# this should be set by `./build_files/cmake/platform/` instead.
-IF(DEFINED LIBDIR)
-  SET(_moltenvk_SEARCH_DIRS ${_moltenvk_SEARCH_DIRS} ${LIBDIR}/moltenvk)
-ENDIF()
 
 FIND_PATH(MOLTENVK_INCLUDE_DIR
   NAMES

build_files/cmake/Modules/FindOptiX.cmake

@@ -17,13 +17,9 @@ ENDIF()
 
 SET(_optix_SEARCH_DIRS
   ${OPTIX_ROOT_DIR}
+  "$ENV{PROGRAMDATA}/NVIDIA Corporation/OptiX SDK 7.3.0"
 )
 
-# TODO: Which environment uses this?
-if(DEFINED ENV{PROGRAMDATA})
-  list(APPEND _optix_SEARCH_DIRS "$ENV{PROGRAMDATA}/NVIDIA Corporation/OptiX SDK 7.3.0")
-endif()
-
 FIND_PATH(OPTIX_INCLUDE_DIR
   NAMES
     optix.h

build_files/cmake/Modules/FindPythonLibsUnix.cmake

@@ -67,8 +67,6 @@ ENDIF()
 
 STRING(REPLACE "." "" PYTHON_VERSION_NO_DOTS ${PYTHON_VERSION})
 
-SET(_PYTHON_ABI_FLAGS "")
-
 SET(_python_SEARCH_DIRS
   ${PYTHON_ROOT_DIR}
   "$ENV{HOME}/py${PYTHON_VERSION_NO_DOTS}"

build_files/cmake/Modules/FindShaderC.cmake

@@ -1,63 +0,0 @@
-# SPDX-License-Identifier: BSD-3-Clause
-# Copyright 2023 Blender Foundation.
-
-# - Find ShaderC libraries
-# Find the ShaderC includes and libraries
-# This module defines
-#  SHADERC_INCLUDE_DIRS, where to find MoltenVK headers, Set when
-#                        SHADERC_INCLUDE_DIR is found.
-#  SHADERC_LIBRARIES, libraries to link against to use ShaderC.
-#  SHADERC_ROOT_DIR, The base directory to search for ShaderC.
-#                    This can also be an environment variable.
-#  SHADERC_FOUND, If false, do not try to use ShaderC.
-#
-
-# If SHADERC_ROOT_DIR was defined in the environment, use it.
-IF(NOT SHADERC_ROOT_DIR AND NOT $ENV{SHADERC_ROOT_DIR} STREQUAL "")
-  SET(SHADERC_ROOT_DIR $ENV{SHADERC_ROOT_DIR})
-ENDIF()
-
-SET(_shaderc_SEARCH_DIRS
-  ${SHADERC_ROOT_DIR}
-)
-
-# FIXME: These finder modules typically don't use LIBDIR,
-# this should be set by `./build_files/cmake/platform/` instead.
-IF(DEFINED LIBDIR)
-  SET(_shaderc_SEARCH_DIRS ${_shaderc_SEARCH_DIRS} ${LIBDIR}/shaderc)
-ENDIF()
-
-FIND_PATH(SHADERC_INCLUDE_DIR
-  NAMES
-    shaderc/shaderc.h
-  HINTS
-    ${_shaderc_SEARCH_DIRS}
-  PATH_SUFFIXES
-    include
-)
-
-FIND_LIBRARY(SHADERC_LIBRARY
-  NAMES
-    shaderc_combined
-  HINTS
-    ${_shaderc_SEARCH_DIRS}
-  PATH_SUFFIXES
-    lib
-)
-
-# handle the QUIETLY and REQUIRED arguments and set SHADERC_FOUND to TRUE if
-# all listed variables are TRUE
-INCLUDE(FindPackageHandleStandardArgs)
-FIND_PACKAGE_HANDLE_STANDARD_ARGS(ShaderC DEFAULT_MSG SHADERC_LIBRARY SHADERC_INCLUDE_DIR)
-
-IF(SHADERC_FOUND)
-  SET(SHADERC_LIBRARIES ${SHADERC_LIBRARY})
-  SET(SHADERC_INCLUDE_DIRS ${SHADERC_INCLUDE_DIR})
-ENDIF()
-
-MARK_AS_ADVANCED(
-  SHADERC_INCLUDE_DIR
-  SHADERC_LIBRARY
-)
-
-UNSET(_shaderc_SEARCH_DIRS)

build_files/cmake/Modules/FindVulkan.cmake

@@ -1,63 +0,0 @@
-# SPDX-License-Identifier: BSD-3-Clause
-# Copyright 2023 Blender Foundation.
-
-# - Find Vulkan libraries
-# Find the Vulkan includes and libraries
-# This module defines
-#  VULKAN_INCLUDE_DIRS, where to find Vulkan headers, Set when
-#                       VULKAN_INCLUDE_DIR is found.
-#  VULKAN_LIBRARIES, libraries to link against to use Vulkan.
-#  VULKAN_ROOT_DIR, The base directory to search for Vulkan.
-#                    This can also be an environment variable.
-#  VULKAN_FOUND, If false, do not try to use Vulkan.
-#
-
-# If VULKAN_ROOT_DIR was defined in the environment, use it.
-IF(NOT VULKAN_ROOT_DIR AND NOT $ENV{VULKAN_ROOT_DIR} STREQUAL "")
-  SET(VULKAN_ROOT_DIR $ENV{VULKAN_ROOT_DIR})
-ENDIF()
-
-SET(_vulkan_SEARCH_DIRS
-  ${VULKAN_ROOT_DIR}
-)
-
-# FIXME: These finder modules typically don't use LIBDIR,
-# this should be set by `./build_files/cmake/platform/` instead.
-IF(DEFINED LIBDIR)
-  SET(_vulkan_SEARCH_DIRS ${_vulkan_SEARCH_DIRS} ${LIBDIR}/vulkan)
-ENDIF()
-
-FIND_PATH(VULKAN_INCLUDE_DIR
-  NAMES
-    vulkan/vulkan.h
-  HINTS
-    ${_vulkan_SEARCH_DIRS}
-  PATH_SUFFIXES
-    include
-)
-
-FIND_LIBRARY(VULKAN_LIBRARY
-  NAMES
-    vulkan
-  HINTS
-    ${_vulkan_SEARCH_DIRS}
-  PATH_SUFFIXES
-    lib
-)
-
-# handle the QUIETLY and REQUIRED arguments and set VULKAN_FOUND to TRUE if
-# all listed variables are TRUE
-INCLUDE(FindPackageHandleStandardArgs)
-FIND_PACKAGE_HANDLE_STANDARD_ARGS(Vulkan DEFAULT_MSG VULKAN_LIBRARY VULKAN_INCLUDE_DIR)
-
-IF(VULKAN_FOUND)
-  SET(VULKAN_LIBRARIES ${VULKAN_LIBRARY})
-  SET(VULKAN_INCLUDE_DIRS ${VULKAN_INCLUDE_DIR})
-ENDIF()
-
-MARK_AS_ADVANCED(
-  VULKAN_INCLUDE_DIR
-  VULKAN_LIBRARY
-)
-
-UNSET(_vulkan_SEARCH_DIRS)

build_files/cmake/buildinfo.cmake

@@ -23,19 +23,19 @@ if(EXISTS ${SOURCE_DIR}/.git)
 
   if(MY_WC_BRANCH STREQUAL "HEAD")
     # Detached HEAD, check whether commit hash is reachable
-    # in the main branch
+    # in the master branch
     execute_process(COMMAND git rev-parse --short=12 HEAD
                     WORKING_DIRECTORY ${SOURCE_DIR}
                     OUTPUT_VARIABLE MY_WC_HASH
                     OUTPUT_STRIP_TRAILING_WHITESPACE)
 
-    execute_process(COMMAND git branch --list main blender-v* --contains ${MY_WC_HASH}
+    execute_process(COMMAND git branch --list master blender-v* --contains ${MY_WC_HASH}
                     WORKING_DIRECTORY ${SOURCE_DIR}
                     OUTPUT_VARIABLE _git_contains_check
                     OUTPUT_STRIP_TRAILING_WHITESPACE)
 
     if(NOT _git_contains_check STREQUAL "")
-      set(MY_WC_BRANCH "main")
+      set(MY_WC_BRANCH "master")
     else()
       execute_process(COMMAND git show-ref --tags -d
                       WORKING_DIRECTORY ${SOURCE_DIR}
@@ -48,7 +48,7 @@ if(EXISTS ${SOURCE_DIR}/.git)
                       OUTPUT_STRIP_TRAILING_WHITESPACE)
 
       if(_git_tag_hashes MATCHES "${_git_head_hash}")
-        set(MY_WC_BRANCH "main")
+        set(MY_WC_BRANCH "master")
       else()
         execute_process(COMMAND git branch --contains ${MY_WC_HASH}
                         WORKING_DIRECTORY ${SOURCE_DIR}

build_files/cmake/cmake_print_build_options.py

@@ -6,80 +6,18 @@
 import re
 import sys
 
-from typing import Optional
-
 cmakelists_file = sys.argv[-1]
 
 
-def count_backslashes_before_pos(file_data: str, pos: int) -> int:
-    slash_count = 0
-    pos -= 1
-    while pos >= 0:
-        if file_data[pos] != '\\':
-            break
-        pos -= 1
-        slash_count += 1
-    return slash_count
-
-
-def extract_cmake_string_at_pos(file_data: str, pos_beg: int) -> Optional[str]:
-    assert file_data[pos_beg - 1] == '"'
-
-    pos = pos_beg
-    # Dummy assignment.
-    pos_end = pos_beg
-    while True:
-        pos_next = file_data.find('"', pos)
-        if pos_next == -1:
-            raise Exception("Un-terminated string (parse error?)")
-
-        count_slashes = count_backslashes_before_pos(file_data, pos_next)
-        if (count_slashes % 2) == 0:
-            pos_end = pos_next
-            # Found the closing quote.
-            break
-
-        # The quote was back-slash escaped, step over it.
-        pos = pos_next + 1
-        file_data[pos_next]
-
-    assert file_data[pos_end] == '"'
-
-    if pos_beg == pos_end:
-        return None
-
-    # See: https://cmake.org/cmake/help/latest/manual/cmake-language.7.html#escape-sequences
-    text = file_data[pos_beg: pos_end].replace(
-        # Handle back-slash literals.
-        "\\\\", "\\",
-    ).replace(
-        # Handle tabs.
-        "\\t", "\t",
-    ).replace(
-        # Handle escaped quotes.
-        "\\\"", "\"",
-    ).replace(
-        # Handle tabs.
-        "\\;", ";",
-    ).replace(
-        # Handle trailing newlines.
-        "\\\n", "",
-    )
-
-    return text
-
-
-def main() -> None:
+def main():
     options = []
-    with open(cmakelists_file, 'r', encoding="utf-8") as fh:
-        file_data = fh.read()
-        for m in re.finditer(r"^\s*option\s*\(\s*(WITH_[a-zA-Z0-9_]+)\s+(\")", file_data, re.MULTILINE):
-            option_name = m.group(1)
-            option_descr = extract_cmake_string_at_pos(file_data, m.span(2)[1])
-            if option_descr is None:
-                # Possibly a parsing error, at least show something.
-                option_descr = "(UNDOCUMENTED)"
-            options.append("{:s}: {:s}".format(option_name, option_descr))
+    for l in open(cmakelists_file, 'r').readlines():
+        if not l.lstrip().startswith('#'):
+            l_option = re.sub(r'.*\boption\s*\(\s*(WITH_[a-zA-Z0-9_]+)\s+\"(.*)\"\s*.*', r'\g<1> - \g<2>', l)
+            if l_option != l:
+                l_option = l_option.strip()
+                if l_option.startswith('WITH_'):
+                    options.append(l_option)
 
     print('\n'.join(options))
 

build_files/cmake/config/blender_release.cmake

@@ -85,7 +85,7 @@ if(NOT APPLE)
   set(WITH_CYCLES_DEVICE_OPTIX    ON  CACHE BOOL "" FORCE)
   set(WITH_CYCLES_CUDA_BINARIES   ON  CACHE BOOL "" FORCE)
   set(WITH_CYCLES_CUBIN_COMPILER  OFF CACHE BOOL "" FORCE)
-  set(WITH_CYCLES_HIP_BINARIES    OFF CACHE BOOL "" FORCE)
+  set(WITH_CYCLES_HIP_BINARIES    ON  CACHE BOOL "" FORCE)
   set(WITH_CYCLES_DEVICE_ONEAPI   ON  CACHE BOOL "" FORCE)
   set(WITH_CYCLES_ONEAPI_BINARIES ON  CACHE BOOL "" FORCE)
 endif()

build_files/cmake/example_scripts/cmake_linux_install.sh

@@ -11,11 +11,11 @@
 mkdir ~/blender-git
 cd ~/blender-git
 
-git clone https://projects.blender.org/blender/blender.git
+git clone http://git.blender.org/blender.git
 cd blender
 git submodule update --init --recursive
-git submodule foreach git checkout main
-git submodule foreach git pull --rebase origin main
+git submodule foreach git checkout master
+git submodule foreach git pull --rebase origin master
 
 # create build dir
 mkdir ~/blender-git/build-cmake
@@ -35,7 +35,7 @@ ln -s ~/blender-git/build-cmake/bin/blender ~/blender-git/blender/blender.bin
 echo ""
 echo "* Useful Commands *"
 echo "   Run Blender: ~/blender-git/blender/blender.bin"
-echo "   Update Blender: git pull --rebase; git submodule foreach git pull --rebase origin main"
+echo "   Update Blender: git pull --rebase; git submodule foreach git pull --rebase origin master"
 echo "   Reconfigure Blender: cd ~/blender-git/build-cmake ; cmake ."
 echo "   Build Blender: cd ~/blender-git/build-cmake ; make"
 echo ""

build_files/cmake/macros.cmake

@@ -544,15 +544,13 @@ endfunction()
 function(setup_platform_linker_libs
   target
   )
-  # jemalloc must be early in the list, to be before pthread (see #57998).
+  # jemalloc must be early in the list, to be before pthread (see T57998)
   if(WITH_MEM_JEMALLOC)
     target_link_libraries(${target} ${JEMALLOC_LIBRARIES})
   endif()
 
   if(WIN32 AND NOT UNIX)
-    if(DEFINED PTHREADS_LIBRARIES)
-      target_link_libraries(${target} ${PTHREADS_LIBRARIES})
-    endif()
+    target_link_libraries(${target} ${PTHREADS_LIBRARIES})
   endif()
 
   # target_link_libraries(${target} ${PLATFORM_LINKLIBS} ${CMAKE_DL_LIBS})
@@ -1117,7 +1115,7 @@ function(find_python_package
     # endif()
     # Not set, so initialize.
   else()
-    string(REPLACE "." ";" _PY_VER_SPLIT "${PYTHON_VERSION}")
+   string(REPLACE "." ";" _PY_VER_SPLIT "${PYTHON_VERSION}")
     list(GET _PY_VER_SPLIT 0 _PY_VER_MAJOR)
 
     # re-cache
@@ -1264,7 +1262,7 @@ endmacro()
 
 # Utility to gather and install precompiled shared libraries.
 macro(add_bundled_libraries library_dir)
-  if(DEFINED LIBDIR)
+  if(EXISTS ${LIBDIR})
     set(_library_dir ${LIBDIR}/${library_dir})
     if(WIN32)
       file(GLOB _all_library_versions ${_library_dir}/*\.dll)
@@ -1277,7 +1275,7 @@ macro(add_bundled_libraries library_dir)
     list(APPEND PLATFORM_BUNDLED_LIBRARY_DIRS ${_library_dir})
     unset(_all_library_versions)
     unset(_library_dir)
-  endif()
+ endif()
 endmacro()
 
 macro(windows_install_shared_manifest)

build_files/cmake/platform/platform_apple.cmake

@@ -97,8 +97,20 @@ add_bundled_libraries(materialx/lib)
 
 if(WITH_VULKAN_BACKEND)
   find_package(MoltenVK REQUIRED)
-  find_package(ShaderC REQUIRED)
-  find_package(Vulkan REQUIRED)
+
+  if(EXISTS ${LIBDIR}/vulkan)
+    set(VULKAN_FOUND On)
+    set(VULKAN_ROOT_DIR ${LIBDIR}/vulkan/macOS)
+    set(VULKAN_INCLUDE_DIR ${VULKAN_ROOT_DIR}/include)
+    set(VULKAN_LIBRARY ${VULKAN_ROOT_DIR}/lib/libvulkan.1.dylib)
+    set(SHADERC_LIBRARY ${VULKAN_ROOT_DIR}/lib/libshaderc_combined.a)
+
+    set(VULKAN_INCLUDE_DIRS ${VULKAN_INCLUDE_DIR} ${MOLTENVK_INCLUDE_DIRS})
+    set(VULKAN_LIBRARIES ${VULKAN_LIBRARY} ${SHADERC_LIBRARY} ${MOLTENVK_LIBRARIES})
+  else()
+    message(WARNING "Vulkan SDK was not found, disabling WITH_VULKAN_BACKEND")
+    set(WITH_VULKAN_BACKEND OFF)
+  endif()
 endif()
 
 if(WITH_OPENSUBDIV)
@@ -440,7 +452,7 @@ string(APPEND PLATFORM_LINKFLAGS " -stdlib=libc++")
 # Make stack size more similar to Embree, required for Embree.
 string(APPEND PLATFORM_LINKFLAGS_EXECUTABLE " -Wl,-stack_size,0x100000")
 
-# Suppress ranlib "has no symbols" warnings (workaround for #48250).
+# Suppress ranlib "has no symbols" warnings (workaround for T48250)
 set(CMAKE_C_ARCHIVE_CREATE   "<CMAKE_AR> Scr <TARGET> <LINK_FLAGS> <OBJECTS>")
 set(CMAKE_CXX_ARCHIVE_CREATE "<CMAKE_AR> Scr <TARGET> <LINK_FLAGS> <OBJECTS>")
 # llvm-ranlib doesn't support this flag. Xcode's libtool does.

build_files/cmake/platform/platform_old_libs_update.cmake

@@ -1,12 +1,7 @@
 # SPDX-License-Identifier: GPL-2.0-or-later
 # Copyright 2022 Blender Foundation. All rights reserved.
 
-# Auto update existing CMake caches for new libraries.
-
-# Assert that `LIBDIR` is defined.
-if(NOT (DEFINED LIBDIR))
-  message(FATAL_ERROR "Logical error, expected 'LIBDIR' to be defined!")
-endif()
+# Auto update existing CMake caches for new libraries
 
 # Clear cached variables whose name matches `pattern`.
 function(unset_cache_variables pattern)

build_files/cmake/platform/platform_unix.cmake

@@ -4,52 +4,38 @@
 # Libraries configuration for any *nix system including Linux and Unix (excluding APPLE).
 
 # Detect precompiled library directory
+if(NOT DEFINED LIBDIR)
+  # Path to a locally compiled libraries.
+  set(LIBDIR_NAME ${CMAKE_SYSTEM_NAME}_${CMAKE_SYSTEM_PROCESSOR})
+  string(TOLOWER ${LIBDIR_NAME} LIBDIR_NAME)
+  set(LIBDIR_NATIVE_ABI ${CMAKE_SOURCE_DIR}/../lib/${LIBDIR_NAME})
 
-if(NOT WITH_LIBS_PRECOMPILED)
-  unset(LIBDIR)
-else()
-  if(NOT DEFINED LIBDIR)
-    # Path to a locally compiled libraries.
-    set(LIBDIR_NAME ${CMAKE_SYSTEM_NAME}_${CMAKE_SYSTEM_PROCESSOR})
-    string(TOLOWER ${LIBDIR_NAME} LIBDIR_NAME)
-    set(LIBDIR_NATIVE_ABI ${CMAKE_SOURCE_DIR}/../lib/${LIBDIR_NAME})
+  # Path to precompiled libraries with known glibc 2.28 ABI.
+  set(LIBDIR_GLIBC228_ABI ${CMAKE_SOURCE_DIR}/../lib/linux_x86_64_glibc_228)
 
-    # Path to precompiled libraries with known glibc 2.28 ABI.
-    set(LIBDIR_GLIBC228_ABI ${CMAKE_SOURCE_DIR}/../lib/linux_x86_64_glibc_228)
-
-    # Choose the best suitable libraries.
-    if(EXISTS ${LIBDIR_NATIVE_ABI})
-      set(LIBDIR ${LIBDIR_NATIVE_ABI})
+  # Choose the best suitable libraries.
+  if(EXISTS ${LIBDIR_NATIVE_ABI})
+    set(LIBDIR ${LIBDIR_NATIVE_ABI})
+    set(WITH_LIBC_MALLOC_HOOK_WORKAROUND True)
+  elseif(EXISTS ${LIBDIR_GLIBC228_ABI})
+    set(LIBDIR ${LIBDIR_GLIBC228_ABI})
+    if(WITH_MEM_JEMALLOC)
+      # jemalloc provides malloc hooks.
+      set(WITH_LIBC_MALLOC_HOOK_WORKAROUND False)
+    else()
       set(WITH_LIBC_MALLOC_HOOK_WORKAROUND True)
-    elseif(EXISTS ${LIBDIR_GLIBC228_ABI})
-      set(LIBDIR ${LIBDIR_GLIBC228_ABI})
-      if(WITH_MEM_JEMALLOC)
-        # jemalloc provides malloc hooks.
-        set(WITH_LIBC_MALLOC_HOOK_WORKAROUND False)
-      else()
-        set(WITH_LIBC_MALLOC_HOOK_WORKAROUND True)
-      endif()
     endif()
-
-    # Avoid namespace pollustion.
-    unset(LIBDIR_NATIVE_ABI)
-    unset(LIBDIR_GLIBC228_ABI)
   endif()
 
-  if(NOT (EXISTS ${LIBDIR}))
-    message(STATUS
-      "Unable to find LIBDIR: ${LIBDIR}, system libraries may be used "
-      "(disable WITH_LIBS_PRECOMPILED to suppress this message)."
-    )
-    unset(LIBDIR)
-  endif()
+  # Avoid namespace pollustion.
+  unset(LIBDIR_NATIVE_ABI)
+  unset(LIBDIR_GLIBC228_ABI)
 endif()
 
-
 # Support restoring this value once pre-compiled libraries have been handled.
 set(WITH_STATIC_LIBS_INIT ${WITH_STATIC_LIBS})
 
-if(DEFINED LIBDIR)
+if(EXISTS ${LIBDIR})
   message(STATUS "Using pre-compiled LIBDIR: ${LIBDIR}")
 
   file(GLOB LIB_SUBDIRS ${LIBDIR}/*)
@@ -99,7 +85,7 @@ endmacro()
 # These are libraries that may be precompiled. For this we disable searching in
 # the system directories so that we don't accidentally use them instead.
 
-if(DEFINED LIBDIR)
+if(EXISTS ${LIBDIR})
   without_system_libs_begin()
 endif()
 
@@ -111,7 +97,6 @@ find_package_wrapper(Epoxy REQUIRED)
 
 if(WITH_VULKAN_BACKEND)
   find_package_wrapper(Vulkan REQUIRED)
-  find_package_wrapper(ShaderC REQUIRED)
 endif()
 
 function(check_freetype_for_brotli)
@@ -129,7 +114,7 @@ endfunction()
 if(NOT WITH_SYSTEM_FREETYPE)
   # FreeType compiled with Brotli compression for woff2.
   find_package_wrapper(Freetype REQUIRED)
-  if(DEFINED LIBDIR)
+  if(EXISTS ${LIBDIR})
     find_package_wrapper(Brotli REQUIRED)
 
     # NOTE: This is done on WIN32 & APPLE but fails on some Linux systems.
@@ -156,7 +141,7 @@ if(WITH_PYTHON)
   if(WITH_PYTHON_MODULE AND NOT WITH_INSTALL_PORTABLE)
     # Installing into `site-packages`, warn when installing into `./../lib/`
     # which script authors almost certainly don't want.
-    if(DEFINED LIBDIR)
+    if(EXISTS ${LIBDIR})
       path_is_prefix(LIBDIR PYTHON_SITE_PACKAGES _is_prefix)
       if(_is_prefix)
         message(WARNING "
@@ -232,7 +217,7 @@ if(WITH_CODEC_SNDFILE)
 endif()
 
 if(WITH_CODEC_FFMPEG)
-  if(DEFINED LIBDIR)
+  if(EXISTS ${LIBDIR})
     set(FFMPEG_ROOT_DIR ${LIBDIR}/ffmpeg)
     # Override FFMPEG components to also include static library dependencies
     # included with precompiled libraries, and to ensure correct link order.
@@ -247,7 +232,7 @@ if(WITH_CODEC_FFMPEG)
       vpx
       x264
       xvidcore)
-    if((DEFINED LIBDIR) AND (EXISTS ${LIBDIR}/ffmpeg/lib/libaom.a))
+    if(EXISTS ${LIBDIR}/ffmpeg/lib/libaom.a)
       list(APPEND FFMPEG_FIND_COMPONENTS aom)
     endif()
   elseif(FFMPEG)
@@ -445,13 +430,10 @@ if(WITH_OPENIMAGEIO)
     ${PNG_LIBRARIES}
     ${JPEG_LIBRARIES}
     ${ZLIB_LIBRARIES}
+    ${BOOST_LIBRARIES}
   )
-
   set(OPENIMAGEIO_DEFINITIONS "")
 
-  if(WITH_BOOST)
-    list(APPEND OPENIMAGEIO_LIBRARIES "${BOOST_LIBRARIES}")
-  endif()
   if(WITH_IMAGE_TIFF)
     list(APPEND OPENIMAGEIO_LIBRARIES "${TIFF_LIBRARY}")
   endif()
@@ -469,7 +451,7 @@ add_bundled_libraries(openimageio/lib)
 if(WITH_OPENCOLORIO)
   find_package_wrapper(OpenColorIO 2.0.0)
 
-  set(OPENCOLORIO_DEFINITIONS "")
+  set(OPENCOLORIO_DEFINITIONS)
   set_and_warn_library_found("OpenColorIO" OPENCOLORIO_FOUND WITH_OPENCOLORIO)
 endif()
 add_bundled_libraries(opencolorio/lib)
@@ -484,7 +466,7 @@ if(WITH_OPENIMAGEDENOISE)
 endif()
 
 if(WITH_LLVM)
-  if(DEFINED LIBDIR)
+  if(EXISTS ${LIBDIR})
     set(LLVM_STATIC ON)
   endif()
 
@@ -498,7 +480,7 @@ if(WITH_LLVM)
     endif()
 
     # Symbol conflicts with same UTF library used by OpenCollada
-    if(DEFINED LIBDIR)
+    if(EXISTS ${LIBDIR})
       if(WITH_OPENCOLLADA AND (${LLVM_VERSION} VERSION_LESS "4.0.0"))
         list(REMOVE_ITEM OPENCOLLADA_LIBRARIES ${OPENCOLLADA_UTF_LIBRARY})
       endif()
@@ -554,7 +536,7 @@ if(WITH_CYCLES AND WITH_CYCLES_PATH_GUIDING)
   endif()
 endif()
 
-if(DEFINED LIBDIR)
+if(EXISTS ${LIBDIR})
   without_system_libs_end()
 endif()
 
@@ -569,14 +551,9 @@ else()
 endif()
 
 find_package(Threads REQUIRED)
-# `FindThreads` documentation notes that this may be empty
-# with the system libraries provide threading functionality.
-if(CMAKE_THREAD_LIBS_INIT)
-  list(APPEND PLATFORM_LINKLIBS ${CMAKE_THREAD_LIBS_INIT})
-  # used by other platforms
-  set(PTHREADS_LIBRARIES ${CMAKE_THREAD_LIBS_INIT})
-endif()
-
+list(APPEND PLATFORM_LINKLIBS ${CMAKE_THREAD_LIBS_INIT})
+# used by other platforms
+set(PTHREADS_LIBRARIES ${CMAKE_THREAD_LIBS_INIT})
 
 if(CMAKE_DL_LIBS)
   list(APPEND PLATFORM_LINKLIBS ${CMAKE_DL_LIBS})
@@ -598,7 +575,7 @@ add_definitions(-D_LARGEFILE_SOURCE -D_FILE_OFFSET_BITS=64 -D_LARGEFILE64_SOURCE
 #
 # Keep last, so indirectly linked libraries don't override our own pre-compiled libs.
 
-if(DEFINED LIBDIR)
+if(EXISTS ${LIBDIR})
   # Clear the prefix path as it causes the `LIBDIR` to override system locations.
   unset(CMAKE_PREFIX_PATH)
 
@@ -654,7 +631,7 @@ if(WITH_GHOST_WAYLAND)
   # When dynamically linked WAYLAND is used and `${LIBDIR}/wayland` is present,
   # there is no need to search for the libraries as they are not needed for building.
   # Only the headers are needed which can reference the known paths.
-  if((DEFINED LIBDIR) AND (EXISTS "${LIBDIR}/wayland" AND WITH_GHOST_WAYLAND_DYNLOAD))
+  if(EXISTS "${LIBDIR}/wayland" AND WITH_GHOST_WAYLAND_DYNLOAD)
     set(_use_system_wayland OFF)
   else()
     set(_use_system_wayland ON)
@@ -718,7 +695,7 @@ if(WITH_GHOST_WAYLAND)
       add_definitions(-DWITH_GHOST_WAYLAND_LIBDECOR)
     endif()
 
-    if((DEFINED LIBDIR) AND (EXISTS "${LIBDIR}/wayland/bin/wayland-scanner"))
+    if(EXISTS "${LIBDIR}/wayland/bin/wayland-scanner")
       set(WAYLAND_SCANNER "${LIBDIR}/wayland/bin/wayland-scanner")
     else()
       pkg_get_variable(WAYLAND_SCANNER wayland-scanner wayland_scanner)

build_files/cmake/platform/platform_win32.cmake

@@ -121,7 +121,7 @@ if(WITH_WINDOWS_BUNDLE_CRT)
   include(InstallRequiredSystemLibraries)
 
   # ucrtbase(d).dll cannot be in the manifest, due to the way windows 10 handles
-  # redirects for this dll, for details see #88813.
+  # redirects for this dll, for details see T88813.
   foreach(lib ${CMAKE_INSTALL_SYSTEM_RUNTIME_LIBS})
     string(FIND ${lib} "ucrtbase" pos)
     if(NOT pos EQUAL -1)
@@ -295,7 +295,7 @@ unset(MATERIALX_LIB_FOLDER_EXISTS)
 if(NOT MSVC_CLANG                  AND # Available with MSVC 15.7+ but not for CLANG.
    NOT WITH_WINDOWS_SCCACHE        AND # And not when sccache is enabled
    NOT VS_CLANG_TIDY)                  # Clang-tidy does not like these options
-  add_compile_options(/experimental:external /external:I "${LIBDIR}" /external:W0)
+  add_compile_options(/experimental:external /external:templates- /external:I "${LIBDIR}" /external:W0)
 endif()
 
 # Add each of our libraries to our cmake_prefix_path so find_package() could work

build_files/config/pipeline_config.yaml

@@ -5,16 +5,16 @@
 update-code:
     git:
         submodules:
-        -   branch: main
+        -   branch: master
             commit_id: HEAD
             path: release/scripts/addons
-        -   branch: main
+        -   branch: master
             commit_id: HEAD
             path: release/scripts/addons_contrib
-        -   branch: main
+        -   branch: master
             commit_id: HEAD
             path: release/datafiles/locale
-        -   branch: main
+        -   branch: master
             commit_id: HEAD
             path: source/tools
     svn:
@@ -43,10 +43,6 @@ update-code:
             branch: trunk
             commit_id: HEAD
             path: lib/benchmarks
-        assets:
-            branch: trunk
-            commit_id: HEAD
-            path: lib/assets
 
 #
 # Buildbot only configs
@@ -63,7 +59,7 @@ buildbot:
     optix:
         version: '7.3.0'
     ocloc:
-        version: '101.4032'
+        version: '101.3430'
     cmake:
         default:
             version: any

build_files/utils/make_bpy_wheel.py

@@ -24,7 +24,7 @@ import os
 import re
 import platform
 import string
-import setuptools
+import setuptools  # type: ignore
 import sys
 
 from typing import (
@@ -58,7 +58,7 @@ Each Blender release supports one Python version, and the package is only compat
 ## Source Code
 
 * [Releases](https://download.blender.org/source/)
-* Repository: [projects.blender.org/blender/blender.git](https://projects.blender.org/blender/blender)
+* Repository: [git.blender.org/blender.git](https://git.blender.org/gitweb/gitweb.cgi/blender.git)
 
 ## Credits
 
@@ -208,7 +208,7 @@ def main() -> None:
         return paths
 
     # Ensure this wheel is marked platform specific.
-    class BinaryDistribution(setuptools.dist.Distribution):
+    class BinaryDistribution(setuptools.dist.Distribution):  # type: ignore
         def has_ext_modules(self) -> bool:
             return True
 

build_files/utils/make_test.py

@@ -13,10 +13,10 @@ import sys
 import make_utils
 from make_utils import call
 
-# Parse arguments.
+# Parse arguments
 
 
-def parse_arguments() -> argparse.Namespace:
+def parse_arguments():
     parser = argparse.ArgumentParser()
     parser.add_argument("--ctest-command", default="ctest")
     parser.add_argument("--cmake-command", default="cmake")

build_files/utils/make_update.py

@@ -42,7 +42,6 @@ def parse_arguments() -> argparse.Namespace:
     parser.add_argument("--svn-branch", default=None)
     parser.add_argument("--git-command", default="git")
     parser.add_argument("--use-linux-libraries", action="store_true")
-    parser.add_argument("--architecture", type=str, choices=("x86_64", "amd64", "arm64",))
     return parser.parse_args()
 
 
@@ -52,17 +51,6 @@ def get_blender_git_root() -> str:
 # Setup for precompiled libraries and tests from svn.
 
 
-def get_effective_architecture(args: argparse.Namespace):
-    if args.architecture:
-        return args.architecture
-
-    # Check platform.version to detect arm64 with x86_64 python binary.
-    if "ARM64" in platform.version():
-        return "arm64"
-
-    return platform.machine().lower()
-
-
 def svn_update(args: argparse.Namespace, release_version: Optional[str]) -> None:
     svn_non_interactive = [args.svn_command, '--non-interactive']
 
@@ -70,11 +58,11 @@ def svn_update(args: argparse.Namespace, release_version: Optional[str]) -> None
     svn_url = make_utils.svn_libraries_base_url(release_version, args.svn_branch)
 
     # Checkout precompiled libraries
-    architecture = get_effective_architecture(args)
     if sys.platform == 'darwin':
-        if architecture == 'arm64':
+        # Check platform.version to detect arm64 with x86_64 python binary.
+        if platform.machine() == 'arm64' or ('ARM64' in platform.version()):
             lib_platform = "darwin_arm64"
-        elif architecture == 'x86_64':
+        elif platform.machine() == 'x86_64':
             lib_platform = "darwin"
         else:
             lib_platform = None
@@ -116,30 +104,17 @@ def svn_update(args: argparse.Namespace, release_version: Optional[str]) -> None
             svn_url_tests = svn_url + lib_tests
             call(svn_non_interactive + ["checkout", svn_url_tests, lib_tests_dirpath])
 
-    lib_assets = "assets"
-    lib_assets_dirpath = os.path.join(lib_dirpath, lib_assets)
-
-    if not os.path.exists(lib_assets_dirpath):
-        print_stage("Checking out Assets")
-
-        if make_utils.command_missing(args.svn_command):
-            sys.stderr.write("svn not found, can't checkout assets\n")
-            sys.exit(1)
-
-        svn_url_assets = svn_url + lib_assets
-        call(svn_non_interactive + ["checkout", svn_url_assets, lib_assets_dirpath])
-
-    # Update precompiled libraries, assets and tests
+    # Update precompiled libraries and tests
 
     if not os.path.isdir(lib_dirpath):
         print("Library path: %r, not found, skipping" % lib_dirpath)
     else:
         paths_local_and_remote = []
         if os.path.exists(os.path.join(lib_dirpath, ".svn")):
-            print_stage("Updating Precompiled Libraries, Assets and Tests (one repository)")
+            print_stage("Updating Precompiled Libraries and Tests (one repository)")
             paths_local_and_remote.append((lib_dirpath, svn_url))
         else:
-            print_stage("Updating Precompiled Libraries, Assets and Tests (multiple repositories)")
+            print_stage("Updating Precompiled Libraries and Tests (multiple repositories)")
             # Separate paths checked out.
             for dirname in os.listdir(lib_dirpath):
                 if dirname.startswith("."):
@@ -182,7 +157,7 @@ def git_update_skip(args: argparse.Namespace, check_remote_exists: bool = True)
         return "rebase or merge in progress, complete it first"
 
     # Abort if uncommitted changes.
-    changes = check_output([args.git_command, 'status', '--porcelain', '--untracked-files=no', '--ignore-submodules'])
+    changes = check_output([args.git_command, 'status', '--porcelain', '--untracked-files=no'])
     if len(changes) != 0:
         return "you have unstaged changes"
 
@@ -214,8 +189,8 @@ def submodules_update(
         sys.exit(1)
 
     # Update submodules to appropriate given branch,
-    # falling back to main if none is given and/or found in a sub-repository.
-    branch_fallback = "main"
+    # falling back to master if none is given and/or found in a sub-repository.
+    branch_fallback = "master"
     if not branch:
         branch = branch_fallback
 
@@ -268,15 +243,14 @@ if __name__ == "__main__":
     blender_skip_msg = ""
     submodules_skip_msg = ""
 
-    blender_version = make_utils. parse_blender_version()
-    if blender_version.cycle != 'alpha':
-        major = blender_version.version // 100
-        minor = blender_version.version % 100
-        branch = f"blender-v{major}.{minor}-release"
-        release_version = f"{major}.{minor}"
-    else:
-        branch = 'main'
-        release_version = None
+    # Test if we are building a specific release version.
+    branch = make_utils.git_branch(args.git_command)
+    if branch == 'HEAD':
+        sys.stderr.write('Blender git repository is in detached HEAD state, must be in a branch\n')
+        sys.exit(1)
+
+    tag = make_utils.git_tag(args.git_command)
+    release_version = make_utils.git_branch_release_version(branch, tag)
 
     if not args.no_libraries:
         svn_update(args, release_version)

build_files/windows/check_submodules.cmd

@@ -3,9 +3,9 @@ if NOT exist "%BLENDER_DIR%\source\tools\.git" (
 	if not "%GIT%" == "" (
 		"%GIT%" submodule update --init --recursive --progress
 		if errorlevel 1 goto FAIL
-		"%GIT%" submodule foreach git checkout main
+		"%GIT%" submodule foreach git checkout master
 		if errorlevel 1 goto FAIL
-		"%GIT%" submodule foreach git pull --rebase origin main
+		"%GIT%" submodule foreach git pull --rebase origin master
 		if errorlevel 1 goto FAIL
 		goto EOF
 	) else (

doc/doxygen/Doxyfile

@@ -38,7 +38,7 @@ PROJECT_NAME           = Blender
 # could be handy for archiving the generated documentation or if some version
 # control system is used.
 
-PROJECT_NUMBER         = V3.6
+PROJECT_NUMBER         = V3.5
 
 # Using the PROJECT_BRIEF tag one can provide an optional one line description
 # for a project that appears at the top of each page and should give viewer a

doc/python_api/examples/blf.py

@@ -37,7 +37,7 @@ def draw_callback_px(self, context):
     # BLF drawing routine
     font_id = font_info["font_id"]
     blf.position(font_id, 2, 80, 0)
-    blf.size(font_id, 50)
+    blf.size(font_id, 50, 72)
     blf.draw(font_id, "Hello World")
 
 

doc/python_api/sphinx_doc_gen.py

@@ -476,7 +476,7 @@ MODULE_GROUPING = {
 
 # -------------------------------BLENDER----------------------------------------
 
-# Converting bytes to strings, due to #30154.
+# converting bytes to strings, due to T30154
 BLENDER_REVISION = str(bpy.app.build_hash, 'utf_8')
 BLENDER_REVISION_TIMESTAMP = bpy.app.build_commit_timestamp
 
@@ -487,7 +487,7 @@ BLENDER_VERSION_DOTS = "%d.%d" % (bpy.app.version[0], bpy.app.version[1])
 if BLENDER_REVISION != "Unknown":
     # SHA1 Git hash
     BLENDER_VERSION_HASH = BLENDER_REVISION
-    BLENDER_VERSION_HASH_HTML_LINK = "<a href=https://projects.blender.org/blender/blender/commit/%s>%s</a>" % (
+    BLENDER_VERSION_HASH_HTML_LINK = "<a href=https://developer.blender.org/rB%s>%s</a>" % (
         BLENDER_VERSION_HASH, BLENDER_VERSION_HASH,
     )
     BLENDER_VERSION_DATE = time.strftime("%d/%m/%Y", time.localtime(BLENDER_REVISION_TIMESTAMP))
@@ -647,7 +647,7 @@ def undocumented_message(module_name, type_name, identifier):
         module_name, type_name, identifier,
     )
 
-    return "Undocumented, consider `contributing <https://developer.blender.org/>`__."
+    return "Undocumented, consider `contributing <https://developer.blender.org/T51061>`__."
 
 
 def range_str(val):
@@ -1816,9 +1816,9 @@ def pyrna2sphinx(basepath):
 
     # operators
     def write_ops():
-        API_BASEURL = "https://projects.blender.org/blender/blender/src/branch/main/release/scripts"
-        API_BASEURL_ADDON = "https://projects.blender.org/blender/blender-addons"
-        API_BASEURL_ADDON_CONTRIB = "https://projects.blender.org/blender/blender-addons-contrib"
+        API_BASEURL = "https://developer.blender.org/diffusion/B/browse/master/release/scripts"
+        API_BASEURL_ADDON = "https://developer.blender.org/diffusion/BA"
+        API_BASEURL_ADDON_CONTRIB = "https://developer.blender.org/diffusion/BAC"
 
         op_modules = {}
         op = None
@@ -2098,8 +2098,6 @@ def write_rst_types_index(basepath):
         fw(title_string("Types (bpy.types)", "="))
         fw(".. module:: bpy.types\n\n")
         fw(".. toctree::\n")
-        # Only show top-level entries (avoids unreasonably large pages).
-        fw("   :maxdepth: 1\n")
         fw("   :glob:\n\n")
         fw("   bpy.types.*\n\n")
 
@@ -2126,8 +2124,6 @@ def write_rst_ops_index(basepath):
         write_example_ref("", fw, "bpy.ops")
         fw(".. toctree::\n")
         fw("   :caption: Submodules\n")
-        # Only show top-level entries (avoids unreasonably large pages).
-        fw("   :maxdepth: 1\n")
         fw("   :glob:\n\n")
         fw("   bpy.ops.*\n\n")
         file.close()
@@ -2200,7 +2196,7 @@ def write_rst_enum_items(basepath, key, key_no_prefix, enum_items):
     Write a single page for a static enum in RST.
 
     This helps avoiding very large lists being in-lined in many places which is an issue
-    especially with icons in ``bpy.types.UILayout``. See #87008.
+    especially with icons in ``bpy.types.UILayout``. See T87008.
     """
     filepath = os.path.join(basepath, "%s.rst" % key_no_prefix)
     with open(filepath, "w", encoding="utf-8") as fh:

doc/python_api/static/js/version_switch.js

@@ -156,7 +156,7 @@ var Popover = function() {
     },
     getNamed : function(v) {
       $.each(all_versions, function(ix, title) {
-        if (ix === "master" || ix === "main" || ix === "latest") {
+        if (ix === "master" || ix === "latest") {
           var m = title.match(/\d\.\d[\w\d\.]*/)[0];
           if (parseFloat(m) == v) {
             v = ix;

extern/hipew/README.blender

@@ -1,5 +1,5 @@
 Project: Blender
-URL: https://projects.blender.org/blender/blender.git
+URL: https://git.blender.org/blender.git
 License: Apache 2.0
 Upstream version: N/A
 Local modifications: None

extern/mantaflow/CMakeLists.txt

@@ -13,12 +13,10 @@ endif()
 
 # Exporting functions from the blender binary gives linker warnings on Apple arm64 systems.
 # Silence them here.
-if(APPLE)
-  if("${CMAKE_OSX_ARCHITECTURES}" STREQUAL "arm64")
-    if(CMAKE_COMPILER_IS_GNUCXX OR "${CMAKE_CXX_COMPILER_ID}" MATCHES "Clang")
-      string(APPEND CMAKE_C_FLAGS " -fvisibility=hidden")
-      string(APPEND CMAKE_CXX_FLAGS " -fvisibility=hidden")
-    endif()
+if(APPLE AND ("${CMAKE_OSX_ARCHITECTURES}" STREQUAL "arm64"))
+  if(CMAKE_COMPILER_IS_GNUCXX OR "${CMAKE_CXX_COMPILER_ID}" MATCHES "Clang")
+    string(APPEND CMAKE_C_FLAGS " -fvisibility=hidden")
+    string(APPEND CMAKE_CXX_FLAGS " -fvisibility=hidden")
   endif()
 endif()
 
@@ -263,11 +261,9 @@ set(LIB
 
 blender_add_lib(extern_mantaflow "${SRC}" "${INC}" "${INC_SYS}" "${LIB}")
 
-if(WITH_OPENVDB)
-  # The VDB libs above are only added to as INTERFACE libs by blender_add_lib,
-  # meaning extern_mantaflow itself actually does not have a dependency on the
-  # openvdb libraries, and CMAKE is free to link the vdb libs before
-  # extern_mantaflow causing linker errors on linux. By explicitly declaring
-  # a dependency here, cmake will do the right thing.
-  target_link_libraries(extern_mantaflow PRIVATE ${OPENVDB_LIBRARIES})
-endif()
+# The VDB libs above are only added to as INTERFACE libs by blender_add_lib,
+# meaning extern_mantaflow itself actually does not have a dependency on the
+# openvdb libraries, and CMAKE is free to link the vdb libs before
+# extern_mantaflow causing linker errors on linux. By explicitly declaring
+# a dependency here, cmake will do the right thing.
+target_link_libraries(extern_mantaflow PRIVATE ${OPENVDB_LIBRARIES})

extern/vulkan_memory_allocator/CMakeLists.txt

@@ -7,7 +7,6 @@ set(INC
 
 set(INC_SYS
   ${VULKAN_INCLUDE_DIRS}
-  ${MOLTENVK_INCLUDE_DIRS}
 )
 
 set(SRC

extern/vulkan_memory_allocator/patches/remove_compilation_warning.diff

@@ -1,15 +0,0 @@
-diff --git a/extern/vulkan_memory_allocator/vk_mem_alloc.h b/extern/vulkan_memory_allocator/vk_mem_alloc.h
-index 60f572038c0..63a9994ba46 100644
---- a/extern/vulkan_memory_allocator/vk_mem_alloc.h
-+++ b/extern/vulkan_memory_allocator/vk_mem_alloc.h
-@@ -13371,8 +13371,8 @@ bool VmaDefragmentationContext_T::IncrementCounters(VkDeviceSize bytes)
-     // Early return when max found
-     if (++m_PassStats.allocationsMoved >= m_MaxPassAllocations || m_PassStats.bytesMoved >= m_MaxPassBytes)
-     {
--        VMA_ASSERT(m_PassStats.allocationsMoved == m_MaxPassAllocations ||
--            m_PassStats.bytesMoved == m_MaxPassBytes && "Exceeded maximal pass threshold!");
-+        VMA_ASSERT((m_PassStats.allocationsMoved == m_MaxPassAllocations ||
-+            m_PassStats.bytesMoved == m_MaxPassBytes) && "Exceeded maximal pass threshold!");
-         return true;
-     }
-     return false;

intern/cycles/blender/addon/properties.py

@@ -12,7 +12,6 @@ from bpy.props import (
     PointerProperty,
     StringProperty,
 )
-from bpy.app.translations import pgettext_iface as iface_
 
 from math import pi
 
@@ -1665,51 +1664,30 @@ class CyclesPreferences(bpy.types.AddonPreferences):
             col.label(text="No compatible GPUs found for Cycles", icon='INFO')
 
             if device_type == 'CUDA':
-                compute_capability = "3.0"
-                col.label(text=iface_("Requires NVIDIA GPU with compute capability %s") % compute_capability,
-                          icon='BLANK1', translate=False)
+                col.label(text="Requires NVIDIA GPU with compute capability 3.0", icon='BLANK1')
             elif device_type == 'OPTIX':
-                compute_capability = "5.0"
-                driver_version = "470"
-                col.label(text=iface_("Requires NVIDIA GPU with compute capability %s") % compute_capability,
-                          icon='BLANK1', translate=False)
-                col.label(text="and NVIDIA driver version %s or newer" % driver_version,
-                          icon='BLANK1', translate=False)
+                col.label(text="Requires NVIDIA GPU with compute capability 5.0", icon='BLANK1')
+                col.label(text="and NVIDIA driver version 470 or newer", icon='BLANK1')
             elif device_type == 'HIP':
-                if True:
-                    col.label(text="HIP temporarily disabled due to compiler bugs", icon='BLANK1')
-                else:
-                    import sys
-                    if sys.platform[:3] == "win":
-                        driver_version = "21.Q4"
-                        col.label(text="Requires AMD GPU with Vega or RDNA architecture", icon='BLANK1')
-                        col.label(text=iface_("and AMD Radeon Pro %s driver or newer") % driver_version,
-                                  icon='BLANK1', translate=False)
-                    elif sys.platform.startswith("linux"):
-                        driver_version = "22.10"
-                        col.label(text="Requires AMD GPU with Vega or RDNA architecture", icon='BLANK1')
-                        col.label(text=iface_("and AMD driver version %s or newer") % driver_version, icon='BLANK1',
-                                  translate=False)
+                import sys
+                if sys.platform[:3] == "win":
+                    col.label(text="Requires AMD GPU with Vega or RDNA architecture", icon='BLANK1')
+                    col.label(text="and AMD Radeon Pro 21.Q4 driver or newer", icon='BLANK1')
+                elif sys.platform.startswith("linux"):
+                    col.label(text="Requires AMD GPU with Vega or RDNA architecture", icon='BLANK1')
+                    col.label(text="and AMD driver version 22.10 or newer", icon='BLANK1')
             elif device_type == 'ONEAPI':
                 import sys
                 if sys.platform.startswith("win"):
-                    driver_version = "101.4032"
                     col.label(text="Requires Intel GPU with Xe-HPG architecture", icon='BLANK1')
-                    col.label(text=iface_("and Windows driver version %s or newer") % driver_version,
-                              icon='BLANK1', translate=False)
+                    col.label(text="and Windows driver version 101.3430 or newer", icon='BLANK1')
                 elif sys.platform.startswith("linux"):
-                    driver_version = "1.3.24931"
                     col.label(text="Requires Intel GPU with Xe-HPG architecture and", icon='BLANK1')
-                    col.label(text=iface_("  - intel-level-zero-gpu version %s or newer") % driver_version,
-                              icon='BLANK1', translate=False)
+                    col.label(text="  - intel-level-zero-gpu version 1.3.23904 or newer", icon='BLANK1')
                     col.label(text="  - oneAPI Level-Zero Loader", icon='BLANK1')
             elif device_type == 'METAL':
-                silicon_mac_version = "12.2"
-                amd_mac_version = "12.3"
-                col.label(text=iface_("Requires Apple Silicon with macOS %s or newer") % silicon_mac_version,
-                          icon='BLANK1', translate=False)
-                col.label(text=iface_("or AMD with macOS %s or newer") % amd_mac_version, icon='BLANK1',
-                          translate=False)
+                col.label(text="Requires Apple Silicon with macOS 12.2 or newer", icon='BLANK1')
+                col.label(text="or AMD with macOS 12.3 or newer", icon='BLANK1')
             return
 
         for device in devices:
@@ -1745,21 +1723,12 @@ class CyclesPreferences(bpy.types.AddonPreferences):
 
         if compute_device_type == 'METAL':
             import platform
-            import re
-            is_navi_2 = False
-            for device in devices:
-                if re.search(r"((RX)|(Pro)|(PRO))\s+W?6\d00X", device.name):
-                    is_navi_2 = True
-                    break
-
-            # MetalRT only works on Apple Silicon and Navi2.
-            is_arm64 = platform.machine() == 'arm64'
-            if is_arm64 or is_navi_2:
+            # MetalRT only works on Apple Silicon at present, pending argument encoding fixes on AMD
+            # Kernel specialization is only viable on Apple Silicon at present due to relative compilation speed
+            if platform.machine() == 'arm64':
                 col = layout.column()
                 col.use_property_split = True
-                # Kernel specialization is only supported on Apple Silicon
-                if is_arm64:
-                    col.prop(self, "kernel_optimization_level")
+                col.prop(self, "kernel_optimization_level")
                 col.prop(self, "use_metalrt")
 
     def draw(self, context):

intern/cycles/blender/addon/ui.py

@@ -20,7 +20,7 @@ class CyclesPresetPanel(PresetPanel, Panel):
     @staticmethod
     def post_cb(context):
         # Modify an arbitrary built-in scene property to force a depsgraph
-        # update, because add-on properties don't. (see #62325)
+        # update, because add-on properties don't. (see T62325)
         render = context.scene.render
         render.filter_size = render.filter_size
 

intern/cycles/blender/display_driver.cpp

@@ -105,12 +105,11 @@ GPUShader *BlenderFallbackDisplayShader::bind(int width, int height)
 
   /* Bind shader now to enable uniform assignment. */
   GPU_shader_bind(shader_program_);
-  int slot = 0;
-  GPU_shader_uniform_int_ex(shader_program_, image_texture_location_, 1, 1, &slot);
+  GPU_shader_uniform_int(shader_program_, image_texture_location_, 0);
   float size[2];
   size[0] = width;
   size[1] = height;
-  GPU_shader_uniform_float_ex(shader_program_, fullscreen_location_, 2, 1, size);
+  GPU_shader_uniform_vector(shader_program_, fullscreen_location_, 2, 1, size);
   return shader_program_;
 }
 

intern/cycles/blender/image.cpp

@@ -20,7 +20,7 @@ BlenderImageLoader::BlenderImageLoader(BL::Image b_image,
     : b_image(b_image),
       frame(frame),
       tile_number(tile_number),
-      /* Don't free cache for preview render to avoid race condition from #93560, to be fixed
+      /* Don't free cache for preview render to avoid race condition from T93560, to be fixed
        * properly later as we are close to release. */
       free_cache(!is_preview_render && !b_image.has_data())
 {
@@ -72,7 +72,7 @@ bool BlenderImageLoader::load_metadata(const ImageDeviceFeatures &, ImageMetaDat
     metadata.colorspace = u_colorspace_raw;
   }
   else {
-    /* In some cases (e.g. #94135), the colorspace setting in Blender gets updated as part of the
+    /* In some cases (e.g. T94135), the colorspace setting in Blender gets updated as part of the
      * metadata queries in this function, so update the colorspace setting here. */
     PointerRNA colorspace_ptr = b_image.colorspace_settings().ptr;
     metadata.colorspace = get_enum_identifier(colorspace_ptr, "name");

intern/cycles/blender/light.cpp

@@ -24,7 +24,7 @@ void BlenderSync::sync_light(BL::Object &b_parent,
   Light *light = light_map.find(key);
 
   /* Check if the transform was modified, in case a linked collection is moved we do not get a
-   * specific depsgraph update (#88515). This also mimics the behavior for Objects. */
+   * specific depsgraph update (T88515). This also mimics the behavior for Objects. */
   const bool tfm_updated = (light && light->get_tfm() != tfm);
 
   /* Update if either object or light data changed. */
@@ -48,8 +48,6 @@ void BlenderSync::sync_light(BL::Object &b_parent,
     case BL::Light::type_SPOT: {
       BL::SpotLight b_spot_light(b_light);
       light->set_size(b_spot_light.shadow_soft_size());
-      light->set_axisu(transform_get_column(&tfm, 0));
-      light->set_axisv(transform_get_column(&tfm, 1));
       light->set_light_type(LIGHT_SPOT);
       light->set_spot_angle(b_spot_light.spot_size());
       light->set_spot_smooth(b_spot_light.spot_blend());

intern/cycles/blender/python.cpp

@@ -94,7 +94,7 @@ void python_thread_state_restore(void **python_thread_state)
   *python_thread_state = NULL;
 }
 
-static const char *PyC_UnicodeAsBytes(PyObject *py_str, PyObject **coerce)
+static const char *PyC_UnicodeAsByte(PyObject *py_str, PyObject **coerce)
 {
   const char *result = PyUnicode_AsUTF8(py_str);
   if (result) {
@@ -131,8 +131,8 @@ static PyObject *init_func(PyObject * /*self*/, PyObject *args)
   }
 
   PyObject *path_coerce = nullptr, *user_path_coerce = nullptr;
-  path_init(PyC_UnicodeAsBytes(path, &path_coerce),
-            PyC_UnicodeAsBytes(user_path, &user_path_coerce));
+  path_init(PyC_UnicodeAsByte(path, &path_coerce),
+            PyC_UnicodeAsByte(user_path, &user_path_coerce));
   Py_XDECREF(path_coerce);
   Py_XDECREF(user_path_coerce);
 

intern/cycles/blender/session.cpp

@@ -404,7 +404,7 @@ void BlenderSession::render(BL::Depsgraph &b_depsgraph_)
      * point we know that we've got everything to render current view layer.
      */
     /* At the moment we only free if we are not doing multi-view
-     * (or if we are rendering the last view). See #58142/D4239 for discussion.
+     * (or if we are rendering the last view). See T58142/D4239 for discussion.
      */
     if (view_index == num_views - 1) {
       free_blender_memory_if_possible();

intern/cycles/blender/sync.cpp

@@ -766,7 +766,7 @@ void BlenderSync::free_data_after_sync(BL::Depsgraph &b_depsgraph)
       (BlenderSession::headless || is_interface_locked) &&
       /* Baking re-uses the depsgraph multiple times, clearing crashes
        * reading un-evaluated mesh data which isn't aligned with the
-       * geometry we're baking, see #71012. */
+       * geometry we're baking, see T71012. */
       !scene->bake_manager->get_baking() &&
       /* Persistent data must main caches for performance and correctness. */
       !is_persistent_data;

intern/cycles/cmake/external_libs.cmake

@@ -42,15 +42,12 @@ endif()
 ###########################################################################
 
 if(WITH_CYCLES_HIP_BINARIES AND WITH_CYCLES_DEVICE_HIP)
-  set(WITH_CYCLES_HIP_BINARIES OFF)
-  message(STATUS "HIP temporarily disabled due to compiler bugs")
+  find_package(HIP)
+  set_and_warn_library_found("HIP compiler" HIP_FOUND WITH_CYCLES_HIP_BINARIES)
 
-  # find_package(HIP)
-  # set_and_warn_library_found("HIP compiler" HIP_FOUND WITH_CYCLES_HIP_BINARIES)
-
-  # if(HIP_FOUND)
-  #   message(STATUS "Found HIP ${HIP_HIPCC_EXECUTABLE} (${HIP_VERSION})")
-  # endif()
+  if(HIP_FOUND)
+    message(STATUS "Found HIP ${HIP_HIPCC_EXECUTABLE} (${HIP_VERSION})")
+  endif()
 endif()
 
 if(NOT WITH_HIP_DYNLOAD)

intern/cycles/cmake/macros.cmake

@@ -111,10 +111,8 @@ macro(cycles_external_libraries_append libraries)
   endif()
   if(WITH_OPENIMAGEDENOISE)
     list(APPEND ${libraries} ${OPENIMAGEDENOISE_LIBRARIES})
-    if(APPLE)
-      if("${CMAKE_OSX_ARCHITECTURES}" STREQUAL "arm64")
-        list(APPEND ${libraries} "-framework Accelerate")
-      endif()
+    if(APPLE AND "${CMAKE_OSX_ARCHITECTURES}" STREQUAL "arm64")
+      list(APPEND ${libraries} "-framework Accelerate")
     endif()
   endif()
   if(WITH_ALEMBIC)
@@ -138,15 +136,7 @@ macro(cycles_external_libraries_append libraries)
     ${PYTHON_LIBRARIES}
     ${ZLIB_LIBRARIES}
     ${CMAKE_DL_LIBS}
-  )
-
-  if(DEFINED PTHREADS_LIBRARIES)
-    list(APPEND ${libraries}
-      ${PTHREADS_LIBRARIES}
-    )
-  endif()
-
-  list(APPEND ${libraries}
+    ${PTHREADS_LIBRARIES}
     ${PLATFORM_LINKLIBS}
   )
 

intern/cycles/device/cuda/device_impl.cpp

@@ -53,12 +53,8 @@ void CUDADevice::set_error(const string &error)
 }
 
 CUDADevice::CUDADevice(const DeviceInfo &info, Stats &stats, Profiler &profiler)
-    : GPUDevice(info, stats, profiler)
+    : Device(info, stats, profiler), texture_info(this, "texture_info", MEM_GLOBAL)
 {
-  /* Verify that base class types can be used with specific backend types */
-  static_assert(sizeof(texMemObject) == sizeof(CUtexObject));
-  static_assert(sizeof(arrayMemObject) == sizeof(CUarray));
-
   first_error = true;
 
   cuDevId = info.num;
@@ -69,6 +65,12 @@ CUDADevice::CUDADevice(const DeviceInfo &info, Stats &stats, Profiler &profiler)
 
   need_texture_info = false;
 
+  device_texture_headroom = 0;
+  device_working_headroom = 0;
+  move_texture_to_host = false;
+  map_host_limit = 0;
+  map_host_used = 0;
+  can_map_host = 0;
   pitch_alignment = 0;
 
   /* Initialize CUDA. */
@@ -89,9 +91,8 @@ CUDADevice::CUDADevice(const DeviceInfo &info, Stats &stats, Profiler &profiler)
   /* CU_CTX_MAP_HOST for mapping host memory when out of device memory.
    * CU_CTX_LMEM_RESIZE_TO_MAX for reserving local memory ahead of render,
    * so we can predict which memory to map to host. */
-  int value;
-  cuda_assert(cuDeviceGetAttribute(&value, CU_DEVICE_ATTRIBUTE_CAN_MAP_HOST_MEMORY, cuDevice));
-  can_map_host = value != 0;
+  cuda_assert(
+      cuDeviceGetAttribute(&can_map_host, CU_DEVICE_ATTRIBUTE_CAN_MAP_HOST_MEMORY, cuDevice));
 
   cuda_assert(cuDeviceGetAttribute(
       &pitch_alignment, CU_DEVICE_ATTRIBUTE_TEXTURE_PITCH_ALIGNMENT, cuDevice));
@@ -498,57 +499,311 @@ void CUDADevice::reserve_local_memory(const uint kernel_features)
 #  endif
 }
 
-void CUDADevice::get_device_memory_info(size_t &total, size_t &free)
+void CUDADevice::init_host_memory()
+{
+  /* Limit amount of host mapped memory, because allocating too much can
+   * cause system instability. Leave at least half or 4 GB of system
+   * memory free, whichever is smaller. */
+  size_t default_limit = 4 * 1024 * 1024 * 1024LL;
+  size_t system_ram = system_physical_ram();
+
+  if (system_ram > 0) {
+    if (system_ram / 2 > default_limit) {
+      map_host_limit = system_ram - default_limit;
+    }
+    else {
+      map_host_limit = system_ram / 2;
+    }
+  }
+  else {
+    VLOG_WARNING << "Mapped host memory disabled, failed to get system RAM";
+    map_host_limit = 0;
+  }
+
+  /* Amount of device memory to keep is free after texture memory
+   * and working memory allocations respectively. We set the working
+   * memory limit headroom lower so that some space is left after all
+   * texture memory allocations. */
+  device_working_headroom = 32 * 1024 * 1024LL;   // 32MB
+  device_texture_headroom = 128 * 1024 * 1024LL;  // 128MB
+
+  VLOG_INFO << "Mapped host memory limit set to " << string_human_readable_number(map_host_limit)
+            << " bytes. (" << string_human_readable_size(map_host_limit) << ")";
+}
+
+void CUDADevice::load_texture_info()
+{
+  if (need_texture_info) {
+    /* Unset flag before copying, so this does not loop indefinitely if the copy below calls
+     * into 'move_textures_to_host' (which calls 'load_texture_info' again). */
+    need_texture_info = false;
+    texture_info.copy_to_device();
+  }
+}
+
+void CUDADevice::move_textures_to_host(size_t size, bool for_texture)
+{
+  /* Break out of recursive call, which can happen when moving memory on a multi device. */
+  static bool any_device_moving_textures_to_host = false;
+  if (any_device_moving_textures_to_host) {
+    return;
+  }
+
+  /* Signal to reallocate textures in host memory only. */
+  move_texture_to_host = true;
+
+  while (size > 0) {
+    /* Find suitable memory allocation to move. */
+    device_memory *max_mem = NULL;
+    size_t max_size = 0;
+    bool max_is_image = false;
+
+    thread_scoped_lock lock(cuda_mem_map_mutex);
+    foreach (CUDAMemMap::value_type &pair, cuda_mem_map) {
+      device_memory &mem = *pair.first;
+      CUDAMem *cmem = &pair.second;
+
+      /* Can only move textures allocated on this device (and not those from peer devices).
+       * And need to ignore memory that is already on the host. */
+      if (!mem.is_resident(this) || cmem->use_mapped_host) {
+        continue;
+      }
+
+      bool is_texture = (mem.type == MEM_TEXTURE || mem.type == MEM_GLOBAL) &&
+                        (&mem != &texture_info);
+      bool is_image = is_texture && (mem.data_height > 1);
+
+      /* Can't move this type of memory. */
+      if (!is_texture || cmem->array) {
+        continue;
+      }
+
+      /* For other textures, only move image textures. */
+      if (for_texture && !is_image) {
+        continue;
+      }
+
+      /* Try to move largest allocation, prefer moving images. */
+      if (is_image > max_is_image || (is_image == max_is_image && mem.device_size > max_size)) {
+        max_is_image = is_image;
+        max_size = mem.device_size;
+        max_mem = &mem;
+      }
+    }
+    lock.unlock();
+
+    /* Move to host memory. This part is mutex protected since
+     * multiple CUDA devices could be moving the memory. The
+     * first one will do it, and the rest will adopt the pointer. */
+    if (max_mem) {
+      VLOG_WORK << "Move memory from device to host: " << max_mem->name;
+
+      static thread_mutex move_mutex;
+      thread_scoped_lock lock(move_mutex);
+
+      any_device_moving_textures_to_host = true;
+
+      /* Potentially need to call back into multi device, so pointer mapping
+       * and peer devices are updated. This is also necessary since the device
+       * pointer may just be a key here, so cannot be accessed and freed directly.
+       * Unfortunately it does mean that memory is reallocated on all other
+       * devices as well, which is potentially dangerous when still in use (since
+       * a thread rendering on another devices would only be caught in this mutex
+       * if it so happens to do an allocation at the same time as well. */
+      max_mem->device_copy_to();
+      size = (max_size >= size) ? 0 : size - max_size;
+
+      any_device_moving_textures_to_host = false;
+    }
+    else {
+      break;
+    }
+  }
+
+  /* Unset flag before texture info is reloaded, since it should stay in device memory. */
+  move_texture_to_host = false;
+
+  /* Update texture info array with new pointers. */
+  load_texture_info();
+}
+
+CUDADevice::CUDAMem *CUDADevice::generic_alloc(device_memory &mem, size_t pitch_padding)
 {
   CUDAContextScope scope(this);
 
+  CUdeviceptr device_pointer = 0;
+  size_t size = mem.memory_size() + pitch_padding;
+
+  CUresult mem_alloc_result = CUDA_ERROR_OUT_OF_MEMORY;
+  const char *status = "";
+
+  /* First try allocating in device memory, respecting headroom. We make
+   * an exception for texture info. It is small and frequently accessed,
+   * so treat it as working memory.
+   *
+   * If there is not enough room for working memory, we will try to move
+   * textures to host memory, assuming the performance impact would have
+   * been worse for working memory. */
+  bool is_texture = (mem.type == MEM_TEXTURE || mem.type == MEM_GLOBAL) && (&mem != &texture_info);
+  bool is_image = is_texture && (mem.data_height > 1);
+
+  size_t headroom = (is_texture) ? device_texture_headroom : device_working_headroom;
+
+  size_t total = 0, free = 0;
   cuMemGetInfo(&free, &total);
+
+  /* Move textures to host memory if needed. */
+  if (!move_texture_to_host && !is_image && (size + headroom) >= free && can_map_host) {
+    move_textures_to_host(size + headroom - free, is_texture);
+    cuMemGetInfo(&free, &total);
+  }
+
+  /* Allocate in device memory. */
+  if (!move_texture_to_host && (size + headroom) < free) {
+    mem_alloc_result = cuMemAlloc(&device_pointer, size);
+    if (mem_alloc_result == CUDA_SUCCESS) {
+      status = " in device memory";
+    }
+  }
+
+  /* Fall back to mapped host memory if needed and possible. */
+
+  void *shared_pointer = 0;
+
+  if (mem_alloc_result != CUDA_SUCCESS && can_map_host && mem.type != MEM_DEVICE_ONLY) {
+    if (mem.shared_pointer) {
+      /* Another device already allocated host memory. */
+      mem_alloc_result = CUDA_SUCCESS;
+      shared_pointer = mem.shared_pointer;
+    }
+    else if (map_host_used + size < map_host_limit) {
+      /* Allocate host memory ourselves. */
+      mem_alloc_result = cuMemHostAlloc(
+          &shared_pointer, size, CU_MEMHOSTALLOC_DEVICEMAP | CU_MEMHOSTALLOC_WRITECOMBINED);
+
+      assert((mem_alloc_result == CUDA_SUCCESS && shared_pointer != 0) ||
+             (mem_alloc_result != CUDA_SUCCESS && shared_pointer == 0));
+    }
+
+    if (mem_alloc_result == CUDA_SUCCESS) {
+      cuda_assert(cuMemHostGetDevicePointer_v2(&device_pointer, shared_pointer, 0));
+      map_host_used += size;
+      status = " in host memory";
+    }
+  }
+
+  if (mem_alloc_result != CUDA_SUCCESS) {
+    if (mem.type == MEM_DEVICE_ONLY) {
+      status = " failed, out of device memory";
+      set_error("System is out of GPU memory");
+    }
+    else {
+      status = " failed, out of device and host memory";
+      set_error("System is out of GPU and shared host memory");
+    }
+  }
+
+  if (mem.name) {
+    VLOG_WORK << "Buffer allocate: " << mem.name << ", "
+              << string_human_readable_number(mem.memory_size()) << " bytes. ("
+              << string_human_readable_size(mem.memory_size()) << ")" << status;
+  }
+
+  mem.device_pointer = (device_ptr)device_pointer;
+  mem.device_size = size;
+  stats.mem_alloc(size);
+
+  if (!mem.device_pointer) {
+    return NULL;
+  }
+
+  /* Insert into map of allocations. */
+  thread_scoped_lock lock(cuda_mem_map_mutex);
+  CUDAMem *cmem = &cuda_mem_map[&mem];
+  if (shared_pointer != 0) {
+    /* Replace host pointer with our host allocation. Only works if
+     * CUDA memory layout is the same and has no pitch padding. Also
+     * does not work if we move textures to host during a render,
+     * since other devices might be using the memory. */
+
+    if (!move_texture_to_host && pitch_padding == 0 && mem.host_pointer &&
+        mem.host_pointer != shared_pointer) {
+      memcpy(shared_pointer, mem.host_pointer, size);
+
+      /* A Call to device_memory::host_free() should be preceded by
+       * a call to device_memory::device_free() for host memory
+       * allocated by a device to be handled properly. Two exceptions
+       * are here and a call in OptiXDevice::generic_alloc(), where
+       * the current host memory can be assumed to be allocated by
+       * device_memory::host_alloc(), not by a device */
+
+      mem.host_free();
+      mem.host_pointer = shared_pointer;
+    }
+    mem.shared_pointer = shared_pointer;
+    mem.shared_counter++;
+    cmem->use_mapped_host = true;
+  }
+  else {
+    cmem->use_mapped_host = false;
+  }
+
+  return cmem;
 }
 
-bool CUDADevice::alloc_device(void *&device_pointer, size_t size)
+void CUDADevice::generic_copy_to(device_memory &mem)
 {
-  CUDAContextScope scope(this);
+  if (!mem.host_pointer || !mem.device_pointer) {
+    return;
+  }
 
-  CUresult mem_alloc_result = cuMemAlloc((CUdeviceptr *)&device_pointer, size);
-  return mem_alloc_result == CUDA_SUCCESS;
+  /* If use_mapped_host of mem is false, the current device only uses device memory allocated by
+   * cuMemAlloc regardless of mem.host_pointer and mem.shared_pointer, and should copy data from
+   * mem.host_pointer. */
+  thread_scoped_lock lock(cuda_mem_map_mutex);
+  if (!cuda_mem_map[&mem].use_mapped_host || mem.host_pointer != mem.shared_pointer) {
+    const CUDAContextScope scope(this);
+    cuda_assert(
+        cuMemcpyHtoD((CUdeviceptr)mem.device_pointer, mem.host_pointer, mem.memory_size()));
+  }
 }
 
-void CUDADevice::free_device(void *device_pointer)
+void CUDADevice::generic_free(device_memory &mem)
 {
-  CUDAContextScope scope(this);
+  if (mem.device_pointer) {
+    CUDAContextScope scope(this);
+    thread_scoped_lock lock(cuda_mem_map_mutex);
+    DCHECK(cuda_mem_map.find(&mem) != cuda_mem_map.end());
+    const CUDAMem &cmem = cuda_mem_map[&mem];
 
-  cuda_assert(cuMemFree((CUdeviceptr)device_pointer));
-}
+    /* If cmem.use_mapped_host is true, reference counting is used
+     * to safely free a mapped host memory. */
 
-bool CUDADevice::alloc_host(void *&shared_pointer, size_t size)
-{
-  CUDAContextScope scope(this);
+    if (cmem.use_mapped_host) {
+      assert(mem.shared_pointer);
+      if (mem.shared_pointer) {
+        assert(mem.shared_counter > 0);
+        if (--mem.shared_counter == 0) {
+          if (mem.host_pointer == mem.shared_pointer) {
+            mem.host_pointer = 0;
+          }
+          cuMemFreeHost(mem.shared_pointer);
+          mem.shared_pointer = 0;
+        }
+      }
+      map_host_used -= mem.device_size;
+    }
+    else {
+      /* Free device memory. */
+      cuda_assert(cuMemFree(mem.device_pointer));
+    }
 
-  CUresult mem_alloc_result = cuMemHostAlloc(
-      &shared_pointer, size, CU_MEMHOSTALLOC_DEVICEMAP | CU_MEMHOSTALLOC_WRITECOMBINED);
-  return mem_alloc_result == CUDA_SUCCESS;
-}
+    stats.mem_free(mem.device_size);
+    mem.device_pointer = 0;
+    mem.device_size = 0;
 
-void CUDADevice::free_host(void *shared_pointer)
-{
-  CUDAContextScope scope(this);
-
-  cuMemFreeHost(shared_pointer);
-}
-
-bool CUDADevice::transform_host_pointer(void *&device_pointer, void *&shared_pointer)
-{
-  CUDAContextScope scope(this);
-
-  cuda_assert(cuMemHostGetDevicePointer_v2((CUdeviceptr *)&device_pointer, shared_pointer, 0));
-  return true;
-}
-
-void CUDADevice::copy_host_to_device(void *device_pointer, void *host_pointer, size_t size)
-{
-  const CUDAContextScope scope(this);
-
-  cuda_assert(cuMemcpyHtoD((CUdeviceptr)device_pointer, host_pointer, size));
+    cuda_mem_map.erase(cuda_mem_map.find(&mem));
+  }
 }
 
 void CUDADevice::mem_alloc(device_memory &mem)
@@ -613,8 +868,8 @@ void CUDADevice::mem_zero(device_memory &mem)
 
   /* If use_mapped_host of mem is false, mem.device_pointer currently refers to device memory
    * regardless of mem.host_pointer and mem.shared_pointer. */
-  thread_scoped_lock lock(device_mem_map_mutex);
-  if (!device_mem_map[&mem].use_mapped_host || mem.host_pointer != mem.shared_pointer) {
+  thread_scoped_lock lock(cuda_mem_map_mutex);
+  if (!cuda_mem_map[&mem].use_mapped_host || mem.host_pointer != mem.shared_pointer) {
     const CUDAContextScope scope(this);
     cuda_assert(cuMemsetD8((CUdeviceptr)mem.device_pointer, 0, mem.memory_size()));
   }
@@ -739,19 +994,19 @@ void CUDADevice::tex_alloc(device_texture &mem)
       return;
   }
 
-  Mem *cmem = NULL;
+  CUDAMem *cmem = NULL;
   CUarray array_3d = NULL;
   size_t src_pitch = mem.data_width * dsize * mem.data_elements;
   size_t dst_pitch = src_pitch;
 
   if (!mem.is_resident(this)) {
-    thread_scoped_lock lock(device_mem_map_mutex);
-    cmem = &device_mem_map[&mem];
+    thread_scoped_lock lock(cuda_mem_map_mutex);
+    cmem = &cuda_mem_map[&mem];
     cmem->texobject = 0;
 
     if (mem.data_depth > 1) {
       array_3d = (CUarray)mem.device_pointer;
-      cmem->array = reinterpret_cast<arrayMemObject>(array_3d);
+      cmem->array = array_3d;
     }
     else if (mem.data_height > 0) {
       dst_pitch = align_up(src_pitch, pitch_alignment);
@@ -795,10 +1050,10 @@ void CUDADevice::tex_alloc(device_texture &mem)
     mem.device_size = size;
     stats.mem_alloc(size);
 
-    thread_scoped_lock lock(device_mem_map_mutex);
-    cmem = &device_mem_map[&mem];
+    thread_scoped_lock lock(cuda_mem_map_mutex);
+    cmem = &cuda_mem_map[&mem];
     cmem->texobject = 0;
-    cmem->array = reinterpret_cast<arrayMemObject>(array_3d);
+    cmem->array = array_3d;
   }
   else if (mem.data_height > 0) {
     /* 2D texture, using pitch aligned linear memory. */
@@ -882,8 +1137,8 @@ void CUDADevice::tex_alloc(device_texture &mem)
     texDesc.filterMode = filter_mode;
     texDesc.flags = CU_TRSF_NORMALIZED_COORDINATES;
 
-    thread_scoped_lock lock(device_mem_map_mutex);
-    cmem = &device_mem_map[&mem];
+    thread_scoped_lock lock(cuda_mem_map_mutex);
+    cmem = &cuda_mem_map[&mem];
 
     cuda_assert(cuTexObjectCreate(&cmem->texobject, &resDesc, &texDesc, NULL));
 
@@ -898,9 +1153,9 @@ void CUDADevice::tex_free(device_texture &mem)
 {
   if (mem.device_pointer) {
     CUDAContextScope scope(this);
-    thread_scoped_lock lock(device_mem_map_mutex);
-    DCHECK(device_mem_map.find(&mem) != device_mem_map.end());
-    const Mem &cmem = device_mem_map[&mem];
+    thread_scoped_lock lock(cuda_mem_map_mutex);
+    DCHECK(cuda_mem_map.find(&mem) != cuda_mem_map.end());
+    const CUDAMem &cmem = cuda_mem_map[&mem];
 
     if (cmem.texobject) {
       /* Free bindless texture. */
@@ -909,16 +1164,16 @@ void CUDADevice::tex_free(device_texture &mem)
 
     if (!mem.is_resident(this)) {
       /* Do not free memory here, since it was allocated on a different device. */
-      device_mem_map.erase(device_mem_map.find(&mem));
+      cuda_mem_map.erase(cuda_mem_map.find(&mem));
     }
     else if (cmem.array) {
       /* Free array. */
-      cuArrayDestroy(reinterpret_cast<CUarray>(cmem.array));
+      cuArrayDestroy(cmem.array);
       stats.mem_free(mem.device_size);
       mem.device_pointer = 0;
       mem.device_size = 0;
 
-      device_mem_map.erase(device_mem_map.find(&mem));
+      cuda_mem_map.erase(cuda_mem_map.find(&mem));
     }
     else {
       lock.unlock();

intern/cycles/device/cuda/device_impl.h

@@ -21,7 +21,7 @@ CCL_NAMESPACE_BEGIN
 
 class DeviceQueue;
 
-class CUDADevice : public GPUDevice {
+class CUDADevice : public Device {
 
   friend class CUDAContextScope;
 
@@ -29,11 +29,36 @@ class CUDADevice : public GPUDevice {
   CUdevice cuDevice;
   CUcontext cuContext;
   CUmodule cuModule;
+  size_t device_texture_headroom;
+  size_t device_working_headroom;
+  bool move_texture_to_host;
+  size_t map_host_used;
+  size_t map_host_limit;
+  int can_map_host;
   int pitch_alignment;
   int cuDevId;
   int cuDevArchitecture;
   bool first_error;
 
+  struct CUDAMem {
+    CUDAMem() : texobject(0), array(0), use_mapped_host(false)
+    {
+    }
+
+    CUtexObject texobject;
+    CUarray array;
+
+    /* If true, a mapped host memory in shared_pointer is being used. */
+    bool use_mapped_host;
+  };
+  typedef map<device_memory *, CUDAMem> CUDAMemMap;
+  CUDAMemMap cuda_mem_map;
+  thread_mutex cuda_mem_map_mutex;
+
+  /* Bindless Textures */
+  device_vector<TextureInfo> texture_info;
+  bool need_texture_info;
+
   CUDADeviceKernels kernels;
 
   static bool have_precompiled_kernels();
@@ -63,13 +88,17 @@ class CUDADevice : public GPUDevice {
 
   void reserve_local_memory(const uint kernel_features);
 
-  virtual void get_device_memory_info(size_t &total, size_t &free) override;
-  virtual bool alloc_device(void *&device_pointer, size_t size) override;
-  virtual void free_device(void *device_pointer) override;
-  virtual bool alloc_host(void *&shared_pointer, size_t size) override;
-  virtual void free_host(void *shared_pointer) override;
-  virtual bool transform_host_pointer(void *&device_pointer, void *&shared_pointer) override;
-  virtual void copy_host_to_device(void *device_pointer, void *host_pointer, size_t size) override;
+  void init_host_memory();
+
+  void load_texture_info();
+
+  void move_textures_to_host(size_t size, bool for_texture);
+
+  CUDAMem *generic_alloc(device_memory &mem, size_t pitch_padding = 0);
+
+  void generic_copy_to(device_memory &mem);
+
+  void generic_free(device_memory &mem);
 
   void mem_alloc(device_memory &mem) override;
 

intern/cycles/device/device.cpp

@@ -452,320 +452,6 @@ void *Device::get_cpu_osl_memory()
   return nullptr;
 }
 
-GPUDevice::~GPUDevice() noexcept(false)
-{
-}
-
-bool GPUDevice::load_texture_info()
-{
-  if (need_texture_info) {
-    /* Unset flag before copying, so this does not loop indefinitely if the copy below calls
-     * into 'move_textures_to_host' (which calls 'load_texture_info' again). */
-    need_texture_info = false;
-    texture_info.copy_to_device();
-    return true;
-  }
-  else {
-    return false;
-  }
-}
-
-void GPUDevice::init_host_memory(size_t preferred_texture_headroom,
-                                 size_t preferred_working_headroom)
-{
-  /* Limit amount of host mapped memory, because allocating too much can
-   * cause system instability. Leave at least half or 4 GB of system
-   * memory free, whichever is smaller. */
-  size_t default_limit = 4 * 1024 * 1024 * 1024LL;
-  size_t system_ram = system_physical_ram();
-
-  if (system_ram > 0) {
-    if (system_ram / 2 > default_limit) {
-      map_host_limit = system_ram - default_limit;
-    }
-    else {
-      map_host_limit = system_ram / 2;
-    }
-  }
-  else {
-    VLOG_WARNING << "Mapped host memory disabled, failed to get system RAM";
-    map_host_limit = 0;
-  }
-
-  /* Amount of device memory to keep free after texture memory
-   * and working memory allocations respectively. We set the working
-   * memory limit headroom lower than the working one so there
-   * is space left for it. */
-  device_working_headroom = preferred_working_headroom > 0 ? preferred_working_headroom :
-                                                             32 * 1024 * 1024LL;  // 32MB
-  device_texture_headroom = preferred_texture_headroom > 0 ? preferred_texture_headroom :
-                                                             128 * 1024 * 1024LL;  // 128MB
-
-  VLOG_INFO << "Mapped host memory limit set to " << string_human_readable_number(map_host_limit)
-            << " bytes. (" << string_human_readable_size(map_host_limit) << ")";
-}
-
-void GPUDevice::move_textures_to_host(size_t size, bool for_texture)
-{
-  /* Break out of recursive call, which can happen when moving memory on a multi device. */
-  static bool any_device_moving_textures_to_host = false;
-  if (any_device_moving_textures_to_host) {
-    return;
-  }
-
-  /* Signal to reallocate textures in host memory only. */
-  move_texture_to_host = true;
-
-  while (size > 0) {
-    /* Find suitable memory allocation to move. */
-    device_memory *max_mem = NULL;
-    size_t max_size = 0;
-    bool max_is_image = false;
-
-    thread_scoped_lock lock(device_mem_map_mutex);
-    foreach (MemMap::value_type &pair, device_mem_map) {
-      device_memory &mem = *pair.first;
-      Mem *cmem = &pair.second;
-
-      /* Can only move textures allocated on this device (and not those from peer devices).
-       * And need to ignore memory that is already on the host. */
-      if (!mem.is_resident(this) || cmem->use_mapped_host) {
-        continue;
-      }
-
-      bool is_texture = (mem.type == MEM_TEXTURE || mem.type == MEM_GLOBAL) &&
-                        (&mem != &texture_info);
-      bool is_image = is_texture && (mem.data_height > 1);
-
-      /* Can't move this type of memory. */
-      if (!is_texture || cmem->array) {
-        continue;
-      }
-
-      /* For other textures, only move image textures. */
-      if (for_texture && !is_image) {
-        continue;
-      }
-
-      /* Try to move largest allocation, prefer moving images. */
-      if (is_image > max_is_image || (is_image == max_is_image && mem.device_size > max_size)) {
-        max_is_image = is_image;
-        max_size = mem.device_size;
-        max_mem = &mem;
-      }
-    }
-    lock.unlock();
-
-    /* Move to host memory. This part is mutex protected since
-     * multiple backend devices could be moving the memory. The
-     * first one will do it, and the rest will adopt the pointer. */
-    if (max_mem) {
-      VLOG_WORK << "Move memory from device to host: " << max_mem->name;
-
-      static thread_mutex move_mutex;
-      thread_scoped_lock lock(move_mutex);
-
-      any_device_moving_textures_to_host = true;
-
-      /* Potentially need to call back into multi device, so pointer mapping
-       * and peer devices are updated. This is also necessary since the device
-       * pointer may just be a key here, so cannot be accessed and freed directly.
-       * Unfortunately it does mean that memory is reallocated on all other
-       * devices as well, which is potentially dangerous when still in use (since
-       * a thread rendering on another devices would only be caught in this mutex
-       * if it so happens to do an allocation at the same time as well. */
-      max_mem->device_copy_to();
-      size = (max_size >= size) ? 0 : size - max_size;
-
-      any_device_moving_textures_to_host = false;
-    }
-    else {
-      break;
-    }
-  }
-
-  /* Unset flag before texture info is reloaded, since it should stay in device memory. */
-  move_texture_to_host = false;
-
-  /* Update texture info array with new pointers. */
-  load_texture_info();
-}
-
-GPUDevice::Mem *GPUDevice::generic_alloc(device_memory &mem, size_t pitch_padding)
-{
-  void *device_pointer = 0;
-  size_t size = mem.memory_size() + pitch_padding;
-
-  bool mem_alloc_result = false;
-  const char *status = "";
-
-  /* First try allocating in device memory, respecting headroom. We make
-   * an exception for texture info. It is small and frequently accessed,
-   * so treat it as working memory.
-   *
-   * If there is not enough room for working memory, we will try to move
-   * textures to host memory, assuming the performance impact would have
-   * been worse for working memory. */
-  bool is_texture = (mem.type == MEM_TEXTURE || mem.type == MEM_GLOBAL) && (&mem != &texture_info);
-  bool is_image = is_texture && (mem.data_height > 1);
-
-  size_t headroom = (is_texture) ? device_texture_headroom : device_working_headroom;
-
-  size_t total = 0, free = 0;
-  get_device_memory_info(total, free);
-
-  /* Move textures to host memory if needed. */
-  if (!move_texture_to_host && !is_image && (size + headroom) >= free && can_map_host) {
-    move_textures_to_host(size + headroom - free, is_texture);
-    get_device_memory_info(total, free);
-  }
-
-  /* Allocate in device memory. */
-  if (!move_texture_to_host && (size + headroom) < free) {
-    mem_alloc_result = alloc_device(device_pointer, size);
-    if (mem_alloc_result) {
-      device_mem_in_use += size;
-      status = " in device memory";
-    }
-  }
-
-  /* Fall back to mapped host memory if needed and possible. */
-
-  void *shared_pointer = 0;
-
-  if (!mem_alloc_result && can_map_host && mem.type != MEM_DEVICE_ONLY) {
-    if (mem.shared_pointer) {
-      /* Another device already allocated host memory. */
-      mem_alloc_result = true;
-      shared_pointer = mem.shared_pointer;
-    }
-    else if (map_host_used + size < map_host_limit) {
-      /* Allocate host memory ourselves. */
-      mem_alloc_result = alloc_host(shared_pointer, size);
-
-      assert((mem_alloc_result && shared_pointer != 0) ||
-             (!mem_alloc_result && shared_pointer == 0));
-    }
-
-    if (mem_alloc_result) {
-      assert(transform_host_pointer(device_pointer, shared_pointer));
-      map_host_used += size;
-      status = " in host memory";
-    }
-  }
-
-  if (!mem_alloc_result) {
-    if (mem.type == MEM_DEVICE_ONLY) {
-      status = " failed, out of device memory";
-      set_error("System is out of GPU memory");
-    }
-    else {
-      status = " failed, out of device and host memory";
-      set_error("System is out of GPU and shared host memory");
-    }
-  }
-
-  if (mem.name) {
-    VLOG_WORK << "Buffer allocate: " << mem.name << ", "
-              << string_human_readable_number(mem.memory_size()) << " bytes. ("
-              << string_human_readable_size(mem.memory_size()) << ")" << status;
-  }
-
-  mem.device_pointer = (device_ptr)device_pointer;
-  mem.device_size = size;
-  stats.mem_alloc(size);
-
-  if (!mem.device_pointer) {
-    return NULL;
-  }
-
-  /* Insert into map of allocations. */
-  thread_scoped_lock lock(device_mem_map_mutex);
-  Mem *cmem = &device_mem_map[&mem];
-  if (shared_pointer != 0) {
-    /* Replace host pointer with our host allocation. Only works if
-     * memory layout is the same and has no pitch padding. Also
-     * does not work if we move textures to host during a render,
-     * since other devices might be using the memory. */
-
-    if (!move_texture_to_host && pitch_padding == 0 && mem.host_pointer &&
-        mem.host_pointer != shared_pointer) {
-      memcpy(shared_pointer, mem.host_pointer, size);
-
-      /* A Call to device_memory::host_free() should be preceded by
-       * a call to device_memory::device_free() for host memory
-       * allocated by a device to be handled properly. Two exceptions
-       * are here and a call in OptiXDevice::generic_alloc(), where
-       * the current host memory can be assumed to be allocated by
-       * device_memory::host_alloc(), not by a device */
-
-      mem.host_free();
-      mem.host_pointer = shared_pointer;
-    }
-    mem.shared_pointer = shared_pointer;
-    mem.shared_counter++;
-    cmem->use_mapped_host = true;
-  }
-  else {
-    cmem->use_mapped_host = false;
-  }
-
-  return cmem;
-}
-
-void GPUDevice::generic_free(device_memory &mem)
-{
-  if (mem.device_pointer) {
-    thread_scoped_lock lock(device_mem_map_mutex);
-    DCHECK(device_mem_map.find(&mem) != device_mem_map.end());
-    const Mem &cmem = device_mem_map[&mem];
-
-    /* If cmem.use_mapped_host is true, reference counting is used
-     * to safely free a mapped host memory. */
-
-    if (cmem.use_mapped_host) {
-      assert(mem.shared_pointer);
-      if (mem.shared_pointer) {
-        assert(mem.shared_counter > 0);
-        if (--mem.shared_counter == 0) {
-          if (mem.host_pointer == mem.shared_pointer) {
-            mem.host_pointer = 0;
-          }
-          free_host(mem.shared_pointer);
-          mem.shared_pointer = 0;
-        }
-      }
-      map_host_used -= mem.device_size;
-    }
-    else {
-      /* Free device memory. */
-      free_device((void *)mem.device_pointer);
-      device_mem_in_use -= mem.device_size;
-    }
-
-    stats.mem_free(mem.device_size);
-    mem.device_pointer = 0;
-    mem.device_size = 0;
-
-    device_mem_map.erase(device_mem_map.find(&mem));
-  }
-}
-
-void GPUDevice::generic_copy_to(device_memory &mem)
-{
-  if (!mem.host_pointer || !mem.device_pointer) {
-    return;
-  }
-
-  /* If use_mapped_host of mem is false, the current device only uses device memory allocated by
-   * backend device allocation regardless of mem.host_pointer and mem.shared_pointer, and should
-   * copy data from mem.host_pointer. */
-  thread_scoped_lock lock(device_mem_map_mutex);
-  if (!device_mem_map[&mem].use_mapped_host || mem.host_pointer != mem.shared_pointer) {
-    copy_host_to_device((void *)mem.device_pointer, mem.host_pointer, mem.memory_size());
-  }
-}
-
 /* DeviceInfo */
 
 CCL_NAMESPACE_END

intern/cycles/device/device.h

@@ -309,93 +309,6 @@ class Device {
   static uint devices_initialized_mask;
 };
 
-/* Device, which is GPU, with some common functionality for GPU backends */
-class GPUDevice : public Device {
- protected:
-  GPUDevice(const DeviceInfo &info_, Stats &stats_, Profiler &profiler_)
-      : Device(info_, stats_, profiler_),
-        texture_info(this, "texture_info", MEM_GLOBAL),
-        need_texture_info(false),
-        can_map_host(false),
-        map_host_used(0),
-        map_host_limit(0),
-        device_texture_headroom(0),
-        device_working_headroom(0),
-        device_mem_map(),
-        device_mem_map_mutex(),
-        move_texture_to_host(false),
-        device_mem_in_use(0)
-  {
-  }
-
- public:
-  virtual ~GPUDevice() noexcept(false);
-
-  /* For GPUs that can use bindless textures in some way or another. */
-  device_vector<TextureInfo> texture_info;
-  bool need_texture_info;
-  /* Returns true if the texture info was copied to the device (meaning, some more
-   * re-initialization might be needed). */
-  virtual bool load_texture_info();
-
- protected:
-  /* Memory allocation, only accessed through device_memory. */
-  friend class device_memory;
-
-  bool can_map_host;
-  size_t map_host_used;
-  size_t map_host_limit;
-  size_t device_texture_headroom;
-  size_t device_working_headroom;
-  typedef unsigned long long texMemObject;
-  typedef unsigned long long arrayMemObject;
-  struct Mem {
-    Mem() : texobject(0), array(0), use_mapped_host(false)
-    {
-    }
-
-    texMemObject texobject;
-    arrayMemObject array;
-
-    /* If true, a mapped host memory in shared_pointer is being used. */
-    bool use_mapped_host;
-  };
-  typedef map<device_memory *, Mem> MemMap;
-  MemMap device_mem_map;
-  thread_mutex device_mem_map_mutex;
-  bool move_texture_to_host;
-  /* Simple counter which will try to track amount of used device memory */
-  size_t device_mem_in_use;
-
-  virtual void init_host_memory(size_t preferred_texture_headroom = 0,
-                                size_t preferred_working_headroom = 0);
-  virtual void move_textures_to_host(size_t size, bool for_texture);
-
-  /* Allocation, deallocation and copy functions, with corresponding
-   * support of device/host allocations. */
-  virtual GPUDevice::Mem *generic_alloc(device_memory &mem, size_t pitch_padding = 0);
-  virtual void generic_free(device_memory &mem);
-  virtual void generic_copy_to(device_memory &mem);
-
-  /* total - amount of device memory, free - amount of available device memory */
-  virtual void get_device_memory_info(size_t &total, size_t &free) = 0;
-
-  virtual bool alloc_device(void *&device_pointer, size_t size) = 0;
-
-  virtual void free_device(void *device_pointer) = 0;
-
-  virtual bool alloc_host(void *&shared_pointer, size_t size) = 0;
-
-  virtual void free_host(void *shared_pointer) = 0;
-
-  /* This function should return device pointer corresponding to shared pointer, which
-   * is host buffer, allocated in `alloc_host`. The function should `true`, if such
-   * address transformation is possible and `false` otherwise. */
-  virtual bool transform_host_pointer(void *&device_pointer, void *&shared_pointer) = 0;
-
-  virtual void copy_host_to_device(void *device_pointer, void *host_pointer, size_t size) = 0;
-};
-
 CCL_NAMESPACE_END
 
 #endif /* __DEVICE_H__ */

intern/cycles/device/hip/device_impl.cpp

@@ -53,12 +53,8 @@ void HIPDevice::set_error(const string &error)
 }
 
 HIPDevice::HIPDevice(const DeviceInfo &info, Stats &stats, Profiler &profiler)
-    : GPUDevice(info, stats, profiler)
+    : Device(info, stats, profiler), texture_info(this, "texture_info", MEM_GLOBAL)
 {
-  /* Verify that base class types can be used with specific backend types */
-  static_assert(sizeof(texMemObject) == sizeof(hipTextureObject_t));
-  static_assert(sizeof(arrayMemObject) == sizeof(hArray));
-
   first_error = true;
 
   hipDevId = info.num;
@@ -69,6 +65,12 @@ HIPDevice::HIPDevice(const DeviceInfo &info, Stats &stats, Profiler &profiler)
 
   need_texture_info = false;
 
+  device_texture_headroom = 0;
+  device_working_headroom = 0;
+  move_texture_to_host = false;
+  map_host_limit = 0;
+  map_host_used = 0;
+  can_map_host = 0;
   pitch_alignment = 0;
 
   /* Initialize HIP. */
@@ -89,9 +91,7 @@ HIPDevice::HIPDevice(const DeviceInfo &info, Stats &stats, Profiler &profiler)
   /* hipDeviceMapHost for mapping host memory when out of device memory.
    * hipDeviceLmemResizeToMax for reserving local memory ahead of render,
    * so we can predict which memory to map to host. */
-  int value;
-  hip_assert(hipDeviceGetAttribute(&value, hipDeviceAttributeCanMapHostMemory, hipDevice));
-  can_map_host = value != 0;
+  hip_assert(hipDeviceGetAttribute(&can_map_host, hipDeviceAttributeCanMapHostMemory, hipDevice));
 
   hip_assert(
       hipDeviceGetAttribute(&pitch_alignment, hipDeviceAttributeTexturePitchAlignment, hipDevice));
@@ -460,58 +460,305 @@ void HIPDevice::reserve_local_memory(const uint kernel_features)
 #  endif
 }
 
-void HIPDevice::get_device_memory_info(size_t &total, size_t &free)
+void HIPDevice::init_host_memory()
+{
+  /* Limit amount of host mapped memory, because allocating too much can
+   * cause system instability. Leave at least half or 4 GB of system
+   * memory free, whichever is smaller. */
+  size_t default_limit = 4 * 1024 * 1024 * 1024LL;
+  size_t system_ram = system_physical_ram();
+
+  if (system_ram > 0) {
+    if (system_ram / 2 > default_limit) {
+      map_host_limit = system_ram - default_limit;
+    }
+    else {
+      map_host_limit = system_ram / 2;
+    }
+  }
+  else {
+    VLOG_WARNING << "Mapped host memory disabled, failed to get system RAM";
+    map_host_limit = 0;
+  }
+
+  /* Amount of device memory to keep is free after texture memory
+   * and working memory allocations respectively. We set the working
+   * memory limit headroom lower so that some space is left after all
+   * texture memory allocations. */
+  device_working_headroom = 32 * 1024 * 1024LL;   // 32MB
+  device_texture_headroom = 128 * 1024 * 1024LL;  // 128MB
+
+  VLOG_INFO << "Mapped host memory limit set to " << string_human_readable_number(map_host_limit)
+            << " bytes. (" << string_human_readable_size(map_host_limit) << ")";
+}
+
+void HIPDevice::load_texture_info()
+{
+  if (need_texture_info) {
+    /* Unset flag before copying, so this does not loop indefinitely if the copy below calls
+     * into 'move_textures_to_host' (which calls 'load_texture_info' again). */
+    need_texture_info = false;
+    texture_info.copy_to_device();
+  }
+}
+
+void HIPDevice::move_textures_to_host(size_t size, bool for_texture)
+{
+  /* Break out of recursive call, which can happen when moving memory on a multi device. */
+  static bool any_device_moving_textures_to_host = false;
+  if (any_device_moving_textures_to_host) {
+    return;
+  }
+
+  /* Signal to reallocate textures in host memory only. */
+  move_texture_to_host = true;
+
+  while (size > 0) {
+    /* Find suitable memory allocation to move. */
+    device_memory *max_mem = NULL;
+    size_t max_size = 0;
+    bool max_is_image = false;
+
+    thread_scoped_lock lock(hip_mem_map_mutex);
+    foreach (HIPMemMap::value_type &pair, hip_mem_map) {
+      device_memory &mem = *pair.first;
+      HIPMem *cmem = &pair.second;
+
+      /* Can only move textures allocated on this device (and not those from peer devices).
+       * And need to ignore memory that is already on the host. */
+      if (!mem.is_resident(this) || cmem->use_mapped_host) {
+        continue;
+      }
+
+      bool is_texture = (mem.type == MEM_TEXTURE || mem.type == MEM_GLOBAL) &&
+                        (&mem != &texture_info);
+      bool is_image = is_texture && (mem.data_height > 1);
+
+      /* Can't move this type of memory. */
+      if (!is_texture || cmem->array) {
+        continue;
+      }
+
+      /* For other textures, only move image textures. */
+      if (for_texture && !is_image) {
+        continue;
+      }
+
+      /* Try to move largest allocation, prefer moving images. */
+      if (is_image > max_is_image || (is_image == max_is_image && mem.device_size > max_size)) {
+        max_is_image = is_image;
+        max_size = mem.device_size;
+        max_mem = &mem;
+      }
+    }
+    lock.unlock();
+
+    /* Move to host memory. This part is mutex protected since
+     * multiple HIP devices could be moving the memory. The
+     * first one will do it, and the rest will adopt the pointer. */
+    if (max_mem) {
+      VLOG_WORK << "Move memory from device to host: " << max_mem->name;
+
+      static thread_mutex move_mutex;
+      thread_scoped_lock lock(move_mutex);
+
+      any_device_moving_textures_to_host = true;
+
+      /* Potentially need to call back into multi device, so pointer mapping
+       * and peer devices are updated. This is also necessary since the device
+       * pointer may just be a key here, so cannot be accessed and freed directly.
+       * Unfortunately it does mean that memory is reallocated on all other
+       * devices as well, which is potentially dangerous when still in use (since
+       * a thread rendering on another devices would only be caught in this mutex
+       * if it so happens to do an allocation at the same time as well. */
+      max_mem->device_copy_to();
+      size = (max_size >= size) ? 0 : size - max_size;
+
+      any_device_moving_textures_to_host = false;
+    }
+    else {
+      break;
+    }
+  }
+
+  /* Unset flag before texture info is reloaded, since it should stay in device memory. */
+  move_texture_to_host = false;
+
+  /* Update texture info array with new pointers. */
+  load_texture_info();
+}
+
+HIPDevice::HIPMem *HIPDevice::generic_alloc(device_memory &mem, size_t pitch_padding)
 {
   HIPContextScope scope(this);
 
+  hipDeviceptr_t device_pointer = 0;
+  size_t size = mem.memory_size() + pitch_padding;
+
+  hipError_t mem_alloc_result = hipErrorOutOfMemory;
+  const char *status = "";
+
+  /* First try allocating in device memory, respecting headroom. We make
+   * an exception for texture info. It is small and frequently accessed,
+   * so treat it as working memory.
+   *
+   * If there is not enough room for working memory, we will try to move
+   * textures to host memory, assuming the performance impact would have
+   * been worse for working memory. */
+  bool is_texture = (mem.type == MEM_TEXTURE || mem.type == MEM_GLOBAL) && (&mem != &texture_info);
+  bool is_image = is_texture && (mem.data_height > 1);
+
+  size_t headroom = (is_texture) ? device_texture_headroom : device_working_headroom;
+
+  size_t total = 0, free = 0;
   hipMemGetInfo(&free, &total);
+
+  /* Move textures to host memory if needed. */
+  if (!move_texture_to_host && !is_image && (size + headroom) >= free && can_map_host) {
+    move_textures_to_host(size + headroom - free, is_texture);
+    hipMemGetInfo(&free, &total);
+  }
+
+  /* Allocate in device memory. */
+  if (!move_texture_to_host && (size + headroom) < free) {
+    mem_alloc_result = hipMalloc(&device_pointer, size);
+    if (mem_alloc_result == hipSuccess) {
+      status = " in device memory";
+    }
+  }
+
+  /* Fall back to mapped host memory if needed and possible. */
+
+  void *shared_pointer = 0;
+
+  if (mem_alloc_result != hipSuccess && can_map_host) {
+    if (mem.shared_pointer) {
+      /* Another device already allocated host memory. */
+      mem_alloc_result = hipSuccess;
+      shared_pointer = mem.shared_pointer;
+    }
+    else if (map_host_used + size < map_host_limit) {
+      /* Allocate host memory ourselves. */
+      mem_alloc_result = hipHostMalloc(
+          &shared_pointer, size, hipHostMallocMapped | hipHostMallocWriteCombined);
+
+      assert((mem_alloc_result == hipSuccess && shared_pointer != 0) ||
+             (mem_alloc_result != hipSuccess && shared_pointer == 0));
+    }
+
+    if (mem_alloc_result == hipSuccess) {
+      hip_assert(hipHostGetDevicePointer(&device_pointer, shared_pointer, 0));
+      map_host_used += size;
+      status = " in host memory";
+    }
+  }
+
+  if (mem_alloc_result != hipSuccess) {
+    status = " failed, out of device and host memory";
+    set_error("System is out of GPU and shared host memory");
+  }
+
+  if (mem.name) {
+    VLOG_WORK << "Buffer allocate: " << mem.name << ", "
+              << string_human_readable_number(mem.memory_size()) << " bytes. ("
+              << string_human_readable_size(mem.memory_size()) << ")" << status;
+  }
+
+  mem.device_pointer = (device_ptr)device_pointer;
+  mem.device_size = size;
+  stats.mem_alloc(size);
+
+  if (!mem.device_pointer) {
+    return NULL;
+  }
+
+  /* Insert into map of allocations. */
+  thread_scoped_lock lock(hip_mem_map_mutex);
+  HIPMem *cmem = &hip_mem_map[&mem];
+  if (shared_pointer != 0) {
+    /* Replace host pointer with our host allocation. Only works if
+     * HIP memory layout is the same and has no pitch padding. Also
+     * does not work if we move textures to host during a render,
+     * since other devices might be using the memory. */
+
+    if (!move_texture_to_host && pitch_padding == 0 && mem.host_pointer &&
+        mem.host_pointer != shared_pointer) {
+      memcpy(shared_pointer, mem.host_pointer, size);
+
+      /* A Call to device_memory::host_free() should be preceded by
+       * a call to device_memory::device_free() for host memory
+       * allocated by a device to be handled properly. Two exceptions
+       * are here and a call in OptiXDevice::generic_alloc(), where
+       * the current host memory can be assumed to be allocated by
+       * device_memory::host_alloc(), not by a device */
+
+      mem.host_free();
+      mem.host_pointer = shared_pointer;
+    }
+    mem.shared_pointer = shared_pointer;
+    mem.shared_counter++;
+    cmem->use_mapped_host = true;
+  }
+  else {
+    cmem->use_mapped_host = false;
+  }
+
+  return cmem;
 }
 
-bool HIPDevice::alloc_device(void *&device_pointer, size_t size)
+void HIPDevice::generic_copy_to(device_memory &mem)
 {
-  HIPContextScope scope(this);
+  if (!mem.host_pointer || !mem.device_pointer) {
+    return;
+  }
 
-  hipError_t mem_alloc_result = hipMalloc((hipDeviceptr_t *)&device_pointer, size);
-  return mem_alloc_result == hipSuccess;
+  /* If use_mapped_host of mem is false, the current device only uses device memory allocated by
+   * hipMalloc regardless of mem.host_pointer and mem.shared_pointer, and should copy data from
+   * mem.host_pointer. */
+  thread_scoped_lock lock(hip_mem_map_mutex);
+  if (!hip_mem_map[&mem].use_mapped_host || mem.host_pointer != mem.shared_pointer) {
+    const HIPContextScope scope(this);
+    hip_assert(
+        hipMemcpyHtoD((hipDeviceptr_t)mem.device_pointer, mem.host_pointer, mem.memory_size()));
+  }
 }
 
-void HIPDevice::free_device(void *device_pointer)
+void HIPDevice::generic_free(device_memory &mem)
 {
-  HIPContextScope scope(this);
+  if (mem.device_pointer) {
+    HIPContextScope scope(this);
+    thread_scoped_lock lock(hip_mem_map_mutex);
+    DCHECK(hip_mem_map.find(&mem) != hip_mem_map.end());
+    const HIPMem &cmem = hip_mem_map[&mem];
 
-  hip_assert(hipFree((hipDeviceptr_t)device_pointer));
-}
+    /* If cmem.use_mapped_host is true, reference counting is used
+     * to safely free a mapped host memory. */
 
-bool HIPDevice::alloc_host(void *&shared_pointer, size_t size)
-{
-  HIPContextScope scope(this);
+    if (cmem.use_mapped_host) {
+      assert(mem.shared_pointer);
+      if (mem.shared_pointer) {
+        assert(mem.shared_counter > 0);
+        if (--mem.shared_counter == 0) {
+          if (mem.host_pointer == mem.shared_pointer) {
+            mem.host_pointer = 0;
+          }
+          hipHostFree(mem.shared_pointer);
+          mem.shared_pointer = 0;
+        }
+      }
+      map_host_used -= mem.device_size;
+    }
+    else {
+      /* Free device memory. */
+      hip_assert(hipFree(mem.device_pointer));
+    }
 
-  hipError_t mem_alloc_result = hipHostMalloc(
-      &shared_pointer, size, hipHostMallocMapped | hipHostMallocWriteCombined);
+    stats.mem_free(mem.device_size);
+    mem.device_pointer = 0;
+    mem.device_size = 0;
 
-  return mem_alloc_result == hipSuccess;
-}
-
-void HIPDevice::free_host(void *shared_pointer)
-{
-  HIPContextScope scope(this);
-
-  hipHostFree(shared_pointer);
-}
-
-bool HIPDevice::transform_host_pointer(void *&device_pointer, void *&shared_pointer)
-{
-  HIPContextScope scope(this);
-
-  hip_assert(hipHostGetDevicePointer((hipDeviceptr_t *)&device_pointer, shared_pointer, 0));
-  return true;
-}
-
-void HIPDevice::copy_host_to_device(void *device_pointer, void *host_pointer, size_t size)
-{
-  const HIPContextScope scope(this);
-
-  hip_assert(hipMemcpyHtoD((hipDeviceptr_t)device_pointer, host_pointer, size));
+    hip_mem_map.erase(hip_mem_map.find(&mem));
+  }
 }
 
 void HIPDevice::mem_alloc(device_memory &mem)
@@ -576,8 +823,8 @@ void HIPDevice::mem_zero(device_memory &mem)
 
   /* If use_mapped_host of mem is false, mem.device_pointer currently refers to device memory
    * regardless of mem.host_pointer and mem.shared_pointer. */
-  thread_scoped_lock lock(device_mem_map_mutex);
-  if (!device_mem_map[&mem].use_mapped_host || mem.host_pointer != mem.shared_pointer) {
+  thread_scoped_lock lock(hip_mem_map_mutex);
+  if (!hip_mem_map[&mem].use_mapped_host || mem.host_pointer != mem.shared_pointer) {
     const HIPContextScope scope(this);
     hip_assert(hipMemsetD8((hipDeviceptr_t)mem.device_pointer, 0, mem.memory_size()));
   }
@@ -704,19 +951,19 @@ void HIPDevice::tex_alloc(device_texture &mem)
       return;
   }
 
-  Mem *cmem = NULL;
+  HIPMem *cmem = NULL;
   hArray array_3d = NULL;
   size_t src_pitch = mem.data_width * dsize * mem.data_elements;
   size_t dst_pitch = src_pitch;
 
   if (!mem.is_resident(this)) {
-    thread_scoped_lock lock(device_mem_map_mutex);
-    cmem = &device_mem_map[&mem];
+    thread_scoped_lock lock(hip_mem_map_mutex);
+    cmem = &hip_mem_map[&mem];
     cmem->texobject = 0;
 
     if (mem.data_depth > 1) {
       array_3d = (hArray)mem.device_pointer;
-      cmem->array = reinterpret_cast<arrayMemObject>(array_3d);
+      cmem->array = array_3d;
     }
     else if (mem.data_height > 0) {
       dst_pitch = align_up(src_pitch, pitch_alignment);
@@ -760,10 +1007,10 @@ void HIPDevice::tex_alloc(device_texture &mem)
     mem.device_size = size;
     stats.mem_alloc(size);
 
-    thread_scoped_lock lock(device_mem_map_mutex);
-    cmem = &device_mem_map[&mem];
+    thread_scoped_lock lock(hip_mem_map_mutex);
+    cmem = &hip_mem_map[&mem];
     cmem->texobject = 0;
-    cmem->array = reinterpret_cast<arrayMemObject>(array_3d);
+    cmem->array = array_3d;
   }
   else if (mem.data_height > 0) {
     /* 2D texture, using pitch aligned linear memory. */
@@ -848,8 +1095,8 @@ void HIPDevice::tex_alloc(device_texture &mem)
     texDesc.filterMode = filter_mode;
     texDesc.flags = HIP_TRSF_NORMALIZED_COORDINATES;
 
-    thread_scoped_lock lock(device_mem_map_mutex);
-    cmem = &device_mem_map[&mem];
+    thread_scoped_lock lock(hip_mem_map_mutex);
+    cmem = &hip_mem_map[&mem];
 
     hip_assert(hipTexObjectCreate(&cmem->texobject, &resDesc, &texDesc, NULL));
 
@@ -864,9 +1111,9 @@ void HIPDevice::tex_free(device_texture &mem)
 {
   if (mem.device_pointer) {
     HIPContextScope scope(this);
-    thread_scoped_lock lock(device_mem_map_mutex);
-    DCHECK(device_mem_map.find(&mem) != device_mem_map.end());
-    const Mem &cmem = device_mem_map[&mem];
+    thread_scoped_lock lock(hip_mem_map_mutex);
+    DCHECK(hip_mem_map.find(&mem) != hip_mem_map.end());
+    const HIPMem &cmem = hip_mem_map[&mem];
 
     if (cmem.texobject) {
       /* Free bindless texture. */
@@ -875,16 +1122,16 @@ void HIPDevice::tex_free(device_texture &mem)
 
     if (!mem.is_resident(this)) {
       /* Do not free memory here, since it was allocated on a different device. */
-      device_mem_map.erase(device_mem_map.find(&mem));
+      hip_mem_map.erase(hip_mem_map.find(&mem));
     }
     else if (cmem.array) {
       /* Free array. */
-      hipArrayDestroy(reinterpret_cast<hArray>(cmem.array));
+      hipArrayDestroy(cmem.array);
       stats.mem_free(mem.device_size);
       mem.device_pointer = 0;
       mem.device_size = 0;
 
-      device_mem_map.erase(device_mem_map.find(&mem));
+      hip_mem_map.erase(hip_mem_map.find(&mem));
     }
     else {
       lock.unlock();
@@ -906,7 +1153,7 @@ bool HIPDevice::should_use_graphics_interop()
    * possible, but from the empiric measurements it can be considerably slower than using naive
    * pixels copy. */
 
-  /* Disable graphics interop for now, because of driver bug in 21.40. See #92972 */
+  /* Disable graphics interop for now, because of driver bug in 21.40. See T92972 */
 #  if 0
   HIPContextScope scope(this);
 

intern/cycles/device/hip/device_impl.h

@@ -18,7 +18,7 @@ CCL_NAMESPACE_BEGIN
 
 class DeviceQueue;
 
-class HIPDevice : public GPUDevice {
+class HIPDevice : public Device {
 
   friend class HIPContextScope;
 
@@ -26,11 +26,36 @@ class HIPDevice : public GPUDevice {
   hipDevice_t hipDevice;
   hipCtx_t hipContext;
   hipModule_t hipModule;
+  size_t device_texture_headroom;
+  size_t device_working_headroom;
+  bool move_texture_to_host;
+  size_t map_host_used;
+  size_t map_host_limit;
+  int can_map_host;
   int pitch_alignment;
   int hipDevId;
   int hipDevArchitecture;
   bool first_error;
 
+  struct HIPMem {
+    HIPMem() : texobject(0), array(0), use_mapped_host(false)
+    {
+    }
+
+    hipTextureObject_t texobject;
+    hArray array;
+
+    /* If true, a mapped host memory in shared_pointer is being used. */
+    bool use_mapped_host;
+  };
+  typedef map<device_memory *, HIPMem> HIPMemMap;
+  HIPMemMap hip_mem_map;
+  thread_mutex hip_mem_map_mutex;
+
+  /* Bindless Textures */
+  device_vector<TextureInfo> texture_info;
+  bool need_texture_info;
+
   HIPDeviceKernels kernels;
 
   static bool have_precompiled_kernels();
@@ -56,13 +81,17 @@ class HIPDevice : public GPUDevice {
   virtual bool load_kernels(const uint kernel_features) override;
   void reserve_local_memory(const uint kernel_features);
 
-  virtual void get_device_memory_info(size_t &total, size_t &free) override;
-  virtual bool alloc_device(void *&device_pointer, size_t size) override;
-  virtual void free_device(void *device_pointer) override;
-  virtual bool alloc_host(void *&shared_pointer, size_t size) override;
-  virtual void free_host(void *shared_pointer) override;
-  virtual bool transform_host_pointer(void *&device_pointer, void *&shared_pointer) override;
-  virtual void copy_host_to_device(void *device_pointer, void *host_pointer, size_t size) override;
+  void init_host_memory();
+
+  void load_texture_info();
+
+  void move_textures_to_host(size_t size, bool for_texture);
+
+  HIPMem *generic_alloc(device_memory &mem, size_t pitch_padding = 0);
+
+  void generic_copy_to(device_memory &mem);
+
+  void generic_free(device_memory &mem);
 
   void mem_alloc(device_memory &mem) override;
 

intern/cycles/device/kernel.cpp

@@ -73,10 +73,6 @@ const char *device_kernel_as_string(DeviceKernel kernel)
       return "integrator_terminated_paths_array";
     case DEVICE_KERNEL_INTEGRATOR_SORTED_PATHS_ARRAY:
       return "integrator_sorted_paths_array";
-    case DEVICE_KERNEL_INTEGRATOR_SORT_BUCKET_PASS:
-      return "integrator_sort_bucket_pass";
-    case DEVICE_KERNEL_INTEGRATOR_SORT_WRITE_PASS:
-      return "integrator_sort_write_pass";
     case DEVICE_KERNEL_INTEGRATOR_COMPACT_PATHS_ARRAY:
       return "integrator_compact_paths_array";
     case DEVICE_KERNEL_INTEGRATOR_COMPACT_STATES:

intern/cycles/device/memory.h

@@ -247,8 +247,6 @@ class device_memory {
   bool is_resident(Device *sub_device) const;
 
  protected:
-  friend class Device;
-  friend class GPUDevice;
   friend class CUDADevice;
   friend class OptiXDevice;
   friend class HIPDevice;

intern/cycles/device/metal/bvh.h

@@ -21,7 +21,6 @@ class BVHMetal : public BVH {
 
   API_AVAILABLE(macos(11.0))
   vector<id<MTLAccelerationStructure>> blas_array;
-  vector<uint32_t> blas_lookup;
 
   bool motion_blur = false;
 

intern/cycles/device/metal/bvh.mm

@@ -816,11 +816,6 @@ bool BVHMetal::build_TLAS(Progress &progress,
 
     uint32_t instance_index = 0;
     uint32_t motion_transform_index = 0;
-
-    // allocate look up buffer for wost case scenario
-    uint64_t count = objects.size();
-    blas_lookup.resize(count);
-
     for (Object *ob : objects) {
       /* Skip non-traceable objects */
       if (!ob->is_traceable())
@@ -848,15 +843,12 @@ bool BVHMetal::build_TLAS(Progress &progress,
       /* Set user instance ID to object index */
       int object_index = ob->get_device_index();
       uint32_t user_id = uint32_t(object_index);
-      int currIndex = instance_index++;
-      assert(user_id < blas_lookup.size());
-      blas_lookup[user_id] = accel_struct_index;
 
       /* Bake into the appropriate descriptor */
       if (motion_blur) {
         MTLAccelerationStructureMotionInstanceDescriptor *instances =
             (MTLAccelerationStructureMotionInstanceDescriptor *)[instanceBuf contents];
-        MTLAccelerationStructureMotionInstanceDescriptor &desc = instances[currIndex];
+        MTLAccelerationStructureMotionInstanceDescriptor &desc = instances[instance_index++];
 
         desc.accelerationStructureIndex = accel_struct_index;
         desc.userID = user_id;
@@ -902,7 +894,7 @@ bool BVHMetal::build_TLAS(Progress &progress,
       else {
         MTLAccelerationStructureUserIDInstanceDescriptor *instances =
             (MTLAccelerationStructureUserIDInstanceDescriptor *)[instanceBuf contents];
-        MTLAccelerationStructureUserIDInstanceDescriptor &desc = instances[currIndex];
+        MTLAccelerationStructureUserIDInstanceDescriptor &desc = instances[instance_index++];
 
         desc.accelerationStructureIndex = accel_struct_index;
         desc.userID = user_id;

intern/cycles/device/metal/device.mm

@@ -55,10 +55,6 @@ void device_metal_info(vector<DeviceInfo> &devices)
     info.denoisers = DENOISER_NONE;
     info.id = id;
 
-    if (MetalInfo::get_device_vendor(device) == METAL_GPU_AMD) {
-      info.has_light_tree = false;
-    }
-
     devices.push_back(info);
     device_index++;
   }

intern/cycles/device/metal/device_impl.h

@@ -74,11 +74,6 @@ class MetalDevice : public Device {
   id<MTLBuffer> texture_bindings_3d = nil;
   std::vector<id<MTLTexture>> texture_slot_map;
 
-  /* BLAS encoding & lookup */
-  id<MTLArgumentEncoder> mtlBlasArgEncoder = nil;
-  id<MTLBuffer> blas_buffer = nil;
-  id<MTLBuffer> blas_lookup_buffer = nil;
-
   bool use_metalrt = false;
   MetalPipelineType kernel_specialization_level = PSO_GENERIC;
 
@@ -110,8 +105,6 @@ class MetalDevice : public Device {
 
   bool use_adaptive_compilation();
 
-  bool use_local_atomic_sort() const;
-
   bool make_source_and_check_if_compile_needed(MetalPipelineType pso_type);
 
   void make_source(MetalPipelineType pso_type, const uint kernel_features);

intern/cycles/device/metal/device_impl.mm

@@ -105,7 +105,6 @@ MetalDevice::MetalDevice(const DeviceInfo &info, Stats &stats, Profiler &profile
     }
     case METAL_GPU_AMD: {
       max_threads_per_threadgroup = 128;
-      use_metalrt = info.use_metalrt;
       break;
     }
     case METAL_GPU_APPLE: {
@@ -193,10 +192,6 @@ MetalDevice::MetalDevice(const DeviceInfo &info, Stats &stats, Profiler &profile
         arg_desc_as.dataType = MTLDataTypeInstanceAccelerationStructure;
         arg_desc_as.access = MTLArgumentAccessReadOnly;
 
-        MTLArgumentDescriptor *arg_desc_ptrs = [[MTLArgumentDescriptor alloc] init];
-        arg_desc_ptrs.dataType = MTLDataTypePointer;
-        arg_desc_ptrs.access = MTLArgumentAccessReadOnly;
-
         MTLArgumentDescriptor *arg_desc_ift = [[MTLArgumentDescriptor alloc] init];
         arg_desc_ift.dataType = MTLDataTypeIntersectionFunctionTable;
         arg_desc_ift.access = MTLArgumentAccessReadOnly;
@@ -209,28 +204,14 @@ MetalDevice::MetalDevice(const DeviceInfo &info, Stats &stats, Profiler &profile
         [ancillary_desc addObject:[arg_desc_ift copy]]; /* ift_shadow */
         arg_desc_ift.index = index++;
         [ancillary_desc addObject:[arg_desc_ift copy]]; /* ift_local */
-        arg_desc_ift.index = index++;
-        [ancillary_desc addObject:[arg_desc_ift copy]]; /* ift_local_prim */
-        arg_desc_ptrs.index = index++;
-        [ancillary_desc addObject:[arg_desc_ptrs copy]]; /* blas array */
-        arg_desc_ptrs.index = index++;
-        [ancillary_desc addObject:[arg_desc_ptrs copy]]; /* look up table for blas */
 
         [arg_desc_ift release];
         [arg_desc_as release];
-        [arg_desc_ptrs release];
       }
     }
 
     mtlAncillaryArgEncoder = [mtlDevice newArgumentEncoderWithArguments:ancillary_desc];
 
-    // preparing the blas arg encoder
-    MTLArgumentDescriptor *arg_desc_blas = [[MTLArgumentDescriptor alloc] init];
-    arg_desc_blas.dataType = MTLDataTypeInstanceAccelerationStructure;
-    arg_desc_blas.access = MTLArgumentAccessReadOnly;
-    mtlBlasArgEncoder = [mtlDevice newArgumentEncoderWithArguments:@[ arg_desc_blas ]];
-    [arg_desc_blas release];
-
     for (int i = 0; i < ancillary_desc.count; i++) {
       [ancillary_desc[i] release];
     }
@@ -290,11 +271,6 @@ bool MetalDevice::use_adaptive_compilation()
   return DebugFlags().metal.adaptive_compile;
 }
 
-bool MetalDevice::use_local_atomic_sort() const
-{
-  return DebugFlags().metal.use_local_atomic_sort;
-}
-
 void MetalDevice::make_source(MetalPipelineType pso_type, const uint kernel_features)
 {
   string global_defines;
@@ -302,10 +278,6 @@ void MetalDevice::make_source(MetalPipelineType pso_type, const uint kernel_feat
     global_defines += "#define __KERNEL_FEATURES__ " + to_string(kernel_features) + "\n";
   }
 
-  if (use_local_atomic_sort()) {
-    global_defines += "#define __KERNEL_LOCAL_ATOMIC_SORT__\n";
-  }
-
   if (use_metalrt) {
     global_defines += "#define __METALRT__\n";
     if (motion_blur) {
@@ -355,21 +327,10 @@ void MetalDevice::make_source(MetalPipelineType pso_type, const uint kernel_feat
 #  define KERNEL_STRUCT_BEGIN(name, parent) \
     string_replace_same_length(source, "kernel_data." #parent ".", "kernel_data_" #parent "_");
 
-    bool next_member_is_specialized = true;
-
-#  define KERNEL_STRUCT_MEMBER_DONT_SPECIALIZE next_member_is_specialized = false;
-
     /* Add constants to md5 so that 'get_best_pipeline' is able to return a suitable match. */
 #  define KERNEL_STRUCT_MEMBER(parent, _type, name) \
-    if (next_member_is_specialized) { \
-      baked_constants += string(#parent "." #name "=") + \
-                         to_string(_type(launch_params.data.parent.name)) + "\n"; \
-    } \
-    else { \
-      string_replace( \
-          source, "kernel_data_" #parent "_" #name, "kernel_data." #parent ".__unused_" #name); \
-      next_member_is_specialized = true; \
-    }
+    baked_constants += string(#parent "." #name "=") + \
+                       to_string(_type(launch_params.data.parent.name)) + "\n";
 
 #  include "kernel/data_template.h"
 
@@ -586,7 +547,7 @@ void MetalDevice::erase_allocation(device_memory &mem)
   if (it != metal_mem_map.end()) {
     MetalMem *mmem = it->second.get();
 
-    /* blank out reference to MetalMem* in the launch params (fixes crash #94736) */
+    /* blank out reference to MetalMem* in the launch params (fixes crash T94736) */
     if (mmem->pointer_index >= 0) {
       device_ptr *pointers = (device_ptr *)&launch_params;
       pointers[mmem->pointer_index] = 0;
@@ -1259,33 +1220,6 @@ void MetalDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
     if (@available(macos 11.0, *)) {
       if (bvh->params.top_level) {
         bvhMetalRT = bvh_metal;
-
-        // allocate required buffers for BLAS array
-        uint64_t count = bvhMetalRT->blas_array.size();
-        uint64_t bufferSize = mtlBlasArgEncoder.encodedLength * count;
-        blas_buffer = [mtlDevice newBufferWithLength:bufferSize options:default_storage_mode];
-        stats.mem_alloc(blas_buffer.allocatedSize);
-
-        for (uint64_t i = 0; i < count; ++i) {
-          [mtlBlasArgEncoder setArgumentBuffer:blas_buffer
-                                        offset:i * mtlBlasArgEncoder.encodedLength];
-          [mtlBlasArgEncoder setAccelerationStructure:bvhMetalRT->blas_array[i] atIndex:0];
-        }
-
-        count = bvhMetalRT->blas_lookup.size();
-        bufferSize = sizeof(uint32_t) * count;
-        blas_lookup_buffer = [mtlDevice newBufferWithLength:bufferSize
-                                                    options:default_storage_mode];
-        stats.mem_alloc(blas_lookup_buffer.allocatedSize);
-
-        memcpy([blas_lookup_buffer contents],
-               bvhMetalRT -> blas_lookup.data(),
-               blas_lookup_buffer.allocatedSize);
-
-        if (default_storage_mode == MTLResourceStorageModeManaged) {
-          [blas_buffer didModifyRange:NSMakeRange(0, blas_buffer.length)];
-          [blas_lookup_buffer didModifyRange:NSMakeRange(0, blas_lookup_buffer.length)];
-        }
       }
     }
   }

intern/cycles/device/metal/kernel.h

@@ -19,8 +19,6 @@ enum {
   METALRT_FUNC_SHADOW_BOX,
   METALRT_FUNC_LOCAL_TRI,
   METALRT_FUNC_LOCAL_BOX,
-  METALRT_FUNC_LOCAL_TRI_PRIM,
-  METALRT_FUNC_LOCAL_BOX_PRIM,
   METALRT_FUNC_CURVE_RIBBON,
   METALRT_FUNC_CURVE_RIBBON_SHADOW,
   METALRT_FUNC_CURVE_ALL,
@@ -30,13 +28,7 @@ enum {
   METALRT_FUNC_NUM
 };
 
-enum {
-  METALRT_TABLE_DEFAULT,
-  METALRT_TABLE_SHADOW,
-  METALRT_TABLE_LOCAL,
-  METALRT_TABLE_LOCAL_PRIM,
-  METALRT_TABLE_NUM
-};
+enum { METALRT_TABLE_DEFAULT, METALRT_TABLE_SHADOW, METALRT_TABLE_LOCAL, METALRT_TABLE_NUM };
 
 /* Pipeline State Object types */
 enum MetalPipelineType {

intern/cycles/device/metal/kernel.mm

@@ -49,18 +49,6 @@ struct ShaderCache {
     if (MetalInfo::get_device_vendor(mtlDevice) == METAL_GPU_APPLE) {
       switch (MetalInfo::get_apple_gpu_architecture(mtlDevice)) {
         default:
-        case APPLE_M2_BIG:
-          occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_COMPACT_SHADOW_STATES] = {384, 128};
-          occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_INIT_FROM_CAMERA] = {640, 128};
-          occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_INTERSECT_CLOSEST] = {1024, 64};
-          occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_INTERSECT_SHADOW] = {704, 704};
-          occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_INTERSECT_SUBSURFACE] = {640, 32};
-          occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_QUEUED_PATHS_ARRAY] = {896, 768};
-          occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_SHADE_BACKGROUND] = {512, 128};
-          occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_SHADE_SHADOW] = {32, 32};
-          occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE] = {768, 576};
-          occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_SORTED_PATHS_ARRAY] = {896, 768};
-          break;
         case APPLE_M2:
           occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_COMPACT_SHADOW_STATES] = {32, 32};
           occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_INIT_FROM_CAMERA] = {832, 32};
@@ -87,9 +75,6 @@ struct ShaderCache {
           break;
       }
     }
-
-    occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_SORT_BUCKET_PASS] = {1024, 1024};
-    occupancy_tuning[DEVICE_KERNEL_INTEGRATOR_SORT_WRITE_PASS] = {1024, 1024};
   }
   ~ShaderCache();
 
@@ -463,18 +448,13 @@ static MTLFunctionConstantValues *GetConstantValues(KernelData const *data = nul
   if (!data) {
     data = &zero_data;
   }
-  [constant_values setConstantValue:&zero_data type:MTLDataType_int atIndex:Kernel_DummyConstant];
-
-  bool next_member_is_specialized = true;
-
-#  define KERNEL_STRUCT_MEMBER_DONT_SPECIALIZE next_member_is_specialized = false;
+  int zero_int = 0;
+  [constant_values setConstantValue:&zero_int type:MTLDataType_int atIndex:Kernel_DummyConstant];
 
 #  define KERNEL_STRUCT_MEMBER(parent, _type, name) \
-    [constant_values setConstantValue:next_member_is_specialized ? (void *)&data->parent.name : \
-                                                                   (void *)&zero_data \
+    [constant_values setConstantValue:&data->parent.name \
                                  type:MTLDataType_##_type \
-                              atIndex:KernelData_##parent##_##name]; \
-    next_member_is_specialized = true;
+                              atIndex:KernelData_##parent##_##name];
 
 #  include "kernel/data_template.h"
 
@@ -524,8 +504,6 @@ void MetalKernelPipeline::compile()
           "__anyhit__cycles_metalrt_shadow_all_hit_box",
           "__anyhit__cycles_metalrt_local_hit_tri",
           "__anyhit__cycles_metalrt_local_hit_box",
-          "__anyhit__cycles_metalrt_local_hit_tri_prim",
-          "__anyhit__cycles_metalrt_local_hit_box_prim",
           "__intersection__curve_ribbon",
           "__intersection__curve_ribbon_shadow",
           "__intersection__curve_all",
@@ -616,17 +594,11 @@ void MetalKernelPipeline::compile()
                          rt_intersection_function[METALRT_FUNC_LOCAL_BOX],
                          rt_intersection_function[METALRT_FUNC_LOCAL_BOX],
                          nil];
-    table_functions[METALRT_TABLE_LOCAL_PRIM] = [NSArray
-        arrayWithObjects:rt_intersection_function[METALRT_FUNC_LOCAL_TRI_PRIM],
-                         rt_intersection_function[METALRT_FUNC_LOCAL_BOX_PRIM],
-                         rt_intersection_function[METALRT_FUNC_LOCAL_BOX_PRIM],
-                         nil];
 
     NSMutableSet *unique_functions = [NSMutableSet
         setWithArray:table_functions[METALRT_TABLE_DEFAULT]];
     [unique_functions addObjectsFromArray:table_functions[METALRT_TABLE_SHADOW]];
     [unique_functions addObjectsFromArray:table_functions[METALRT_TABLE_LOCAL]];
-    [unique_functions addObjectsFromArray:table_functions[METALRT_TABLE_LOCAL_PRIM]];
 
     if (kernel_has_intersection(device_kernel)) {
       linked_functions = [[NSArray arrayWithArray:[unique_functions allObjects]]

intern/cycles/device/metal/queue.h

@@ -25,7 +25,6 @@ class MetalDeviceQueue : public DeviceQueue {
   virtual int num_concurrent_states(const size_t) const override;
   virtual int num_concurrent_busy_states(const size_t) const override;
   virtual int num_sort_partition_elements() const override;
-  virtual bool supports_local_atomic_sort() const override;
 
   virtual void init_execution() override;
 

intern/cycles/device/metal/queue.mm

@@ -278,8 +278,7 @@ int MetalDeviceQueue::num_concurrent_states(const size_t state_size) const
   if (metal_device_->device_vendor == METAL_GPU_APPLE) {
     result *= 4;
 
-    /* Increasing the state count doesn't notably benefit M1-family systems.  */
-    if (MetalInfo::get_apple_gpu_architecture(metal_device_->mtlDevice) != APPLE_M1) {
+    if (MetalInfo::get_apple_gpu_architecture(metal_device_->mtlDevice) == APPLE_M2) {
       size_t system_ram = system_physical_ram();
       size_t allocated_so_far = [metal_device_->mtlDevice currentAllocatedSize];
       size_t max_recommended_working_set = [metal_device_->mtlDevice recommendedMaxWorkingSetSize];
@@ -315,11 +314,6 @@ int MetalDeviceQueue::num_sort_partition_elements() const
   return MetalInfo::optimal_sort_partition_elements(metal_device_->mtlDevice);
 }
 
-bool MetalDeviceQueue::supports_local_atomic_sort() const
-{
-  return metal_device_->use_local_atomic_sort();
-}
-
 void MetalDeviceQueue::init_execution()
 {
   /* Synchronize all textures and memory copies before executing task. */
@@ -482,12 +476,6 @@ bool MetalDeviceQueue::enqueue(DeviceKernel kernel,
       if (metal_device_->bvhMetalRT) {
         id<MTLAccelerationStructure> accel_struct = metal_device_->bvhMetalRT->accel_struct;
         [metal_device_->mtlAncillaryArgEncoder setAccelerationStructure:accel_struct atIndex:2];
-        [metal_device_->mtlAncillaryArgEncoder setBuffer:metal_device_->blas_buffer
-                                                  offset:0
-                                                 atIndex:7];
-        [metal_device_->mtlAncillaryArgEncoder setBuffer:metal_device_->blas_lookup_buffer
-                                                  offset:0
-                                                 atIndex:8];
       }
 
       for (int table = 0; table < METALRT_TABLE_NUM; table++) {
@@ -538,10 +526,6 @@ bool MetalDeviceQueue::enqueue(DeviceKernel kernel,
       if (bvhMetalRT) {
         /* Mark all Accelerations resources as used */
         [mtlComputeCommandEncoder useResource:bvhMetalRT->accel_struct usage:MTLResourceUsageRead];
-        [mtlComputeCommandEncoder useResource:metal_device_->blas_buffer
-                                        usage:MTLResourceUsageRead];
-        [mtlComputeCommandEncoder useResource:metal_device_->blas_lookup_buffer
-                                        usage:MTLResourceUsageRead];
         [mtlComputeCommandEncoder useResources:bvhMetalRT->blas_array.data()
                                          count:bvhMetalRT->blas_array.size()
                                          usage:MTLResourceUsageRead];
@@ -568,24 +552,13 @@ bool MetalDeviceQueue::enqueue(DeviceKernel kernel,
       /* See parallel_active_index.h for why this amount of shared memory is needed.
        * Rounded up to 16 bytes for Metal */
       shared_mem_bytes = (int)round_up((num_threads_per_block + 1) * sizeof(int), 16);
+      [mtlComputeCommandEncoder setThreadgroupMemoryLength:shared_mem_bytes atIndex:0];
       break;
 
-    case DEVICE_KERNEL_INTEGRATOR_SORT_BUCKET_PASS:
-    case DEVICE_KERNEL_INTEGRATOR_SORT_WRITE_PASS: {
-      int key_count = metal_device_->launch_params.data.max_shaders;
-      shared_mem_bytes = (int)round_up(key_count * sizeof(int), 16);
-      break;
-    }
-
     default:
       break;
   }
 
-  if (shared_mem_bytes) {
-    assert(shared_mem_bytes <= 32 * 1024);
-    [mtlComputeCommandEncoder setThreadgroupMemoryLength:shared_mem_bytes atIndex:0];
-  }
-
   MTLSize size_threadgroups_per_dispatch = MTLSizeMake(
       divide_up(work_size, num_threads_per_block), 1, 1);
   MTLSize size_threads_per_threadgroup = MTLSizeMake(num_threads_per_block, 1, 1);

intern/cycles/device/metal/util.h

@@ -29,7 +29,6 @@ enum AppleGPUArchitecture {
   APPLE_UNKNOWN,
   APPLE_M1,
   APPLE_M2,
-  APPLE_M2_BIG,
 };
 
 /* Contains static Metal helper functions. */

intern/cycles/device/metal/util.mm

@@ -52,7 +52,7 @@ AppleGPUArchitecture MetalInfo::get_apple_gpu_architecture(id<MTLDevice> device)
     return APPLE_M1;
   }
   else if (strstr(device_name, "M2")) {
-    return get_apple_gpu_core_count(device) <= 10 ? APPLE_M2 : APPLE_M2_BIG;
+    return APPLE_M2;
   }
   return APPLE_UNKNOWN;
 }
@@ -64,12 +64,6 @@ MetalGPUVendor MetalInfo::get_device_vendor(id<MTLDevice> device)
     return METAL_GPU_INTEL;
   }
   else if (strstr(device_name, "AMD")) {
-    /* Setting this env var hides AMD devices thus exposing any integrated Intel devices. */
-    if (auto str = getenv("CYCLES_METAL_FORCE_INTEL")) {
-      if (atoi(str)) {
-        return METAL_GPU_UNKNOWN;
-      }
-    }
     return METAL_GPU_AMD;
   }
   else if (strstr(device_name, "Apple")) {
@@ -102,15 +96,6 @@ vector<id<MTLDevice>> const &MetalInfo::get_usable_devices()
     return usable_devices;
   }
 
-  /* If the system has both an AMD GPU (discrete) and an Intel one (integrated), prefer the AMD
-   * one. This can be overridden with CYCLES_METAL_FORCE_INTEL. */
-  bool has_usable_amd_gpu = false;
-  if (@available(macos 12.3, *)) {
-    for (id<MTLDevice> device in MTLCopyAllDevices()) {
-      has_usable_amd_gpu |= (get_device_vendor(device) == METAL_GPU_AMD);
-    }
-  }
-
   metal_printf("Usable Metal devices:\n");
   for (id<MTLDevice> device in MTLCopyAllDevices()) {
     string device_name = get_device_name(device);
@@ -126,10 +111,8 @@ vector<id<MTLDevice>> const &MetalInfo::get_usable_devices()
     }
 
 #  if defined(MAC_OS_VERSION_13_0)
-    if (!has_usable_amd_gpu) {
-      if (@available(macos 13.0, *)) {
-        usable |= (vendor == METAL_GPU_INTEL);
-      }
+    if (@available(macos 13.0, *)) {
+      usable |= (vendor == METAL_GPU_INTEL);
     }
 #  endif
 

intern/cycles/device/oneapi/device_impl.cpp

@@ -377,7 +377,7 @@ void OneapiDevice::tex_alloc(device_texture &mem)
   generic_alloc(mem);
   generic_copy_to(mem);
 
-  /* Resize if needed. Also, in case of resize - allocate in advance for future allocations. */
+  /* Resize if needed. Also, in case of resize - allocate in advance for future allocs. */
   const uint slot = mem.slot;
   if (slot >= texture_info_.size()) {
     texture_info_.resize(slot + 128);
@@ -631,9 +631,9 @@ bool OneapiDevice::enqueue_kernel(KernelContext *kernel_context,
 /* Compute-runtime (ie. NEO) version is what gets returned by sycl/L0 on Windows
  * since Windows driver 101.3268. */
 /* The same min compute-runtime version is currently required across Windows and Linux.
- * For Windows driver 101.4032, compute-runtime version is 24931. */
-static const int lowest_supported_driver_version_win = 1014032;
-static const int lowest_supported_driver_version_neo = 24931;
+ * For Windows driver 101.3430, compute-runtime version is 23904. */
+static const int lowest_supported_driver_version_win = 1013430;
+static const int lowest_supported_driver_version_neo = 23904;
 
 int OneapiDevice::parse_driver_build_version(const sycl::device &device)
 {

intern/cycles/device/optix/device_impl.cpp

@@ -854,14 +854,12 @@ bool OptiXDevice::load_osl_kernels()
         context, group_descs, 2, &group_options, nullptr, 0, &osl_groups[i * 2]));
   }
 
-  OptixStackSizes stack_size[NUM_PROGRAM_GROUPS] = {};
   vector<OptixStackSizes> osl_stack_size(osl_groups.size());
 
   /* Update SBT with new entries. */
   sbt_data.alloc(NUM_PROGRAM_GROUPS + osl_groups.size());
   for (int i = 0; i < NUM_PROGRAM_GROUPS; ++i) {
     optix_assert(optixSbtRecordPackHeader(groups[i], &sbt_data[i]));
-    optix_assert(optixProgramGroupGetStackSize(groups[i], &stack_size[i]));
   }
   for (size_t i = 0; i < osl_groups.size(); ++i) {
     if (osl_groups[i] != NULL) {
@@ -909,15 +907,13 @@ bool OptiXDevice::load_osl_kernels()
                                      0,
                                      &pipelines[PIP_SHADE]));
 
-    const unsigned int css = std::max(stack_size[PG_RGEN_SHADE_SURFACE_RAYTRACE].cssRG,
-                                      stack_size[PG_RGEN_SHADE_SURFACE_MNEE].cssRG);
     unsigned int dss = 0;
     for (unsigned int i = 0; i < osl_stack_size.size(); ++i) {
       dss = std::max(dss, osl_stack_size[i].dssDC);
     }
 
     optix_assert(optixPipelineSetStackSize(
-        pipelines[PIP_SHADE], 0, dss, css, pipeline_options.usesMotionBlur ? 3 : 2));
+        pipelines[PIP_SHADE], 0, dss, 0, pipeline_options.usesMotionBlur ? 3 : 2));
   }
 
   return !have_error();

intern/cycles/device/queue.h

@@ -112,13 +112,6 @@ class DeviceQueue {
     return 65536;
   }
 
-  /* Does device support local atomic sorting kernels (INTEGRATOR_SORT_BUCKET_PASS and
-   * INTEGRATOR_SORT_WRITE_PASS)? */
-  virtual bool supports_local_atomic_sort() const
-  {
-    return false;
-  }
-
   /* Initialize execution of kernels on this queue.
    *
    * Will, for example, load all data required by the kernels from Device to global or path state.

intern/cycles/graph/CMakeLists.txt

@@ -5,9 +5,6 @@ set(INC
   ..
 )
 
-set(INC_SYS
-)
-
 set(SRC
   node.cpp
   node_type.cpp

intern/cycles/integrator/CMakeLists.txt

@@ -5,9 +5,6 @@ set(INC
   ..
 )
 
-set(INC_SYS
-)
-
 set(SRC
   adaptive_sampling.cpp
   denoiser.cpp

intern/cycles/integrator/path_trace_work_gpu.cpp

@@ -71,8 +71,6 @@ PathTraceWorkGPU::PathTraceWorkGPU(Device *device,
           device, "integrator_shader_mnee_sort_counter", MEM_READ_WRITE),
       integrator_shader_sort_prefix_sum_(
           device, "integrator_shader_sort_prefix_sum", MEM_READ_WRITE),
-      integrator_shader_sort_partition_key_offsets_(
-          device, "integrator_shader_sort_partition_key_offsets", MEM_READ_WRITE),
       integrator_next_main_path_index_(device, "integrator_next_main_path_index", MEM_READ_WRITE),
       integrator_next_shadow_path_index_(
           device, "integrator_next_shadow_path_index", MEM_READ_WRITE),
@@ -209,45 +207,33 @@ void PathTraceWorkGPU::alloc_integrator_sorting()
   integrator_state_gpu_.sort_partition_divisor = (int)divide_up(max_num_paths_,
                                                                 num_sort_partitions_);
 
-  if (num_sort_partitions_ > 1 && queue_->supports_local_atomic_sort()) {
-    /* Allocate array for partitioned shader sorting using local atomics. */
-    const int num_offsets = (device_scene_->data.max_shaders + 1) * num_sort_partitions_;
-    if (integrator_shader_sort_partition_key_offsets_.size() < num_offsets) {
-      integrator_shader_sort_partition_key_offsets_.alloc(num_offsets);
-      integrator_shader_sort_partition_key_offsets_.zero_to_device();
-    }
-    integrator_state_gpu_.sort_partition_key_offsets =
-        (int *)integrator_shader_sort_partition_key_offsets_.device_pointer;
+  /* Allocate arrays for shader sorting. */
+  const int sort_buckets = device_scene_->data.max_shaders * num_sort_partitions_;
+  if (integrator_shader_sort_counter_.size() < sort_buckets) {
+    integrator_shader_sort_counter_.alloc(sort_buckets);
+    integrator_shader_sort_counter_.zero_to_device();
+    integrator_state_gpu_.sort_key_counter[DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE] =
+        (int *)integrator_shader_sort_counter_.device_pointer;
+
+    integrator_shader_sort_prefix_sum_.alloc(sort_buckets);
+    integrator_shader_sort_prefix_sum_.zero_to_device();
   }
-  else {
-    /* Allocate arrays for shader sorting. */
-    const int sort_buckets = device_scene_->data.max_shaders * num_sort_partitions_;
-    if (integrator_shader_sort_counter_.size() < sort_buckets) {
-      integrator_shader_sort_counter_.alloc(sort_buckets);
-      integrator_shader_sort_counter_.zero_to_device();
-      integrator_state_gpu_.sort_key_counter[DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE] =
-          (int *)integrator_shader_sort_counter_.device_pointer;
 
-      integrator_shader_sort_prefix_sum_.alloc(sort_buckets);
-      integrator_shader_sort_prefix_sum_.zero_to_device();
+  if (device_scene_->data.kernel_features & KERNEL_FEATURE_NODE_RAYTRACE) {
+    if (integrator_shader_raytrace_sort_counter_.size() < sort_buckets) {
+      integrator_shader_raytrace_sort_counter_.alloc(sort_buckets);
+      integrator_shader_raytrace_sort_counter_.zero_to_device();
+      integrator_state_gpu_.sort_key_counter[DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE] =
+          (int *)integrator_shader_raytrace_sort_counter_.device_pointer;
     }
+  }
 
-    if (device_scene_->data.kernel_features & KERNEL_FEATURE_NODE_RAYTRACE) {
-      if (integrator_shader_raytrace_sort_counter_.size() < sort_buckets) {
-        integrator_shader_raytrace_sort_counter_.alloc(sort_buckets);
-        integrator_shader_raytrace_sort_counter_.zero_to_device();
-        integrator_state_gpu_.sort_key_counter[DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_RAYTRACE] =
-            (int *)integrator_shader_raytrace_sort_counter_.device_pointer;
-      }
-    }
-
-    if (device_scene_->data.kernel_features & KERNEL_FEATURE_MNEE) {
-      if (integrator_shader_mnee_sort_counter_.size() < sort_buckets) {
-        integrator_shader_mnee_sort_counter_.alloc(sort_buckets);
-        integrator_shader_mnee_sort_counter_.zero_to_device();
-        integrator_state_gpu_.sort_key_counter[DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE] =
-            (int *)integrator_shader_mnee_sort_counter_.device_pointer;
-      }
+  if (device_scene_->data.kernel_features & KERNEL_FEATURE_MNEE) {
+    if (integrator_shader_mnee_sort_counter_.size() < sort_buckets) {
+      integrator_shader_mnee_sort_counter_.alloc(sort_buckets);
+      integrator_shader_mnee_sort_counter_.zero_to_device();
+      integrator_state_gpu_.sort_key_counter[DEVICE_KERNEL_INTEGRATOR_SHADE_SURFACE_MNEE] =
+          (int *)integrator_shader_mnee_sort_counter_.device_pointer;
     }
   }
 }
@@ -465,7 +451,8 @@ void PathTraceWorkGPU::enqueue_path_iteration(DeviceKernel kernel, const int num
     work_size = num_queued;
     d_path_index = queued_paths_.device_pointer;
 
-    compute_sorted_queued_paths(kernel, num_paths_limit);
+    compute_sorted_queued_paths(
+        DEVICE_KERNEL_INTEGRATOR_SORTED_PATHS_ARRAY, kernel, num_paths_limit);
   }
   else if (num_queued < work_size) {
     work_size = num_queued;
@@ -524,26 +511,11 @@ void PathTraceWorkGPU::enqueue_path_iteration(DeviceKernel kernel, const int num
   }
 }
 
-void PathTraceWorkGPU::compute_sorted_queued_paths(DeviceKernel queued_kernel,
+void PathTraceWorkGPU::compute_sorted_queued_paths(DeviceKernel kernel,
+                                                   DeviceKernel queued_kernel,
                                                    const int num_paths_limit)
 {
   int d_queued_kernel = queued_kernel;
-
-  /* Launch kernel to fill the active paths arrays. */
-  if (num_sort_partitions_ > 1 && queue_->supports_local_atomic_sort()) {
-    const int work_size = kernel_max_active_main_path_index(queued_kernel);
-    device_ptr d_queued_paths = queued_paths_.device_pointer;
-
-    int partition_size = (int)integrator_state_gpu_.sort_partition_divisor;
-
-    DeviceKernelArguments args(
-        &work_size, &partition_size, &num_paths_limit, &d_queued_paths, &d_queued_kernel);
-
-    queue_->enqueue(DEVICE_KERNEL_INTEGRATOR_SORT_BUCKET_PASS, 1024 * num_sort_partitions_, args);
-    queue_->enqueue(DEVICE_KERNEL_INTEGRATOR_SORT_WRITE_PASS, 1024 * num_sort_partitions_, args);
-    return;
-  }
-
   device_ptr d_counter = (device_ptr)integrator_state_gpu_.sort_key_counter[d_queued_kernel];
   device_ptr d_prefix_sum = integrator_shader_sort_prefix_sum_.device_pointer;
   assert(d_counter != 0 && d_prefix_sum != 0);
@@ -580,7 +552,7 @@ void PathTraceWorkGPU::compute_sorted_queued_paths(DeviceKernel queued_kernel,
                                &d_prefix_sum,
                                &d_queued_kernel);
 
-    queue_->enqueue(DEVICE_KERNEL_INTEGRATOR_SORTED_PATHS_ARRAY, work_size, args);
+    queue_->enqueue(kernel, work_size, args);
   }
 }
 

intern/cycles/integrator/path_trace_work_gpu.h

@@ -70,7 +70,9 @@ class PathTraceWorkGPU : public PathTraceWork {
   void enqueue_path_iteration(DeviceKernel kernel, const int num_paths_limit = INT_MAX);
 
   void compute_queued_paths(DeviceKernel kernel, DeviceKernel queued_kernel);
-  void compute_sorted_queued_paths(DeviceKernel queued_kernel, const int num_paths_limit);
+  void compute_sorted_queued_paths(DeviceKernel kernel,
+                                   DeviceKernel queued_kernel,
+                                   const int num_paths_limit);
 
   void compact_main_paths(const int num_active_paths);
   void compact_shadow_paths();
@@ -133,7 +135,6 @@ class PathTraceWorkGPU : public PathTraceWork {
   device_vector<int> integrator_shader_raytrace_sort_counter_;
   device_vector<int> integrator_shader_mnee_sort_counter_;
   device_vector<int> integrator_shader_sort_prefix_sum_;
-  device_vector<int> integrator_shader_sort_partition_key_offsets_;
   /* Path split. */
   device_vector<int> integrator_next_main_path_index_;
   device_vector<int> integrator_next_shadow_path_index_;

intern/cycles/integrator/render_scheduler.cpp

@@ -886,7 +886,7 @@ int RenderScheduler::get_num_samples_during_navigation(int resolution_divider) c
 {
   /* Special trick for fast navigation: schedule multiple samples during fast navigation
    * (which will prefer to use lower resolution to keep up with refresh rate). This gives more
-   * usable visual feedback for artists. */
+   * usable visual feedback for artists. There are a couple of tricks though. */
 
   if (is_denoise_active_during_update()) {
     /* When denoising is used during navigation prefer using a higher resolution with less samples
@@ -896,12 +896,25 @@ int RenderScheduler::get_num_samples_during_navigation(int resolution_divider) c
     return 1;
   }
 
-  /* Schedule samples equal to the resolution divider up to a maximum of 4.
-   * The idea is to have enough information on the screen by increasing the sample count as the
-   * resolution is decreased. */
-  /* NOTE: Changing this formula will change the formula in
-   * `RenderScheduler::calculate_resolution_divider_for_time()`. */
-  return min(max(1, resolution_divider / pixel_size_), 4);
+  if (resolution_divider <= pixel_size_) {
+    /* When resolution divider is at or below pixel size, schedule one sample. This doesn't effect
+     * the sample count at this resolution division, but instead assists in the calculation of
+     * the resolution divider. */
+    return 1;
+  }
+
+  if (resolution_divider == pixel_size_ * 2) {
+    /* When resolution divider is the previous step to the final resolution, schedule two samples.
+     * This is so that rendering on lower resolution does not exceed time that it takes to render
+     * first sample at the full resolution. */
+    return 2;
+  }
+
+  /* Always render 4 samples, even if scene is configured for less.
+   * The idea here is to have enough information on the screen. Resolution divider of 2 allows us
+   * to have 4 time extra samples, so overall worst case timing is the same as the final resolution
+   * at one sample. */
+  return 4;
 }
 
 bool RenderScheduler::work_need_adaptive_filter() const
@@ -1087,10 +1100,9 @@ void RenderScheduler::update_start_resolution_divider()
   /* TODO(sergey): Need to add hysteresis to avoid resolution divider bouncing around when actual
    * render time is somewhere on a boundary between two resolutions. */
 
-  /* Don't let resolution drop below the desired one. It's better to be slow than provide an
-   * unreadable viewport render. */
-  start_resolution_divider_ = min(resolution_divider_for_update,
-                                  default_start_resolution_divider_);
+  /* Never increase resolution to higher than the pixel size (which is possible if the scene is
+   * simple and compute device is fast). */
+  start_resolution_divider_ = max(resolution_divider_for_update, pixel_size_);
 
   VLOG_WORK << "Calculated resolution divider is " << start_resolution_divider_;
 }
@@ -1175,24 +1187,24 @@ void RenderScheduler::check_time_limit_reached()
 
 int RenderScheduler::calculate_resolution_divider_for_time(double desired_time, double actual_time)
 {
-  const double ratio_between_times = actual_time / desired_time;
+  /* TODO(sergey): There should a non-iterative analytical formula here. */
 
-  /* We can pass `ratio_between_times` to `get_num_samples_during_navigation()` to get our
-   * navigation samples because the equation for calculating the resolution divider is as follows:
-   * `actual_time / desired_time = sqr(resolution_divider) / sample_count`.
-   * While `resolution_divider` is less than or equal to 4, `resolution_divider = sample_count`
-   * (This relationship is determined in `get_num_samples_during_navigation()`). With some
-   * substitution we end up with `actual_time / desired_time = resolution_divider` while the
-   * resolution divider is less than or equal to 4. Once the resolution divider increases above 4,
-   * the relationship of `actual_time / desired_time = resolution_divider` is no longer true,
-   * however the sample count retrieved from `get_num_samples_during_navigation()` is still
-   * accurate if we continue using this assumption. It should be noted that the interaction between
-   * `pixel_size`, sample count, and resolution divider are automatically accounted for and that's
-   * why `pixel_size` isn't included in any of the equations. */
-  const int navigation_samples = get_num_samples_during_navigation(
-      ceil_to_int(ratio_between_times));
+  int resolution_divider = 1;
 
-  return ceil_to_int(sqrt(navigation_samples * ratio_between_times));
+  /* This algorithm iterates through resolution dividers until a divider is found that achieves
+   * the desired render time. A limit of default_start_resolution_divider_ is put in place as the
+   * maximum resolution divider to avoid an unreadable viewport due to a low resolution.
+   * pre_resolution_division_samples and post_resolution_division_samples are used in this
+   * calculation to better predict the performance impact of changing resolution divisions as
+   * the sample count can also change between resolution divisions. */
+  while (actual_time > desired_time && resolution_divider < default_start_resolution_divider_) {
+    int pre_resolution_division_samples = get_num_samples_during_navigation(resolution_divider);
+    resolution_divider = resolution_divider * 2;
+    int post_resolution_division_samples = get_num_samples_during_navigation(resolution_divider);
+    actual_time /= 4.0 * pre_resolution_division_samples / post_resolution_division_samples;
+  }
+
+  return resolution_divider;
 }
 
 int calculate_resolution_divider_for_resolution(int width, int height, int resolution)