realize non-standard virtual arrays

CMakeLists.txt

@@ -419,8 +419,6 @@ mark_as_advanced(WITH_CYCLES_NATIVE_ONLY)
 
 option(WITH_CYCLES_DEVICE_CUDA              "Enable Cycles CUDA compute support" ON)
 option(WITH_CYCLES_DEVICE_OPTIX             "Enable Cycles OptiX support" ON)
-option(WITH_CYCLES_DEVICE_HIP               "Enable Cycles HIP support" OFF)
-mark_as_advanced(WITH_CYCLES_DEVICE_HIP)
 mark_as_advanced(WITH_CYCLES_DEVICE_CUDA)
 
 option(WITH_CUDA_DYNLOAD "Dynamically load CUDA libraries at runtime" ON)
@@ -823,11 +821,6 @@ if(NOT WITH_CUDA_DYNLOAD)
   endif()
 endif()
 
-if(WITH_CYCLES_DEVICE_HIP)
-  # Currently HIP must be dynamically loaded, this may change in future toolkits
-  set(WITH_HIP_DYNLOAD ON)
-endif()
-
 #-----------------------------------------------------------------------------
 # Check check if submodules are cloned
 
@@ -1857,9 +1850,6 @@ elseif(WITH_CYCLES_STANDALONE)
   if(WITH_CUDA_DYNLOAD)
     add_subdirectory(extern/cuew)
   endif()
-  if(WITH_HIP_DYNLOAD)
-    add_subdirectory(extern/hipew)
-  endif()
   if(NOT WITH_SYSTEM_GLEW)
     add_subdirectory(extern/glew)
   endif()

extern/CMakeLists.txt

@@ -67,12 +67,9 @@ endif()
 
 if(WITH_CYCLES OR WITH_COMPOSITOR OR WITH_OPENSUBDIV)
   add_subdirectory(clew)
-  if((WITH_CYCLES_DEVICE_CUDA OR WITH_CYCLES_DEVICE_OPTIX) AND WITH_CUDA_DYNLOAD)
+  if(WITH_CUDA_DYNLOAD)
     add_subdirectory(cuew)
   endif()
-  if(WITH_CYCLES_DEVICE_HIP AND WITH_HIP_DYNLOAD)
-    add_subdirectory(hipew)
-  endif()
 endif()
 
 if(WITH_GHOST_X11 AND WITH_GHOST_XDND)

extern/hipew/CMakeLists.txt

@@ -1,39 +0,0 @@
-# ***** BEGIN GPL LICENSE BLOCK *****
-#
-# This program is free software; you can redistribute it and/or
-# modify it under the terms of the GNU General Public License
-# as published by the Free Software Foundation; either version 2
-# of the License, or (at your option) any later version.
-#
-# This program is distributed in the hope that it will be useful,
-# but WITHOUT ANY WARRANTY; without even the implied warranty of
-# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
-# GNU General Public License for more details.
-#
-# You should have received a copy of the GNU General Public License
-# along with this program; if not, write to the Free Software Foundation,
-# Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
-#
-# The Original Code is Copyright (C) 2021, Blender Foundation
-# All rights reserved.
-# ***** END GPL LICENSE BLOCK *****
-
-set(INC
-  .
-  include
-)
-
-set(INC_SYS
-
-)
-
-set(SRC
-  src/hipew.c
-
-  include/hipew.h
-)
-
-set(LIB
-)
-
-blender_add_lib(extern_hipew "${SRC}" "${INC}" "${INC_SYS}" "${LIB}")

extern/hipew/src/hipew.c

@@ -1,533 +0,0 @@
-/*
- * Copyright 2011-2021 Blender Foundation
- *
- * Licensed under the Apache License, Version 2.0 (the "License");
- * you may not use this file except in compliance with the License.
- * You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing, software
- * distributed under the License is distributed on an "AS IS" BASIS,
- * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
- * See the License for the specific language governing permissions and
- * limitations under the License
- */
-#ifdef _MSC_VER
-#  if _MSC_VER < 1900
-#    define snprintf _snprintf
-#  endif
-#  define popen _popen
-#  define pclose _pclose
-#  define _CRT_SECURE_NO_WARNINGS
-#endif
-
-#include <hipew.h>
-#include <assert.h>
-#include <stdio.h>
-#include <string.h>
-#include <sys/stat.h>
-
-#ifdef _WIN32
-#  define WIN32_LEAN_AND_MEAN
-#  define VC_EXTRALEAN
-#  include <windows.h>
-
-/* Utility macros. */
-
-typedef HMODULE DynamicLibrary;
-
-#  define dynamic_library_open(path)         LoadLibraryA(path)
-#  define dynamic_library_close(lib)         FreeLibrary(lib)
-#  define dynamic_library_find(lib, symbol)  GetProcAddress(lib, symbol)
-#else
-#  include <dlfcn.h>
-
-typedef void* DynamicLibrary;
-
-#  define dynamic_library_open(path)         dlopen(path, RTLD_NOW)
-#  define dynamic_library_close(lib)         dlclose(lib)
-#  define dynamic_library_find(lib, symbol)  dlsym(lib, symbol)
-#endif
-
-#define _LIBRARY_FIND_CHECKED(lib, name) \
-        name = (t##name *)dynamic_library_find(lib, #name); \
-        assert(name);
-
-#define _LIBRARY_FIND(lib, name) \
-        name = (t##name *)dynamic_library_find(lib, #name);
-
-#define HIP_LIBRARY_FIND_CHECKED(name) \
-        _LIBRARY_FIND_CHECKED(hip_lib, name)
-#define HIP_LIBRARY_FIND(name) _LIBRARY_FIND(hip_lib, name)
-
-
-static DynamicLibrary hip_lib;
-
-/* Function definitions. */
-thipGetErrorName *hipGetErrorName;
-thipInit *hipInit;
-thipDriverGetVersion *hipDriverGetVersion;
-thipGetDevice *hipGetDevice;
-thipGetDeviceCount *hipGetDeviceCount;
-thipDeviceGetName *hipDeviceGetName;
-thipDeviceGetAttribute *hipDeviceGetAttribute;
-thipDeviceComputeCapability *hipDeviceComputeCapability;
-thipDevicePrimaryCtxRetain *hipDevicePrimaryCtxRetain;
-thipDevicePrimaryCtxRelease *hipDevicePrimaryCtxRelease;
-thipDevicePrimaryCtxSetFlags *hipDevicePrimaryCtxSetFlags;
-thipDevicePrimaryCtxGetState *hipDevicePrimaryCtxGetState;
-thipDevicePrimaryCtxReset *hipDevicePrimaryCtxReset;
-thipCtxCreate *hipCtxCreate;
-thipCtxDestroy *hipCtxDestroy;
-thipCtxPushCurrent *hipCtxPushCurrent;
-thipCtxPopCurrent *hipCtxPopCurrent;
-thipCtxSetCurrent *hipCtxSetCurrent;
-thipCtxGetCurrent *hipCtxGetCurrent;
-thipCtxGetDevice *hipCtxGetDevice;
-thipCtxGetFlags *hipCtxGetFlags;
-thipCtxSynchronize *hipCtxSynchronize;
-thipDeviceSynchronize *hipDeviceSynchronize;
-thipCtxGetCacheConfig *hipCtxGetCacheConfig;
-thipCtxSetCacheConfig *hipCtxSetCacheConfig;
-thipCtxGetSharedMemConfig *hipCtxGetSharedMemConfig;
-thipCtxSetSharedMemConfig *hipCtxSetSharedMemConfig;
-thipCtxGetApiVersion *hipCtxGetApiVersion;
-thipModuleLoad *hipModuleLoad;
-thipModuleLoadData *hipModuleLoadData;
-thipModuleLoadDataEx *hipModuleLoadDataEx;
-thipModuleUnload *hipModuleUnload;
-thipModuleGetFunction *hipModuleGetFunction;
-thipModuleGetGlobal *hipModuleGetGlobal;
-thipModuleGetTexRef *hipModuleGetTexRef;
-thipMemGetInfo *hipMemGetInfo;
-thipMalloc *hipMalloc;
-thipMemAllocPitch *hipMemAllocPitch;
-thipFree *hipFree;
-thipMemGetAddressRange *hipMemGetAddressRange;
-thipHostMalloc *hipHostMalloc;
-thipHostFree *hipHostFree;
-thipHostGetDevicePointer *hipHostGetDevicePointer;
-thipHostGetFlags *hipHostGetFlags;
-thipMallocManaged *hipMallocManaged;
-thipDeviceGetByPCIBusId *hipDeviceGetByPCIBusId;
-thipDeviceGetPCIBusId *hipDeviceGetPCIBusId;
-thipMemcpyPeer *hipMemcpyPeer;
-thipMemcpyHtoD *hipMemcpyHtoD;
-thipMemcpyDtoH *hipMemcpyDtoH;
-thipMemcpyDtoD *hipMemcpyDtoD;
-thipDrvMemcpy2DUnaligned *hipDrvMemcpy2DUnaligned;
-thipMemcpyParam2D *hipMemcpyParam2D;
-thipDrvMemcpy3D *hipDrvMemcpy3D;
-thipMemcpyHtoDAsync *hipMemcpyHtoDAsync;
-thipMemcpyDtoHAsync *hipMemcpyDtoHAsync;
-thipMemcpyParam2DAsync *hipMemcpyParam2DAsync;
-thipDrvMemcpy3DAsync *hipDrvMemcpy3DAsync;
-thipMemsetD8 *hipMemsetD8;
-thipMemsetD16 *hipMemsetD16;
-thipMemsetD32 *hipMemsetD32;
-thipMemsetD8Async *hipMemsetD8Async;
-thipMemsetD16Async *hipMemsetD16Async;
-thipMemsetD32Async *hipMemsetD32Async;
-thipArrayCreate *hipArrayCreate;
-thipArrayDestroy *hipArrayDestroy;
-thipArray3DCreate *hipArray3DCreate;
-thipStreamCreateWithFlags *hipStreamCreateWithFlags;
-thipStreamCreateWithPriority *hipStreamCreateWithPriority;
-thipStreamGetPriority *hipStreamGetPriority;
-thipStreamGetFlags *hipStreamGetFlags;
-thipStreamWaitEvent *hipStreamWaitEvent;
-thipStreamAddCallback *hipStreamAddCallback;
-thipStreamQuery *hipStreamQuery;
-thipStreamSynchronize *hipStreamSynchronize;
-thipStreamDestroy *hipStreamDestroy;
-thipEventCreateWithFlags *hipEventCreateWithFlags;
-thipEventRecord *hipEventRecord;
-thipEventQuery *hipEventQuery;
-thipEventSynchronize *hipEventSynchronize;
-thipEventDestroy *hipEventDestroy;
-thipEventElapsedTime *hipEventElapsedTime;
-thipFuncGetAttribute *hipFuncGetAttribute;
-thipFuncSetCacheConfig *hipFuncSetCacheConfig;
-thipModuleLaunchKernel *hipModuleLaunchKernel;
-thipDrvOccupancyMaxActiveBlocksPerMultiprocessor *hipDrvOccupancyMaxActiveBlocksPerMultiprocessor;
-thipDrvOccupancyMaxActiveBlocksPerMultiprocessorWithFlags *hipDrvOccupancyMaxActiveBlocksPerMultiprocessorWithFlags;
-thipModuleOccupancyMaxPotentialBlockSize *hipModuleOccupancyMaxPotentialBlockSize;
-thipTexRefSetArray *hipTexRefSetArray;
-thipTexRefSetAddress *hipTexRefSetAddress;
-thipTexRefSetAddress2D *hipTexRefSetAddress2D;
-thipTexRefSetFormat *hipTexRefSetFormat;
-thipTexRefSetAddressMode *hipTexRefSetAddressMode;
-thipTexRefSetFilterMode *hipTexRefSetFilterMode;
-thipTexRefSetFlags *hipTexRefSetFlags;
-thipTexRefGetAddress *hipTexRefGetAddress;
-thipTexRefGetArray *hipTexRefGetArray;
-thipTexRefGetAddressMode *hipTexRefGetAddressMode;
-thipTexObjectCreate *hipTexObjectCreate;
-thipTexObjectDestroy *hipTexObjectDestroy;
-thipDeviceCanAccessPeer *hipDeviceCanAccessPeer;
-
-thipCtxEnablePeerAccess *hipCtxEnablePeerAccess;
-thipCtxDisablePeerAccess *hipCtxDisablePeerAccess;
-thipDeviceGetP2PAttribute *hipDeviceGetP2PAttribute;
-thipGraphicsUnregisterResource *hipGraphicsUnregisterResource;
-thipGraphicsMapResources *hipGraphicsMapResources;
-thipGraphicsUnmapResources *hipGraphicsUnmapResources;
-thipGraphicsResourceGetMappedPointer *hipGraphicsResourceGetMappedPointer;
-
-thipGraphicsGLRegisterBuffer *hipGraphicsGLRegisterBuffer;
-thipGLGetDevices *hipGLGetDevices;
-
-
-
-static DynamicLibrary dynamic_library_open_find(const char **paths) {
-  int i = 0;
-  while (paths[i] != NULL) {
-      DynamicLibrary lib = dynamic_library_open(paths[i]);
-      if (lib != NULL) {
-        return lib;
-      }
-      ++i;
-  }
-  return NULL;
-}
-
-/* Implementation function. */
-static void hipewHipExit(void) {
-  if (hip_lib != NULL) {
-    /*  Ignore errors. */
-    dynamic_library_close(hip_lib);
-    hip_lib = NULL;
-  }
-}
-
-static int hipewHipInit(void) {
-  /* Library paths. */
-#ifdef _WIN32
-  /* Expected in c:/windows/system or similar, no path needed. */
-  const char *hip_paths[] = {"amdhip64.dll", NULL};
-#elif defined(__APPLE__)
-  /* Default installation path. */
-  const char *hip_paths[] = {"", NULL};
-#else
-  const char *hip_paths[] = {"/opt/rocm/hip/lib/libamdhip64.so", NULL};
-#endif
-  static int initialized = 0;
-  static int result = 0;
-  int error, driver_version;
-
-  if (initialized) {
-    return result;
-  }
-
-  initialized = 1;
-
-  error = atexit(hipewHipExit);
-  if (error) {
-    result = HIPEW_ERROR_ATEXIT_FAILED;
-    return result;
-  }
-
-  /* Load library. */
-  hip_lib = dynamic_library_open_find(hip_paths);
-
-  if (hip_lib == NULL) {
-    result = HIPEW_ERROR_OPEN_FAILED;
-    return result;
-  }
-
-  /* Fetch all function pointers. */
-  HIP_LIBRARY_FIND_CHECKED(hipGetErrorName);
-  HIP_LIBRARY_FIND_CHECKED(hipInit);
-  HIP_LIBRARY_FIND_CHECKED(hipDriverGetVersion);
-  HIP_LIBRARY_FIND_CHECKED(hipGetDevice);
-  HIP_LIBRARY_FIND_CHECKED(hipGetDeviceCount);
-  HIP_LIBRARY_FIND_CHECKED(hipDeviceGetName);
-  HIP_LIBRARY_FIND_CHECKED(hipDeviceGetAttribute);
-  HIP_LIBRARY_FIND_CHECKED(hipDeviceComputeCapability);
-  HIP_LIBRARY_FIND_CHECKED(hipDevicePrimaryCtxRetain);
-  HIP_LIBRARY_FIND_CHECKED(hipDevicePrimaryCtxRelease);
-  HIP_LIBRARY_FIND_CHECKED(hipDevicePrimaryCtxSetFlags);
-  HIP_LIBRARY_FIND_CHECKED(hipDevicePrimaryCtxGetState);
-  HIP_LIBRARY_FIND_CHECKED(hipDevicePrimaryCtxReset);
-  HIP_LIBRARY_FIND_CHECKED(hipCtxCreate);
-  HIP_LIBRARY_FIND_CHECKED(hipCtxDestroy);
-  HIP_LIBRARY_FIND_CHECKED(hipCtxPushCurrent);
-  HIP_LIBRARY_FIND_CHECKED(hipCtxPopCurrent);
-  HIP_LIBRARY_FIND_CHECKED(hipCtxSetCurrent);
-  HIP_LIBRARY_FIND_CHECKED(hipCtxGetCurrent);
-  HIP_LIBRARY_FIND_CHECKED(hipCtxGetDevice);
-  HIP_LIBRARY_FIND_CHECKED(hipCtxGetFlags);
-  HIP_LIBRARY_FIND_CHECKED(hipCtxSynchronize);
-  HIP_LIBRARY_FIND_CHECKED(hipDeviceSynchronize);
-  HIP_LIBRARY_FIND_CHECKED(hipCtxGetCacheConfig);
-  HIP_LIBRARY_FIND_CHECKED(hipCtxSetCacheConfig);
-  HIP_LIBRARY_FIND_CHECKED(hipCtxGetSharedMemConfig);
-  HIP_LIBRARY_FIND_CHECKED(hipCtxSetSharedMemConfig);
-  HIP_LIBRARY_FIND_CHECKED(hipCtxGetApiVersion);
-  HIP_LIBRARY_FIND_CHECKED(hipModuleLoad);
-  HIP_LIBRARY_FIND_CHECKED(hipModuleLoadData);
-  HIP_LIBRARY_FIND_CHECKED(hipModuleLoadDataEx);
-  HIP_LIBRARY_FIND_CHECKED(hipModuleUnload);
-  HIP_LIBRARY_FIND_CHECKED(hipModuleGetFunction);
-  HIP_LIBRARY_FIND_CHECKED(hipModuleGetGlobal);
-  HIP_LIBRARY_FIND_CHECKED(hipModuleGetTexRef);
-  HIP_LIBRARY_FIND_CHECKED(hipMemGetInfo);
-  HIP_LIBRARY_FIND_CHECKED(hipMalloc);
-  HIP_LIBRARY_FIND_CHECKED(hipMemAllocPitch);
-  HIP_LIBRARY_FIND_CHECKED(hipFree);
-  HIP_LIBRARY_FIND_CHECKED(hipMemGetAddressRange);
-  HIP_LIBRARY_FIND_CHECKED(hipHostMalloc);
-  HIP_LIBRARY_FIND_CHECKED(hipHostFree);
-  HIP_LIBRARY_FIND_CHECKED(hipHostGetDevicePointer);
-  HIP_LIBRARY_FIND_CHECKED(hipHostGetFlags);
-  HIP_LIBRARY_FIND_CHECKED(hipMallocManaged);
-  HIP_LIBRARY_FIND_CHECKED(hipDeviceGetByPCIBusId);
-  HIP_LIBRARY_FIND_CHECKED(hipDeviceGetPCIBusId);
-  HIP_LIBRARY_FIND_CHECKED(hipMemcpyPeer);
-  HIP_LIBRARY_FIND_CHECKED(hipMemcpyHtoD);
-  HIP_LIBRARY_FIND_CHECKED(hipMemcpyDtoH);
-  HIP_LIBRARY_FIND_CHECKED(hipMemcpyDtoD);
-  HIP_LIBRARY_FIND_CHECKED(hipMemcpyParam2D);
-  HIP_LIBRARY_FIND_CHECKED(hipDrvMemcpy3D);
-  HIP_LIBRARY_FIND_CHECKED(hipMemcpyHtoDAsync);
-  HIP_LIBRARY_FIND_CHECKED(hipMemcpyDtoHAsync);
-  HIP_LIBRARY_FIND_CHECKED(hipDrvMemcpy2DUnaligned);
-  HIP_LIBRARY_FIND_CHECKED(hipMemcpyParam2DAsync);
-  HIP_LIBRARY_FIND_CHECKED(hipDrvMemcpy3DAsync);
-  HIP_LIBRARY_FIND_CHECKED(hipMemsetD8);
-  HIP_LIBRARY_FIND_CHECKED(hipMemsetD16);
-  HIP_LIBRARY_FIND_CHECKED(hipMemsetD32);
-  HIP_LIBRARY_FIND_CHECKED(hipMemsetD8Async);
-  HIP_LIBRARY_FIND_CHECKED(hipMemsetD16Async);
-  HIP_LIBRARY_FIND_CHECKED(hipMemsetD32Async);
-  HIP_LIBRARY_FIND_CHECKED(hipArrayCreate);
-  HIP_LIBRARY_FIND_CHECKED(hipArrayDestroy);
-  HIP_LIBRARY_FIND_CHECKED(hipArray3DCreate);
-  HIP_LIBRARY_FIND_CHECKED(hipStreamCreateWithFlags);
-  HIP_LIBRARY_FIND_CHECKED(hipStreamCreateWithPriority);
-  HIP_LIBRARY_FIND_CHECKED(hipStreamGetPriority);
-  HIP_LIBRARY_FIND_CHECKED(hipStreamGetFlags);
-  HIP_LIBRARY_FIND_CHECKED(hipStreamWaitEvent);
-  HIP_LIBRARY_FIND_CHECKED(hipStreamAddCallback);
-  HIP_LIBRARY_FIND_CHECKED(hipStreamQuery);
-  HIP_LIBRARY_FIND_CHECKED(hipStreamSynchronize);
-  HIP_LIBRARY_FIND_CHECKED(hipStreamDestroy);
-  HIP_LIBRARY_FIND_CHECKED(hipEventCreateWithFlags);
-  HIP_LIBRARY_FIND_CHECKED(hipEventRecord);
-  HIP_LIBRARY_FIND_CHECKED(hipEventQuery);
-  HIP_LIBRARY_FIND_CHECKED(hipEventSynchronize);
-  HIP_LIBRARY_FIND_CHECKED(hipEventDestroy);
-  HIP_LIBRARY_FIND_CHECKED(hipEventElapsedTime);
-  HIP_LIBRARY_FIND_CHECKED(hipFuncGetAttribute);
-  HIP_LIBRARY_FIND_CHECKED(hipFuncSetCacheConfig);
-  HIP_LIBRARY_FIND_CHECKED(hipModuleLaunchKernel);
-  HIP_LIBRARY_FIND_CHECKED(hipModuleOccupancyMaxPotentialBlockSize);
-  HIP_LIBRARY_FIND_CHECKED(hipTexRefSetArray);
-  HIP_LIBRARY_FIND_CHECKED(hipTexRefSetAddress);
-  HIP_LIBRARY_FIND_CHECKED(hipTexRefSetAddress2D);
-  HIP_LIBRARY_FIND_CHECKED(hipTexRefSetFormat);
-  HIP_LIBRARY_FIND_CHECKED(hipTexRefSetAddressMode);
-  HIP_LIBRARY_FIND_CHECKED(hipTexRefSetFilterMode);
-  HIP_LIBRARY_FIND_CHECKED(hipTexRefSetFlags);
-  HIP_LIBRARY_FIND_CHECKED(hipTexRefGetAddress);
-  HIP_LIBRARY_FIND_CHECKED(hipTexRefGetAddressMode);
-  HIP_LIBRARY_FIND_CHECKED(hipTexObjectCreate);
-  HIP_LIBRARY_FIND_CHECKED(hipTexObjectDestroy);
-  HIP_LIBRARY_FIND_CHECKED(hipDeviceCanAccessPeer);
-  HIP_LIBRARY_FIND_CHECKED(hipCtxEnablePeerAccess);
-  HIP_LIBRARY_FIND_CHECKED(hipCtxDisablePeerAccess);
-  HIP_LIBRARY_FIND_CHECKED(hipDeviceGetP2PAttribute);
-#ifdef _WIN32
-  HIP_LIBRARY_FIND_CHECKED(hipGraphicsUnregisterResource);
-  HIP_LIBRARY_FIND_CHECKED(hipGraphicsMapResources);
-  HIP_LIBRARY_FIND_CHECKED(hipGraphicsUnmapResources);
-  HIP_LIBRARY_FIND_CHECKED(hipGraphicsResourceGetMappedPointer);
-  HIP_LIBRARY_FIND_CHECKED(hipGraphicsGLRegisterBuffer);
-  HIP_LIBRARY_FIND_CHECKED(hipGLGetDevices);
-#endif
-  result = HIPEW_SUCCESS;
-  return result;
-}
-
-
-
-int hipewInit(hipuint32_t flags) {
-  int result = HIPEW_SUCCESS;
-
-  if (flags & HIPEW_INIT_HIP) {
-    result = hipewHipInit();
-    if (result != HIPEW_SUCCESS) {
-      return result;
-    }
-  }
-
-  return result;
-}
-
-
-const char *hipewErrorString(hipError_t result) {
-  switch (result) {
-    case hipSuccess: return "No errors";
-    case hipErrorInvalidValue: return "Invalid value";
-    case hipErrorOutOfMemory: return "Out of memory";
-    case hipErrorNotInitialized: return "Driver not initialized";
-    case hipErrorDeinitialized: return "Driver deinitialized";
-    case hipErrorProfilerDisabled: return "Profiler disabled";
-    case hipErrorProfilerNotInitialized: return "Profiler not initialized";
-    case hipErrorProfilerAlreadyStarted: return "Profiler already started";
-    case hipErrorProfilerAlreadyStopped: return "Profiler already stopped";
-    case hipErrorNoDevice: return "No HIP-capable device available";
-    case hipErrorInvalidDevice: return "Invalid device";
-    case hipErrorInvalidImage: return "Invalid kernel image";
-    case hipErrorInvalidContext: return "Invalid context";
-    case hipErrorContextAlreadyCurrent: return "Context already current";
-    case hipErrorMapFailed: return "Map failed";
-    case hipErrorUnmapFailed: return "Unmap failed";
-    case hipErrorArrayIsMapped: return "Array is mapped";
-    case hipErrorAlreadyMapped: return "Already mapped";
-    case hipErrorNoBinaryForGpu: return "No binary for GPU";
-    case hipErrorAlreadyAcquired: return "Already acquired";
-    case hipErrorNotMapped: return "Not mapped";
-    case hipErrorNotMappedAsArray: return "Mapped resource not available for access as an array";
-    case hipErrorNotMappedAsPointer: return "Mapped resource not available for access as a pointer";
-    case hipErrorECCNotCorrectable: return "Uncorrectable ECC error detected";
-    case hipErrorUnsupportedLimit: return "hipLimit_t not supported by device";
-    case hipErrorContextAlreadyInUse: return "Context already in use";
-    case hipErrorPeerAccessUnsupported: return "Peer access unsupported";
-    case hipErrorInvalidKernelFile: return "Invalid ptx";
-    case hipErrorInvalidGraphicsContext: return "Invalid graphics context";
-    case hipErrorInvalidSource: return "Invalid source";
-    case hipErrorFileNotFound: return "File not found";
-    case hipErrorSharedObjectSymbolNotFound: return "Link to a shared object failed to resolve";
-    case hipErrorSharedObjectInitFailed: return "Shared object initialization failed";
-    case hipErrorOperatingSystem: return "Operating system";
-    case hipErrorInvalidHandle: return "Invalid handle";
-    case hipErrorNotFound: return "Not found";
-    case hipErrorNotReady: return "HIP not ready";
-    case hipErrorIllegalAddress: return "Illegal address";
-    case hipErrorLaunchOutOfResources: return "Launch exceeded resources";
-    case hipErrorLaunchTimeOut: return "Launch exceeded timeout";
-    case hipErrorPeerAccessAlreadyEnabled: return "Peer access already enabled";
-    case hipErrorPeerAccessNotEnabled: return "Peer access not enabled";
-    case hipErrorSetOnActiveProcess: return "Primary context active";
-    case hipErrorAssert: return "Assert";
-    case hipErrorHostMemoryAlreadyRegistered: return "Host memory already registered";
-    case hipErrorHostMemoryNotRegistered: return "Host memory not registered";
-    case hipErrorLaunchFailure: return "Launch failed";
-    case hipErrorCooperativeLaunchTooLarge: return "Cooperative launch too large";
-    case hipErrorNotSupported: return "Not supported";
-    case hipErrorUnknown: return "Unknown error";
-    default: return "Unknown HIP error value";
-  }
-}
-
-static void path_join(const char *path1,
-                      const char *path2,
-                      int maxlen,
-                      char *result) {
-#if defined(WIN32) || defined(_WIN32)
-  const char separator = '\\';
-#else
-  const char separator = '/';
-#endif
-  int n = snprintf(result, maxlen, "%s%c%s", path1, separator, path2);
-  if (n != -1 && n < maxlen) {
-    result[n] = '\0';
-  }
-  else {
-    result[maxlen - 1] = '\0';
-  }
-}
-
-static int path_exists(const char *path) {
-  struct stat st;
-  if (stat(path, &st)) {
-    return 0;
-  }
-  return 1;
-}
-
-const char *hipewCompilerPath(void) {
-    #ifdef _WIN32
-    const char *hipPath = getenv("HIP_ROCCLR_HOME");
-    const char *windowsCommand = "perl ";
-    const char *executable = "bin/hipcc";
-
-    static char hipcc[65536];
-    static char finalCommand[65536];
-    if(hipPath) {
-      path_join(hipPath, executable, sizeof(hipcc), hipcc);
-      if(path_exists(hipcc)) {
-        snprintf(finalCommand, sizeof(hipcc), "%s %s", windowsCommand, hipcc);
-        return finalCommand;
-      } else {
-        printf("Could not find hipcc. Make sure HIP_ROCCLR_HOME points to the directory holding /bin/hipcc");
-      }
-    }
-    #else
-    const char *hipPath =  "opt/rocm/hip/bin";
-    const char *executable = "hipcc";
-
-    static char hipcc[65536];
-    if(hipPath) {
-      path_join(hipPath, executable, sizeof(hipcc), hipcc);
-      if(path_exists(hipcc)){
-        return hipcc;
-      }
-    }
-    #endif
-
-  {
-#ifdef _WIN32
-    FILE *handle = popen("where hipcc", "r");
-#else
-    FILE *handle = popen("which hipcc", "r");
-#endif
-    if (handle) {
-      char buffer[4096] = {0};
-      int len = fread(buffer, 1, sizeof(buffer) - 1, handle);
-      buffer[len] = '\0';
-      pclose(handle);
-      if (buffer[0]) {
-        return "hipcc";
-      }
-    }
-  }
-
-  return NULL;
-}
-
-int hipewCompilerVersion(void) {
-  const char *path = hipewCompilerPath();
-  const char *marker = "Hip compilation tools, release ";
-  FILE *pipe;
-  int major, minor;
-  char *versionstr;
-  char buf[128];
-  char output[65536] = "\0";
-  char command[65536] = "\0";
-
-  if (path == NULL) {
-    return 0;
-  }
-
-  /* get --version output */
-  strcat(command, "\"");
-  strncat(command, path, sizeof(command) - 1);
-  strncat(command, "\" --version", sizeof(command) - strlen(path) - 1);
-  pipe = popen(command, "r");
-  if (!pipe) {
-    fprintf(stderr, "HIP: failed to run compiler to retrieve version");
-    return 0;
-  }
-
-  while (!feof(pipe)) {
-    if (fgets(buf, sizeof(buf), pipe) != NULL) {
-      strncat(output, buf, sizeof(output) - strlen(output) - 1);
-    }
-  }
-
-  pclose(pipe);
-  return 40;
-}

intern/cycles/CMakeLists.txt

@@ -297,7 +297,6 @@ endif()
 
 if(WITH_CYCLES_STANDALONE)
   set(WITH_CYCLES_DEVICE_CUDA TRUE)
-  set(WITH_CYCLES_DEVICE_HIP TRUE)
 endif()
 # TODO(sergey): Consider removing it, only causes confusion in interface.
 set(WITH_CYCLES_DEVICE_MULTI TRUE)

intern/cycles/blender/CMakeLists.txt

@@ -95,9 +95,6 @@ set(ADDON_FILES
 
 add_definitions(${GL_DEFINITIONS})
 
-if(WITH_CYCLES_DEVICE_HIP)
-  add_definitions(-DWITH_HIP)
-endif()
 if(WITH_MOD_FLUID)
   add_definitions(-DWITH_FLUID)
 endif()

intern/cycles/blender/addon/engine.py

@@ -28,7 +28,7 @@ def _configure_argument_parser():
                         action='store_true')
     parser.add_argument("--cycles-device",
                         help="Set the device to use for Cycles, overriding user preferences and the scene setting."
-                             "Valid options are 'CPU', 'CUDA', 'OPTIX', or 'HIP'"
+                             "Valid options are 'CPU', 'CUDA' or 'OPTIX'."
                              "Additionally, you can append '+CPU' to any GPU type for hybrid rendering.",
                         default=None)
     return parser

intern/cycles/blender/addon/properties.py

@@ -111,7 +111,6 @@ enum_device_type = (
     ('CPU', "CPU", "CPU", 0),
     ('CUDA', "CUDA", "CUDA", 1),
     ('OPTIX', "OptiX", "OptiX", 3),
-    ("HIP", "HIP", "HIP", 4)
 )
 
 enum_texture_limit = (
@@ -124,7 +123,7 @@ enum_texture_limit = (
     ('4096', "4096", "Limit texture size to 4096 pixels", 6),
     ('8192', "8192", "Limit texture size to 8192 pixels", 7),
 )
-
+ 
 # NOTE: Identifiers are expected to be an upper case version of identifiers from  `Pass::get_type_enum()`
 enum_view3d_shading_render_pass = (
     ('', "General", ""),
@@ -1267,16 +1266,12 @@ class CyclesPreferences(bpy.types.AddonPreferences):
 
     def get_device_types(self, context):
         import _cycles
-        has_cuda, has_optix, has_hip = _cycles.get_device_types()
-
+        has_cuda, has_optix = _cycles.get_device_types()
         list = [('NONE', "None", "Don't use compute device", 0)]
         if has_cuda:
             list.append(('CUDA', "CUDA", "Use CUDA for GPU acceleration", 1))
         if has_optix:
             list.append(('OPTIX', "OptiX", "Use OptiX for GPU acceleration", 3))
-        if has_hip:
-            list.append(('HIP', "HIP", "Use HIP for GPU acceleration", 4))
-
         return list
 
     compute_device_type: EnumProperty(
@@ -1301,7 +1296,7 @@ class CyclesPreferences(bpy.types.AddonPreferences):
 
     def update_device_entries(self, device_list):
         for device in device_list:
-            if not device[1] in {'CUDA', 'OPTIX', 'CPU', 'HIP'}:
+            if not device[1] in {'CUDA', 'OPTIX', 'CPU'}:
                 continue
             # Try to find existing Device entry
             entry = self.find_existing_device_entry(device)
@@ -1335,7 +1330,7 @@ class CyclesPreferences(bpy.types.AddonPreferences):
             elif entry.type == 'CPU':
                 cpu_devices.append(entry)
         # Extend all GPU devices with CPU.
-        if compute_device_type != 'CPU' and compute_device_type != 'HIP':
+        if compute_device_type != 'CPU':
             devices.extend(cpu_devices)
         return devices
 
@@ -1345,7 +1340,7 @@ class CyclesPreferences(bpy.types.AddonPreferences):
         import _cycles
         # Ensure `self.devices` is not re-allocated when the second call to
         # get_devices_for_type is made, freeing items from the first list.
-        for device_type in ('CUDA', 'OPTIX', 'HIP'):
+        for device_type in ('CUDA', 'OPTIX', 'OPENCL'):
             self.update_device_entries(_cycles.available_devices(device_type))
 
     # Deprecated: use refresh_devices instead.

intern/cycles/blender/addon/ui.py

@@ -99,11 +99,6 @@ def use_cuda(context):
     return (get_device_type(context) == 'CUDA' and cscene.device == 'GPU')
 
 
-def use_hip(context):
-    cscene = context.scene.cycles
-
-    return (get_device_type(context) == 'HIP' and cscene.device == 'GPU')
-
 def use_optix(context):
     cscene = context.scene.cycles
 

intern/cycles/blender/blender_device.cpp

@@ -26,7 +26,6 @@ enum ComputeDevice {
   COMPUTE_DEVICE_CPU = 0,
   COMPUTE_DEVICE_CUDA = 1,
   COMPUTE_DEVICE_OPTIX = 3,
-  COMPUTE_DEVICE_HIP = 4,
 
   COMPUTE_DEVICE_NUM
 };
@@ -82,9 +81,6 @@ DeviceInfo blender_device_info(BL::Preferences &b_preferences, BL::Scene &b_scen
       else if (compute_device == COMPUTE_DEVICE_OPTIX) {
         mask |= DEVICE_MASK_OPTIX;
       }
-      else if (compute_device == COMPUTE_DEVICE_HIP) {
-        mask |= DEVICE_MASK_HIP;
-      }
       vector<DeviceInfo> devices = Device::available_devices(mask);
 
       /* Match device preferences and available devices. */

intern/cycles/blender/blender_geometry.cpp

@@ -80,10 +80,8 @@ Geometry *BlenderSync::sync_geometry(BL::Depsgraph &b_depsgraph,
 {
   /* Test if we can instance or if the object is modified. */
   Geometry::Type geom_type = determine_geom_type(b_ob_info, use_particle_hair);
-  BL::ID b_key_id = (b_ob_info.is_real_object_data() &&
-                     BKE_object_is_modified(b_ob_info.real_object)) ?
-                        b_ob_info.real_object :
-                        b_ob_info.object_data;
+  BL::ID b_key_id = (BKE_object_is_modified(b_ob_info.real_object)) ? b_ob_info.real_object :
+                                                                      b_ob_info.object_data;
   GeometryKey key(b_key_id.ptr.data, geom_type);
 
   /* Find shader indices. */

intern/cycles/blender/blender_gpu_display.cpp

@@ -485,6 +485,12 @@ void BlenderGPUDisplay::do_draw(const GPUDisplayParams &params)
   /* See do_update_begin() for why no locking is required here. */
   const bool transparent = true;  // TODO(sergey): Derive this from Film.
 
+  if (texture_.need_clear) {
+    /* Texture is requested to be cleared and was not yet cleared.
+     * Do early return which should be equivalent of drawing all-zero texture. */
+    return;
+  }
+
   if (!gl_draw_resources_ensure()) {
     return;
   }
@@ -493,16 +499,6 @@ void BlenderGPUDisplay::do_draw(const GPUDisplayParams &params)
     gl_context_mutex_.lock();
   }
 
-  if (texture_.need_clear) {
-    /* Texture is requested to be cleared and was not yet cleared.
-     *
-     * Do early return which should be equivalent of drawing all-zero texture.
-     * Watchout for the lock though so that the clear happening during update is properly
-     * synchronized here. */
-    gl_context_mutex_.unlock();
-    return;
-  }
-
   if (gl_upload_sync_) {
     glWaitSync((GLsync)gl_upload_sync_, 0, GL_TIMEOUT_IGNORED);
   }

intern/cycles/blender/blender_python.cpp

@@ -911,16 +911,14 @@ static PyObject *enable_print_stats_func(PyObject * /*self*/, PyObject * /*args*
 static PyObject *get_device_types_func(PyObject * /*self*/, PyObject * /*args*/)
 {
   vector<DeviceType> device_types = Device::available_types();
-  bool has_cuda = false, has_optix = false, has_hip = false;
+  bool has_cuda = false, has_optix = false;
   foreach (DeviceType device_type, device_types) {
     has_cuda |= (device_type == DEVICE_CUDA);
     has_optix |= (device_type == DEVICE_OPTIX);
-    has_hip |= (device_type == DEVICE_HIP);
   }
-  PyObject *list = PyTuple_New(3);
+  PyObject *list = PyTuple_New(2);
   PyTuple_SET_ITEM(list, 0, PyBool_FromLong(has_cuda));
   PyTuple_SET_ITEM(list, 1, PyBool_FromLong(has_optix));
-  PyTuple_SET_ITEM(list, 2, PyBool_FromLong(has_hip));
   return list;
 }
 
@@ -946,9 +944,6 @@ static PyObject *set_device_override_func(PyObject * /*self*/, PyObject *arg)
   else if (override == "OPTIX") {
     BlenderSession::device_override = DEVICE_MASK_OPTIX;
   }
-  else if (override == "HIP") {
-    BlenderSession::device_override = DEVICE_MASK_HIP;
-  }
   else {
     printf("\nError: %s is not a valid Cycles device.\n", override.c_str());
     Py_RETURN_FALSE;

intern/cycles/blender/blender_session.cpp

@@ -71,8 +71,7 @@ BlenderSession::BlenderSession(BL::RenderEngine &b_engine,
       width(0),
       height(0),
       preview_osl(preview_osl),
-      python_thread_state(NULL),
-      use_developer_ui(false)
+      python_thread_state(NULL)
 {
   /* offline render */
   background = true;
@@ -589,12 +588,6 @@ void BlenderSession::render_frame_finish()
 
   for (string_view filename : full_buffer_files_) {
     session->process_full_buffer_from_disk(filename);
-    if (check_and_report_session_error()) {
-      break;
-    }
-  }
-
-  for (string_view filename : full_buffer_files_) {
     path_remove(filename);
   }
 
@@ -1042,27 +1035,20 @@ void BlenderSession::update_status_progress()
     last_progress = progress;
   }
 
-  check_and_report_session_error();
-}
-
-bool BlenderSession::check_and_report_session_error()
-{
-  if (!session->progress.get_error()) {
-    return false;
+  if (session->progress.get_error()) {
+    string error = session->progress.get_error_message();
+    if (error != last_error) {
+      /* TODO(sergey): Currently C++ RNA API doesn't let us to
+       * use mnemonic name for the variable. Would be nice to
+       * have this figured out.
+       *
+       * For until then, 1 << 5 means RPT_ERROR.
+       */
+      b_engine.report(1 << 5, error.c_str());
+      b_engine.error_set(error.c_str());
+      last_error = error;
+    }
   }
-
-  const string error = session->progress.get_error_message();
-  if (error != last_error) {
-    /* TODO(sergey): Currently C++ RNA API doesn't let us to use mnemonic name for the variable.
-     * Would be nice to have this figured out.
-     *
-     * For until then, 1 << 5 means RPT_ERROR. */
-    b_engine.report(1 << 5, error.c_str());
-    b_engine.error_set(error.c_str());
-    last_error = error;
-  }
-
-  return true;
 }
 
 void BlenderSession::tag_update()

intern/cycles/blender/blender_session.h

@@ -110,7 +110,8 @@ class BlenderSession {
   BL::RenderSettings b_render;
   BL::Depsgraph b_depsgraph;
   /* NOTE: Blender's scene might become invalid after call
-   * #free_blender_memory_if_possible(). */
+   * free_blender_memory_if_possible().
+   */
   BL::Scene b_scene;
   BL::SpaceView3D b_v3d;
   BL::RegionView3D b_rv3d;
@@ -146,11 +147,6 @@ class BlenderSession {
  protected:
   void stamp_view_layer_metadata(Scene *scene, const string &view_layer_name);
 
-  /* Check whether session error happened.
-   * If so, it is reported to the render engine and true is returned.
-   * Otherwise false is returned. */
-  bool check_and_report_session_error();
-
   void builtin_images_load();
 
   /* Is used after each render layer synchronization is done with the goal

intern/cycles/blender/blender_util.h

@@ -90,27 +90,26 @@ static inline BL::Mesh object_to_mesh(BL::BlendData & /*data*/,
   }
 #endif
 
-  BL::Mesh mesh = (b_ob_info.object_data.is_a(&RNA_Mesh)) ? BL::Mesh(b_ob_info.object_data) :
-                                                            BL::Mesh(PointerRNA_NULL);
+  BL::Mesh mesh(PointerRNA_NULL);
+  if (b_ob_info.object_data.is_a(&RNA_Mesh)) {
+    /* TODO: calc_undeformed is not used. */
+    mesh = BL::Mesh(b_ob_info.object_data);
 
-  if (b_ob_info.is_real_object_data()) {
-    if (mesh) {
-      /* Make a copy to split faces if we use autosmooth, otherwise not needed.
-       * Also in edit mode do we need to make a copy, to ensure data layers like
-       * UV are not empty. */
-      if (mesh.is_editmode() ||
-          (mesh.use_auto_smooth() && subdivision_type == Mesh::SUBDIVISION_NONE)) {
-        BL::Depsgraph depsgraph(PointerRNA_NULL);
-        mesh = b_ob_info.real_object.to_mesh(false, depsgraph);
-      }
-    }
-    else {
+    /* Make a copy to split faces if we use autosmooth, otherwise not needed.
+     * Also in edit mode do we need to make a copy, to ensure data layers like
+     * UV are not empty. */
+    if (mesh.is_editmode() ||
+        (mesh.use_auto_smooth() && subdivision_type == Mesh::SUBDIVISION_NONE)) {
       BL::Depsgraph depsgraph(PointerRNA_NULL);
+      assert(b_ob_info.is_real_object_data());
       mesh = b_ob_info.real_object.to_mesh(false, depsgraph);
     }
   }
   else {
-    /* TODO: what to do about non-mesh geometry instances? */
+    BL::Depsgraph depsgraph(PointerRNA_NULL);
+    if (b_ob_info.is_real_object_data()) {
+      mesh = b_ob_info.real_object.to_mesh(false, depsgraph);
+    }
   }
 
 #if 0

intern/cycles/bvh/bvh_embree.cpp

@@ -213,7 +213,7 @@ static void rtc_filter_occluded_func(const RTCFilterFunctionNArguments *args)
       if (ctx->num_hits < ctx->max_hits) {
         Intersection current_isect;
         kernel_embree_convert_hit(kg, ray, hit, &current_isect);
-        for (size_t i = 0; i < ctx->num_hits; ++i) {
+        for (size_t i = 0; i < ctx->max_hits; ++i) {
           if (current_isect.object == ctx->isect_s[i].object &&
               current_isect.prim == ctx->isect_s[i].prim && current_isect.t == ctx->isect_s[i].t) {
             /* This intersection was already recorded, skip it. */

intern/cycles/cmake/external_libs.cmake

@@ -532,13 +532,4 @@ if(WITH_CYCLES_CUDA_BINARIES OR NOT WITH_CUDA_DYNLOAD)
   endif()
 endif()
 
-
-###########################################################################
-# HIP
-###########################################################################
-
-if(NOT WITH_HIP_DYNLOAD)
-  set(WITH_HIP_DYNLOAD ON)
-endif()
-
 unset(_cycles_lib_dir)

intern/cycles/cmake/macros.cmake

@@ -156,16 +156,10 @@ macro(cycles_target_link_libraries target)
     ${PLATFORM_LINKLIBS}
   )
 
-  if(WITH_CYCLES_DEVICE_CUDA OR WITH_CYCLES_DEVICE_OPTIX)
-    if(WITH_CUDA_DYNLOAD)
-      target_link_libraries(${target} extern_cuew)
-    else()
-      target_link_libraries(${target} ${CUDA_CUDA_LIBRARY})
-    endif()
-  endif()
-
-  if(WITH_CYCLES_DEVICE_HIP AND WITH_HIP_DYNLOAD)
-    target_link_libraries(${target} extern_hipew)
+  if(WITH_CUDA_DYNLOAD)
+    target_link_libraries(${target} extern_cuew)
+  else()
+    target_link_libraries(${target} ${CUDA_CUDA_LIBRARY})
   endif()
 
   if(CYCLES_STANDALONE_REPOSITORY)

intern/cycles/device/CMakeLists.txt

@@ -22,25 +22,16 @@ set(INC_SYS
   ../../../extern/clew/include
 )
 
-if(WITH_CYCLES_DEVICE_OPTIX OR WITH_CYCLES_DEVICE_CUDA)
-  if(WITH_CUDA_DYNLOAD)
-    list(APPEND INC
-      ../../../extern/cuew/include
-    )
-    add_definitions(-DWITH_CUDA_DYNLOAD)
-  else()
-    list(APPEND INC_SYS
-      ${CUDA_TOOLKIT_INCLUDE}
-    )
-    add_definitions(-DCYCLES_CUDA_NVCC_EXECUTABLE="${CUDA_NVCC_EXECUTABLE}")
-  endif()
-endif()
-
-if(WITH_CYCLES_DEVICE_HIP AND WITH_HIP_DYNLOAD)
+if(WITH_CUDA_DYNLOAD)
   list(APPEND INC
-    ../../../extern/hipew/include
+    ../../../extern/cuew/include
   )
-  add_definitions(-DWITH_HIP_DYNLOAD)
+  add_definitions(-DWITH_CUDA_DYNLOAD)
+else()
+  list(APPEND INC_SYS
+    ${CUDA_TOOLKIT_INCLUDE}
+  )
+  add_definitions(-DCYCLES_CUDA_NVCC_EXECUTABLE="${CUDA_NVCC_EXECUTABLE}")
 endif()
 
 set(SRC
@@ -79,21 +70,6 @@ set(SRC_CUDA
   cuda/util.h
 )
 
-set(SRC_HIP
-  hip/device.cpp
-  hip/device.h
-  hip/device_impl.cpp
-  hip/device_impl.h
-  hip/graphics_interop.cpp
-  hip/graphics_interop.h
-  hip/kernel.cpp
-  hip/kernel.h
-  hip/queue.cpp
-  hip/queue.h
-  hip/util.cpp
-  hip/util.h
-)
-
 set(SRC_DUMMY
   dummy/device.cpp
   dummy/device.h
@@ -129,21 +105,13 @@ set(LIB
   ${CYCLES_GL_LIBRARIES}
 )
 
-if(WITH_CYCLES_DEVICE_OPTIX OR WITH_CYCLES_DEVICE_CUDA)
-  if(WITH_CUDA_DYNLOAD)
-    list(APPEND LIB
-      extern_cuew
-    )
-  else()
-    list(APPEND LIB
-      ${CUDA_CUDA_LIBRARY}
-    )
-  endif()
-endif()
-
-if(WITH_CYCLES_DEVICE_HIP AND WITH_HIP_DYNLOAD)
+if(WITH_CUDA_DYNLOAD)
   list(APPEND LIB
-    extern_hipew
+    extern_cuew
+  )
+else()
+  list(APPEND LIB
+    ${CUDA_CUDA_LIBRARY}
   )
 endif()
 
@@ -152,9 +120,6 @@ add_definitions(${GL_DEFINITIONS})
 if(WITH_CYCLES_DEVICE_CUDA)
   add_definitions(-DWITH_CUDA)
 endif()
-if(WITH_CYCLES_DEVICE_HIP)
-  add_definitions(-DWITH_HIP)
-endif()
 if(WITH_CYCLES_DEVICE_OPTIX)
   add_definitions(-DWITH_OPTIX)
 endif()
@@ -175,7 +140,6 @@ cycles_add_library(cycles_device "${LIB}"
   ${SRC}
   ${SRC_CPU}
   ${SRC_CUDA}
-  ${SRC_HIP}
   ${SRC_DUMMY}
   ${SRC_MULTI}
   ${SRC_OPTIX}

intern/cycles/device/cuda/queue.cpp

@@ -116,18 +116,18 @@ bool CUDADeviceQueue::enqueue(DeviceKernel kernel, const int work_size, void *ar
   }
 
   /* Launch kernel. */
-  assert_success(cuLaunchKernel(cuda_kernel.function,
-                                num_blocks,
-                                1,
-                                1,
-                                num_threads_per_block,
-                                1,
-                                1,
-                                shared_mem_bytes,
-                                cuda_stream_,
-                                args,
-                                0),
-                 "enqueue");
+  cuda_device_assert(cuda_device_,
+                     cuLaunchKernel(cuda_kernel.function,
+                                    num_blocks,
+                                    1,
+                                    1,
+                                    num_threads_per_block,
+                                    1,
+                                    1,
+                                    shared_mem_bytes,
+                                    cuda_stream_,
+                                    args,
+                                    0));
 
   return !(cuda_device_->have_error());
 }
@@ -139,8 +139,7 @@ bool CUDADeviceQueue::synchronize()
   }
 
   const CUDAContextScope scope(cuda_device_);
-  assert_success(cuStreamSynchronize(cuda_stream_), "synchronize");
-
+  cuda_device_assert(cuda_device_, cuStreamSynchronize(cuda_stream_));
   debug_synchronize();
 
   return !(cuda_device_->have_error());
@@ -163,9 +162,9 @@ void CUDADeviceQueue::zero_to_device(device_memory &mem)
   assert(mem.device_pointer != 0);
 
   const CUDAContextScope scope(cuda_device_);
-  assert_success(
-      cuMemsetD8Async((CUdeviceptr)mem.device_pointer, 0, mem.memory_size(), cuda_stream_),
-      "zero_to_device");
+  cuda_device_assert(
+      cuda_device_,
+      cuMemsetD8Async((CUdeviceptr)mem.device_pointer, 0, mem.memory_size(), cuda_stream_));
 }
 
 void CUDADeviceQueue::copy_to_device(device_memory &mem)
@@ -186,10 +185,10 @@ void CUDADeviceQueue::copy_to_device(device_memory &mem)
 
   /* Copy memory to device. */
   const CUDAContextScope scope(cuda_device_);
-  assert_success(
+  cuda_device_assert(
+      cuda_device_,
       cuMemcpyHtoDAsync(
-          (CUdeviceptr)mem.device_pointer, mem.host_pointer, mem.memory_size(), cuda_stream_),
-      "copy_to_device");
+          (CUdeviceptr)mem.device_pointer, mem.host_pointer, mem.memory_size(), cuda_stream_));
 }
 
 void CUDADeviceQueue::copy_from_device(device_memory &mem)
@@ -205,19 +204,10 @@ void CUDADeviceQueue::copy_from_device(device_memory &mem)
 
   /* Copy memory from device. */
   const CUDAContextScope scope(cuda_device_);
-  assert_success(
+  cuda_device_assert(
+      cuda_device_,
       cuMemcpyDtoHAsync(
-          mem.host_pointer, (CUdeviceptr)mem.device_pointer, mem.memory_size(), cuda_stream_),
-      "copy_from_device");
-}
-
-void CUDADeviceQueue::assert_success(CUresult result, const char *operation)
-{
-  if (result != CUDA_SUCCESS) {
-    const char *name = cuewErrorString(result);
-    cuda_device_->set_error(string_printf(
-        "%s in CUDA queue %s (%s)", name, operation, debug_active_kernels().c_str()));
-  }
+          mem.host_pointer, (CUdeviceptr)mem.device_pointer, mem.memory_size(), cuda_stream_));
 }
 
 unique_ptr<DeviceGraphicsInterop> CUDADeviceQueue::graphics_interop_create()

intern/cycles/device/cuda/queue.h

@@ -60,8 +60,6 @@ class CUDADeviceQueue : public DeviceQueue {
  protected:
   CUDADevice *cuda_device_;
   CUstream cuda_stream_;
-
-  void assert_success(CUresult result, const char *operation);
 };
 
 CCL_NAMESPACE_END

intern/cycles/device/device.cpp

@@ -25,7 +25,6 @@
 #include "device/cpu/device.h"
 #include "device/cuda/device.h"
 #include "device/dummy/device.h"
-#include "device/hip/device.h"
 #include "device/multi/device.h"
 #include "device/optix/device.h"
 
@@ -47,7 +46,6 @@ thread_mutex Device::device_mutex;
 vector<DeviceInfo> Device::cuda_devices;
 vector<DeviceInfo> Device::optix_devices;
 vector<DeviceInfo> Device::cpu_devices;
-vector<DeviceInfo> Device::hip_devices;
 uint Device::devices_initialized_mask = 0;
 
 /* Device */
@@ -98,14 +96,6 @@ Device *Device::create(const DeviceInfo &info, Stats &stats, Profiler &profiler)
         device = device_optix_create(info, stats, profiler);
       break;
 #endif
-
-#ifdef WITH_HIP
-    case DEVICE_HIP:
-      if (device_hip_init())
-        device = device_hip_create(info, stats, profiler);
-      break;
-#endif
-
     default:
       break;
   }
@@ -127,8 +117,6 @@ DeviceType Device::type_from_string(const char *name)
     return DEVICE_OPTIX;
   else if (strcmp(name, "MULTI") == 0)
     return DEVICE_MULTI;
-  else if (strcmp(name, "HIP") == 0)
-    return DEVICE_HIP;
 
   return DEVICE_NONE;
 }
@@ -143,8 +131,6 @@ string Device::string_from_type(DeviceType type)
     return "OPTIX";
   else if (type == DEVICE_MULTI)
     return "MULTI";
-  else if (type == DEVICE_HIP)
-    return "HIP";
 
   return "";
 }
@@ -159,10 +145,6 @@ vector<DeviceType> Device::available_types()
 #ifdef WITH_OPTIX
   types.push_back(DEVICE_OPTIX);
 #endif
-#ifdef WITH_HIP
-  types.push_back(DEVICE_HIP);
-#endif
-
   return types;
 }
 
@@ -204,20 +186,6 @@ vector<DeviceInfo> Device::available_devices(uint mask)
   }
 #endif
 
-#ifdef WITH_HIP
-  if (mask & DEVICE_MASK_HIP) {
-    if (!(devices_initialized_mask & DEVICE_MASK_HIP)) {
-      if (device_hip_init()) {
-        device_hip_info(hip_devices);
-      }
-      devices_initialized_mask |= DEVICE_MASK_HIP;
-    }
-    foreach (DeviceInfo &info, hip_devices) {
-      devices.push_back(info);
-    }
-  }
-#endif
-
   if (mask & DEVICE_MASK_CPU) {
     if (!(devices_initialized_mask & DEVICE_MASK_CPU)) {
       device_cpu_info(cpu_devices);
@@ -258,15 +226,6 @@ string Device::device_capabilities(uint mask)
   }
 #endif
 
-#ifdef WITH_HIP
-  if (mask & DEVICE_MASK_HIP) {
-    if (device_hip_init()) {
-      capabilities += "\nHIP device capabilities:\n";
-      capabilities += device_hip_capabilities();
-    }
-  }
-#endif
-
   return capabilities;
 }
 
@@ -355,7 +314,6 @@ void Device::free_memory()
   devices_initialized_mask = 0;
   cuda_devices.free_memory();
   optix_devices.free_memory();
-  hip_devices.free_memory();
   cpu_devices.free_memory();
 }
 

intern/cycles/device/device.h

@@ -51,7 +51,6 @@ enum DeviceType {
   DEVICE_CUDA,
   DEVICE_MULTI,
   DEVICE_OPTIX,
-  DEVICE_HIP,
   DEVICE_DUMMY,
 };
 
@@ -59,7 +58,6 @@ enum DeviceTypeMask {
   DEVICE_MASK_CPU = (1 << DEVICE_CPU),
   DEVICE_MASK_CUDA = (1 << DEVICE_CUDA),
   DEVICE_MASK_OPTIX = (1 << DEVICE_OPTIX),
-  DEVICE_MASK_HIP = (1 << DEVICE_HIP),
   DEVICE_MASK_ALL = ~0
 };
 
@@ -286,7 +284,6 @@ class Device {
   static vector<DeviceInfo> cuda_devices;
   static vector<DeviceInfo> optix_devices;
   static vector<DeviceInfo> cpu_devices;
-  static vector<DeviceInfo> hip_devices;
   static uint devices_initialized_mask;
 };
 

intern/cycles/device/device_memory.h

@@ -277,7 +277,6 @@ class device_memory {
  protected:
   friend class CUDADevice;
   friend class OptiXDevice;
-  friend class HIPDevice;
 
   /* Only create through subclasses. */
   device_memory(Device *device, const char *name, MemoryType type);

intern/cycles/device/device_queue.cpp

@@ -57,9 +57,8 @@ void DeviceQueue::debug_init_execution()
 {
   if (VLOG_IS_ON(3)) {
     last_sync_time_ = time_dt();
+    last_kernels_enqueued_ = 0;
   }
-
-  last_kernels_enqueued_ = 0;
 }
 
 void DeviceQueue::debug_enqueue(DeviceKernel kernel, const int work_size)
@@ -67,9 +66,8 @@ void DeviceQueue::debug_enqueue(DeviceKernel kernel, const int work_size)
   if (VLOG_IS_ON(3)) {
     VLOG(4) << "GPU queue launch " << device_kernel_as_string(kernel) << ", work_size "
             << work_size;
+    last_kernels_enqueued_ |= (uint64_t(1) << (uint64_t)kernel);
   }
-
-  last_kernels_enqueued_ |= (uint64_t(1) << (uint64_t)kernel);
 }
 
 void DeviceQueue::debug_synchronize()
@@ -82,14 +80,8 @@ void DeviceQueue::debug_synchronize()
     stats_kernel_time_[last_kernels_enqueued_] += elapsed_time;
 
     last_sync_time_ = new_time;
+    last_kernels_enqueued_ = 0;
   }
-
-  last_kernels_enqueued_ = 0;
-}
-
-string DeviceQueue::debug_active_kernels()
-{
-  return device_kernel_mask_as_string(last_kernels_enqueued_);
 }
 
 CCL_NAMESPACE_END

intern/cycles/device/device_queue.h

@@ -21,7 +21,6 @@
 #include "device/device_graphics_interop.h"
 #include "util/util_logging.h"
 #include "util/util_map.h"
-#include "util/util_string.h"
 #include "util/util_unique_ptr.h"
 
 CCL_NAMESPACE_BEGIN
@@ -102,7 +101,6 @@ class DeviceQueue {
   void debug_init_execution();
   void debug_enqueue(DeviceKernel kernel, const int work_size);
   void debug_synchronize();
-  string debug_active_kernels();
 
   /* Combination of kernels enqueued together sync last synchronize. */
   DeviceKernelMask last_kernels_enqueued_;

intern/cycles/device/hip/device.cpp

@@ -1,276 +0,0 @@
-/*
- * Copyright 2011-2021 Blender Foundation
- *
- * Licensed under the Apache License, Version 2.0 (the "License");
- * you may not use this file except in compliance with the License.
- * You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing, software
- * distributed under the License is distributed on an "AS IS" BASIS,
- * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
- * See the License for the specific language governing permissions and
- * limitations under the License.
- */
-
-#include "device/hip/device.h"
-
-#include "util/util_logging.h"
-
-#ifdef WITH_HIP
-#  include "device/device.h"
-#  include "device/hip/device_impl.h"
-
-#  include "util/util_string.h"
-#  include "util/util_windows.h"
-#endif /* WITH_HIP */
-
-CCL_NAMESPACE_BEGIN
-
-bool device_hip_init()
-{
-#if !defined(WITH_HIP)
-  return false;
-#elif defined(WITH_HIP_DYNLOAD)
-  static bool initialized = false;
-  static bool result = false;
-
-  if (initialized)
-    return result;
-
-  initialized = true;
-  int hipew_result = hipewInit(HIPEW_INIT_HIP);
-  if (hipew_result == HIPEW_SUCCESS) {
-    VLOG(1) << "HIPEW initialization succeeded";
-    if (HIPDevice::have_precompiled_kernels()) {
-      VLOG(1) << "Found precompiled kernels";
-      result = true;
-    }
-    else if (hipewCompilerPath() != NULL) {
-      VLOG(1) << "Found HIPCC " << hipewCompilerPath();
-      result = true;
-    }
-    else {
-      VLOG(1) << "Neither precompiled kernels nor HIPCC was found,"
-              << " unable to use HIP";
-    }
-  }
-  else {
-    VLOG(1) << "HIPEW initialization failed: "
-            << ((hipew_result == HIPEW_ERROR_ATEXIT_FAILED) ? "Error setting up atexit() handler" :
-                                                              "Error opening the library");
-  }
-
-  return result;
-#else  /* WITH_HIP_DYNLOAD */
-  return true;
-#endif /* WITH_HIP_DYNLOAD */
-}
-
-Device *device_hip_create(const DeviceInfo &info, Stats &stats, Profiler &profiler)
-{
-#ifdef WITH_HIP
-  return new HIPDevice(info, stats, profiler);
-#else
-  (void)info;
-  (void)stats;
-  (void)profiler;
-
-  LOG(FATAL) << "Request to create HIP device without compiled-in support. Should never happen.";
-
-  return nullptr;
-#endif
-}
-
-#ifdef WITH_HIP
-static hipError_t device_hip_safe_init()
-{
-#  ifdef _WIN32
-  __try {
-    return hipInit(0);
-  }
-  __except (EXCEPTION_EXECUTE_HANDLER) {
-    /* Ignore crashes inside the HIP driver and hope we can
-     * survive even with corrupted HIP installs. */
-    fprintf(stderr, "Cycles HIP: driver crashed, continuing without HIP.\n");
-  }
-
-  return hipErrorNoDevice;
-#  else
-  return hipInit(0);
-#  endif
-}
-#endif /* WITH_HIP */
-
-void device_hip_info(vector<DeviceInfo> &devices)
-{
-#ifdef WITH_HIP
-  hipError_t result = device_hip_safe_init();
-  if (result != hipSuccess) {
-    if (result != hipErrorNoDevice)
-      fprintf(stderr, "HIP hipInit: %s\n", hipewErrorString(result));
-    return;
-  }
-
-  int count = 0;
-  result = hipGetDeviceCount(&count);
-  if (result != hipSuccess) {
-    fprintf(stderr, "HIP hipGetDeviceCount: %s\n", hipewErrorString(result));
-    return;
-  }
-
-  vector<DeviceInfo> display_devices;
-
-  for (int num = 0; num < count; num++) {
-    char name[256];
-
-    result = hipDeviceGetName(name, 256, num);
-    if (result != hipSuccess) {
-      fprintf(stderr, "HIP :hipDeviceGetName: %s\n", hipewErrorString(result));
-      continue;
-    }
-
-    int major;
-    hipDeviceGetAttribute(&major, hipDeviceAttributeComputeCapabilityMajor, num);
-    // TODO : (Arya) What is the last major version we are supporting?
-
-    DeviceInfo info;
-
-    info.type = DEVICE_HIP;
-    info.description = string(name);
-    info.num = num;
-
-    info.has_half_images = (major >= 3);
-    info.has_nanovdb = true;
-    info.denoisers = 0;
-
-    info.has_gpu_queue = true;
-    /* Check if the device has P2P access to any other device in the system. */
-    for (int peer_num = 0; peer_num < count && !info.has_peer_memory; peer_num++) {
-      if (num != peer_num) {
-        int can_access = 0;
-        hipDeviceCanAccessPeer(&can_access, num, peer_num);
-        info.has_peer_memory = (can_access != 0);
-      }
-    }
-
-    int pci_location[3] = {0, 0, 0};
-    hipDeviceGetAttribute(&pci_location[0], hipDeviceAttributePciDomainID, num);
-    hipDeviceGetAttribute(&pci_location[1], hipDeviceAttributePciBusId, num);
-    hipDeviceGetAttribute(&pci_location[2], hipDeviceAttributePciDeviceId, num);
-    info.id = string_printf("HIP_%s_%04x:%02x:%02x",
-                            name,
-                            (unsigned int)pci_location[0],
-                            (unsigned int)pci_location[1],
-                            (unsigned int)pci_location[2]);
-
-    /* If device has a kernel timeout and no compute preemption, we assume
-     * it is connected to a display and will freeze the display while doing
-     * computations. */
-    int timeout_attr = 0, preempt_attr = 0;
-    hipDeviceGetAttribute(&timeout_attr, hipDeviceAttributeKernelExecTimeout, num);
-
-    if (timeout_attr && !preempt_attr) {
-      VLOG(1) << "Device is recognized as display.";
-      info.description += " (Display)";
-      info.display_device = true;
-      display_devices.push_back(info);
-    }
-    else {
-      VLOG(1) << "Device has compute preemption or is not used for display.";
-      devices.push_back(info);
-    }
-    VLOG(1) << "Added device \"" << name << "\" with id \"" << info.id << "\".";
-  }
-
-  if (!display_devices.empty())
-    devices.insert(devices.end(), display_devices.begin(), display_devices.end());
-#else  /* WITH_HIP */
-  (void)devices;
-#endif /* WITH_HIP */
-}
-
-string device_hip_capabilities()
-{
-#ifdef WITH_HIP
-  hipError_t result = device_hip_safe_init();
-  if (result != hipSuccess) {
-    if (result != hipErrorNoDevice) {
-      return string("Error initializing HIP: ") + hipewErrorString(result);
-    }
-    return "No HIP device found\n";
-  }
-
-  int count;
-  result = hipGetDeviceCount(&count);
-  if (result != hipSuccess) {
-    return string("Error getting devices: ") + hipewErrorString(result);
-  }
-
-  string capabilities = "";
-  for (int num = 0; num < count; num++) {
-    char name[256];
-    if (hipDeviceGetName(name, 256, num) != hipSuccess) {
-      continue;
-    }
-    capabilities += string("\t") + name + "\n";
-    int value;
-#  define GET_ATTR(attr) \
-    { \
-      if (hipDeviceGetAttribute(&value, hipDeviceAttribute##attr, num) == hipSuccess) { \
-        capabilities += string_printf("\t\thipDeviceAttribute" #attr "\t\t\t%d\n", value); \
-      } \
-    } \
-    (void)0
-    /* TODO(sergey): Strip all attributes which are not useful for us
-     * or does not depend on the driver.
-     */
-    GET_ATTR(MaxThreadsPerBlock);
-    GET_ATTR(MaxBlockDimX);
-    GET_ATTR(MaxBlockDimY);
-    GET_ATTR(MaxBlockDimZ);
-    GET_ATTR(MaxGridDimX);
-    GET_ATTR(MaxGridDimY);
-    GET_ATTR(MaxGridDimZ);
-    GET_ATTR(MaxSharedMemoryPerBlock);
-    GET_ATTR(TotalConstantMemory);
-    GET_ATTR(WarpSize);
-    GET_ATTR(MaxPitch);
-    GET_ATTR(MaxRegistersPerBlock);
-    GET_ATTR(ClockRate);
-    GET_ATTR(TextureAlignment);
-    GET_ATTR(MultiprocessorCount);
-    GET_ATTR(KernelExecTimeout);
-    GET_ATTR(Integrated);
-    GET_ATTR(CanMapHostMemory);
-    GET_ATTR(ComputeMode);
-    GET_ATTR(MaxTexture1DWidth);
-    GET_ATTR(MaxTexture2DWidth);
-    GET_ATTR(MaxTexture2DHeight);
-    GET_ATTR(MaxTexture3DWidth);
-    GET_ATTR(MaxTexture3DHeight);
-    GET_ATTR(MaxTexture3DDepth);
-    GET_ATTR(ConcurrentKernels);
-    GET_ATTR(EccEnabled);
-    GET_ATTR(MemoryClockRate);
-    GET_ATTR(MemoryBusWidth);
-    GET_ATTR(L2CacheSize);
-    GET_ATTR(MaxThreadsPerMultiProcessor);
-    GET_ATTR(ComputeCapabilityMajor);
-    GET_ATTR(ComputeCapabilityMinor);
-    GET_ATTR(MaxSharedMemoryPerMultiprocessor);
-    GET_ATTR(ManagedMemory);
-    GET_ATTR(IsMultiGpuBoard);
-#  undef GET_ATTR
-    capabilities += "\n";
-  }
-
-  return capabilities;
-
-#else  /* WITH_HIP */
-  return "";
-#endif /* WITH_HIP */
-}
-
-CCL_NAMESPACE_END

intern/cycles/device/hip/device.h

@@ -1,37 +0,0 @@
-/*
- * Copyright 2011-2021 Blender Foundation
- *
- * Licensed under the Apache License, Version 2.0 (the "License");
- * you may not use this file except in compliance with the License.
- * You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing, software
- * distributed under the License is distributed on an "AS IS" BASIS,
- * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
- * See the License for the specific language governing permissions and
- * limitations under the License.
- */
-
-#pragma once
-
-#include "util/util_string.h"
-#include "util/util_vector.h"
-
-CCL_NAMESPACE_BEGIN
-
-class Device;
-class DeviceInfo;
-class Profiler;
-class Stats;
-
-bool device_hip_init();
-
-Device *device_hip_create(const DeviceInfo &info, Stats &stats, Profiler &profiler);
-
-void device_hip_info(vector<DeviceInfo> &devices);
-
-string device_hip_capabilities();
-
-CCL_NAMESPACE_END

intern/cycles/device/hip/device_impl.h

@@ -1,153 +0,0 @@
-/*
- * Copyright 2011-2021 Blender Foundation
- *
- * Licensed under the Apache License, Version 2.0 (the "License");
- * you may not use this file except in compliance with the License.
- * You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing, software
- * distributed under the License is distributed on an "AS IS" BASIS,
- * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
- * See the License for the specific language governing permissions and
- * limitations under the License.
- */
-
-#ifdef WITH_HIP
-
-#  include "device/device.h"
-#  include "device/hip/kernel.h"
-#  include "device/hip/queue.h"
-#  include "device/hip/util.h"
-
-#  include "util/util_map.h"
-
-#  ifdef WITH_HIP_DYNLOAD
-#    include "hipew.h"
-#  else
-#    include "util/util_opengl.h"
-#  endif
-
-CCL_NAMESPACE_BEGIN
-
-class DeviceQueue;
-
-class HIPDevice : public Device {
-
-  friend class HIPContextScope;
-
- public:
-  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();
-
-  virtual bool show_samples() const override;
-
-  virtual BVHLayoutMask get_bvh_layout_mask() const override;
-
-  void set_error(const string &error) override;
-
-  HIPDevice(const DeviceInfo &info, Stats &stats, Profiler &profiler);
-
-  virtual ~HIPDevice();
-
-  bool support_device(const uint /*kernel_features*/);
-
-  bool check_peer_access(Device *peer_device) override;
-
-  bool use_adaptive_compilation();
-
-  virtual string compile_kernel_get_common_cflags(const uint kernel_features);
-
-  string compile_kernel(const uint kernel_features,
-                        const char *name,
-                        const char *base = "hip",
-                        bool force_ptx = false);
-
-  virtual bool load_kernels(const uint kernel_features) override;
-  void reserve_local_memory(const uint kernel_features);
-
-  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;
-
-  void mem_copy_to(device_memory &mem) override;
-
-  void mem_copy_from(device_memory &mem, size_t y, size_t w, size_t h, size_t elem) override;
-
-  void mem_zero(device_memory &mem) override;
-
-  void mem_free(device_memory &mem) override;
-
-  device_ptr mem_alloc_sub_ptr(device_memory &mem, size_t offset, size_t /*size*/) override;
-
-  virtual void const_copy_to(const char *name, void *host, size_t size) override;
-
-  void global_alloc(device_memory &mem);
-
-  void global_free(device_memory &mem);
-
-  void tex_alloc(device_texture &mem);
-
-  void tex_free(device_texture &mem);
-
-  /* Graphics resources interoperability. */
-  virtual bool should_use_graphics_interop() override;
-
-  virtual unique_ptr<DeviceQueue> gpu_queue_create() override;
-
-  int get_num_multiprocessors();
-  int get_max_num_threads_per_multiprocessor();
-
- protected:
-  bool get_device_attribute(hipDeviceAttribute_t attribute, int *value);
-  int get_device_default_attribute(hipDeviceAttribute_t attribute, int default_value);
-};
-
-CCL_NAMESPACE_END
-
-#endif

intern/cycles/device/hip/graphics_interop.cpp

@@ -1,93 +0,0 @@
-/*
- * Copyright 2011-2021 Blender Foundation
- *
- * Licensed under the Apache License, Version 2.0 (the "License");
- * you may not use this file except in compliance with the License.
- * You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing, software
- * distributed under the License is distributed on an "AS IS" BASIS,
- * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
- * See the License for the specific language governing permissions and
- * limitations under the License.
- */
-
-#ifdef WITH_HIP
-
-#  include "device/hip/graphics_interop.h"
-
-#  include "device/hip/device_impl.h"
-#  include "device/hip/util.h"
-
-CCL_NAMESPACE_BEGIN
-
-HIPDeviceGraphicsInterop::HIPDeviceGraphicsInterop(HIPDeviceQueue *queue)
-    : queue_(queue), device_(static_cast<HIPDevice *>(queue->device))
-{
-}
-
-HIPDeviceGraphicsInterop::~HIPDeviceGraphicsInterop()
-{
-  HIPContextScope scope(device_);
-
-  if (hip_graphics_resource_) {
-    hip_device_assert(device_, hipGraphicsUnregisterResource(hip_graphics_resource_));
-  }
-}
-
-void HIPDeviceGraphicsInterop::set_destination(const DeviceGraphicsInteropDestination &destination)
-{
-  const int64_t new_buffer_area = int64_t(destination.buffer_width) * destination.buffer_height;
-
-  if (opengl_pbo_id_ == destination.opengl_pbo_id && buffer_area_ == new_buffer_area) {
-    return;
-  }
-
-  HIPContextScope scope(device_);
-
-  if (hip_graphics_resource_) {
-    hip_device_assert(device_, hipGraphicsUnregisterResource(hip_graphics_resource_));
-  }
-
-  const hipError_t result = hipGraphicsGLRegisterBuffer(
-      &hip_graphics_resource_, destination.opengl_pbo_id, hipGraphicsRegisterFlagsNone);
-  if (result != hipSuccess) {
-    LOG(ERROR) << "Error registering OpenGL buffer: " << hipewErrorString(result);
-  }
-
-  opengl_pbo_id_ = destination.opengl_pbo_id;
-  buffer_area_ = new_buffer_area;
-}
-
-device_ptr HIPDeviceGraphicsInterop::map()
-{
-  if (!hip_graphics_resource_) {
-    return 0;
-  }
-
-  HIPContextScope scope(device_);
-
-  hipDeviceptr_t hip_buffer;
-  size_t bytes;
-
-  hip_device_assert(device_,
-                    hipGraphicsMapResources(1, &hip_graphics_resource_, queue_->stream()));
-  hip_device_assert(
-      device_, hipGraphicsResourceGetMappedPointer(&hip_buffer, &bytes, hip_graphics_resource_));
-
-  return static_cast<device_ptr>(hip_buffer);
-}
-
-void HIPDeviceGraphicsInterop::unmap()
-{
-  HIPContextScope scope(device_);
-
-  hip_device_assert(device_,
-                    hipGraphicsUnmapResources(1, &hip_graphics_resource_, queue_->stream()));
-}
-
-CCL_NAMESPACE_END
-
-#endif

intern/cycles/device/hip/graphics_interop.h

@@ -1,61 +0,0 @@
-/*
- * Copyright 2011-2021 Blender Foundation
- *
- * Licensed under the Apache License, Version 2.0 (the "License");
- * you may not use this file except in compliance with the License.
- * You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing, software
- * distributed under the License is distributed on an "AS IS" BASIS,
- * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
- * See the License for the specific language governing permissions and
- * limitations under the License.
- */
-
-#ifdef WITH_HIP
-
-#  include "device/device_graphics_interop.h"
-
-#  ifdef WITH_HIP_DYNLOAD
-#    include "hipew.h"
-#  endif
-
-CCL_NAMESPACE_BEGIN
-
-class HIPDevice;
-class HIPDeviceQueue;
-
-class HIPDeviceGraphicsInterop : public DeviceGraphicsInterop {
- public:
-  explicit HIPDeviceGraphicsInterop(HIPDeviceQueue *queue);
-
-  HIPDeviceGraphicsInterop(const HIPDeviceGraphicsInterop &other) = delete;
-  HIPDeviceGraphicsInterop(HIPDeviceGraphicsInterop &&other) noexcept = delete;
-
-  ~HIPDeviceGraphicsInterop();
-
-  HIPDeviceGraphicsInterop &operator=(const HIPDeviceGraphicsInterop &other) = delete;
-  HIPDeviceGraphicsInterop &operator=(HIPDeviceGraphicsInterop &&other) = delete;
-
-  virtual void set_destination(const DeviceGraphicsInteropDestination &destination) override;
-
-  virtual device_ptr map() override;
-  virtual void unmap() override;
-
- protected:
-  HIPDeviceQueue *queue_ = nullptr;
-  HIPDevice *device_ = nullptr;
-
-  /* OpenGL PBO which is currently registered as the destination for the CUDA buffer. */
-  uint opengl_pbo_id_ = 0;
-  /* Buffer area in pixels of the corresponding PBO. */
-  int64_t buffer_area_ = 0;
-
-  hipGraphicsResource hip_graphics_resource_ = nullptr;
-};
-
-CCL_NAMESPACE_END
-
-#endif

intern/cycles/device/hip/kernel.cpp

@@ -1,69 +0,0 @@
-/*
- * Copyright 2011-2021 Blender Foundation
- *
- * Licensed under the Apache License, Version 2.0 (the "License");
- * you may not use this file except in compliance with the License.
- * You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing, software
- * distributed under the License is distributed on an "AS IS" BASIS,
- * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
- * See the License for the specific language governing permissions and
- * limitations under the License.
- */
-
-#ifdef WITH_HIP
-
-#  include "device/hip/kernel.h"
-#  include "device/hip/device_impl.h"
-
-CCL_NAMESPACE_BEGIN
-
-void HIPDeviceKernels::load(HIPDevice *device)
-{
-  hipModule_t hipModule = device->hipModule;
-
-  for (int i = 0; i < (int)DEVICE_KERNEL_NUM; i++) {
-    HIPDeviceKernel &kernel = kernels_[i];
-
-    /* No megakernel used for GPU. */
-    if (i == DEVICE_KERNEL_INTEGRATOR_MEGAKERNEL) {
-      continue;
-    }
-
-    const std::string function_name = std::string("kernel_gpu_") +
-                                      device_kernel_as_string((DeviceKernel)i);
-    hip_device_assert(device,
-                      hipModuleGetFunction(&kernel.function, hipModule, function_name.c_str()));
-
-    if (kernel.function) {
-      hip_device_assert(device, hipFuncSetCacheConfig(kernel.function, hipFuncCachePreferL1));
-
-      hip_device_assert(
-          device,
-          hipModuleOccupancyMaxPotentialBlockSize(
-              &kernel.min_blocks, &kernel.num_threads_per_block, kernel.function, 0, 0));
-    }
-    else {
-      LOG(ERROR) << "Unable to load kernel " << function_name;
-    }
-  }
-
-  loaded = true;
-}
-
-const HIPDeviceKernel &HIPDeviceKernels::get(DeviceKernel kernel) const
-{
-  return kernels_[(int)kernel];
-}
-
-bool HIPDeviceKernels::available(DeviceKernel kernel) const
-{
-  return kernels_[(int)kernel].function != nullptr;
-}
-
-CCL_NAMESPACE_END
-
-#endif /* WITH_HIP*/

intern/cycles/device/hip/kernel.h

@@ -1,54 +0,0 @@
-/*
- * Copyright 2011-2021 Blender Foundation
- *
- * Licensed under the Apache License, Version 2.0 (the "License");
- * you may not use this file except in compliance with the License.
- * You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing, software
- * distributed under the License is distributed on an "AS IS" BASIS,
- * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
- * See the License for the specific language governing permissions and
- * limitations under the License.
- */
-
-#pragma once
-
-#ifdef WITH_HIP
-
-#  include "device/device_kernel.h"
-
-#  ifdef WITH_HIP_DYNLOAD
-#    include "hipew.h"
-#  endif
-
-CCL_NAMESPACE_BEGIN
-
-class HIPDevice;
-
-/* HIP kernel and associate occupancy information. */
-class HIPDeviceKernel {
- public:
-  hipFunction_t function = nullptr;
-
-  int num_threads_per_block = 0;
-  int min_blocks = 0;
-};
-
-/* Cache of HIP kernels for each DeviceKernel. */
-class HIPDeviceKernels {
- public:
-  void load(HIPDevice *device);
-  const HIPDeviceKernel &get(DeviceKernel kernel) const;
-  bool available(DeviceKernel kernel) const;
-
- protected:
-  HIPDeviceKernel kernels_[DEVICE_KERNEL_NUM];
-  bool loaded = false;
-};
-
-CCL_NAMESPACE_END
-
-#endif /* WITH_HIP */

intern/cycles/device/hip/queue.cpp

@@ -1,209 +0,0 @@
-/*
- * Copyright 2011-2021 Blender Foundation
- *
- * Licensed under the Apache License, Version 2.0 (the "License");
- * you may not use this file except in compliance with the License.
- * You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing, software
- * distributed under the License is distributed on an "AS IS" BASIS,
- * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
- * See the License for the specific language governing permissions and
- * limitations under the License.
- */
-
-#ifdef WITH_HIP
-
-#  include "device/hip/queue.h"
-
-#  include "device/hip/device_impl.h"
-#  include "device/hip/graphics_interop.h"
-#  include "device/hip/kernel.h"
-
-CCL_NAMESPACE_BEGIN
-
-/* HIPDeviceQueue */
-
-HIPDeviceQueue::HIPDeviceQueue(HIPDevice *device)
-    : DeviceQueue(device), hip_device_(device), hip_stream_(nullptr)
-{
-  const HIPContextScope scope(hip_device_);
-  hip_device_assert(hip_device_, hipStreamCreateWithFlags(&hip_stream_, hipStreamNonBlocking));
-}
-
-HIPDeviceQueue::~HIPDeviceQueue()
-{
-  const HIPContextScope scope(hip_device_);
-  hipStreamDestroy(hip_stream_);
-}
-
-int HIPDeviceQueue::num_concurrent_states(const size_t /*state_size*/) const
-{
-  /* TODO: compute automatically. */
-  /* TODO: must have at least num_threads_per_block. */
-  return 14416128;
-}
-
-int HIPDeviceQueue::num_concurrent_busy_states() const
-{
-  const int max_num_threads = hip_device_->get_num_multiprocessors() *
-                              hip_device_->get_max_num_threads_per_multiprocessor();
-
-  if (max_num_threads == 0) {
-    return 65536;
-  }
-
-  return 4 * max_num_threads;
-}
-
-void HIPDeviceQueue::init_execution()
-{
-  /* Synchronize all textures and memory copies before executing task. */
-  HIPContextScope scope(hip_device_);
-  hip_device_->load_texture_info();
-  hip_device_assert(hip_device_, hipDeviceSynchronize());
-
-  debug_init_execution();
-}
-
-bool HIPDeviceQueue::kernel_available(DeviceKernel kernel) const
-{
-  return hip_device_->kernels.available(kernel);
-}
-
-bool HIPDeviceQueue::enqueue(DeviceKernel kernel, const int work_size, void *args[])
-{
-  if (hip_device_->have_error()) {
-    return false;
-  }
-
-  debug_enqueue(kernel, work_size);
-
-  const HIPContextScope scope(hip_device_);
-  const HIPDeviceKernel &hip_kernel = hip_device_->kernels.get(kernel);
-
-  /* Compute kernel launch parameters. */
-  const int num_threads_per_block = hip_kernel.num_threads_per_block;
-  const int num_blocks = divide_up(work_size, num_threads_per_block);
-
-  int shared_mem_bytes = 0;
-
-  switch (kernel) {
-    case DEVICE_KERNEL_INTEGRATOR_QUEUED_PATHS_ARRAY:
-    case DEVICE_KERNEL_INTEGRATOR_QUEUED_SHADOW_PATHS_ARRAY:
-    case DEVICE_KERNEL_INTEGRATOR_ACTIVE_PATHS_ARRAY:
-    case DEVICE_KERNEL_INTEGRATOR_TERMINATED_PATHS_ARRAY:
-    case DEVICE_KERNEL_INTEGRATOR_SORTED_PATHS_ARRAY:
-    case DEVICE_KERNEL_INTEGRATOR_COMPACT_PATHS_ARRAY:
-      /* See parall_active_index.h for why this amount of shared memory is needed. */
-      shared_mem_bytes = (num_threads_per_block + 1) * sizeof(int);
-      break;
-    default:
-      break;
-  }
-
-  /* Launch kernel. */
-  hip_device_assert(hip_device_,
-                    hipModuleLaunchKernel(hip_kernel.function,
-                                          num_blocks,
-                                          1,
-                                          1,
-                                          num_threads_per_block,
-                                          1,
-                                          1,
-                                          shared_mem_bytes,
-                                          hip_stream_,
-                                          args,
-                                          0));
-  return !(hip_device_->have_error());
-}
-
-bool HIPDeviceQueue::synchronize()
-{
-  if (hip_device_->have_error()) {
-    return false;
-  }
-
-  const HIPContextScope scope(hip_device_);
-  hip_device_assert(hip_device_, hipStreamSynchronize(hip_stream_));
-  debug_synchronize();
-
-  return !(hip_device_->have_error());
-}
-
-void HIPDeviceQueue::zero_to_device(device_memory &mem)
-{
-  assert(mem.type != MEM_GLOBAL && mem.type != MEM_TEXTURE);
-
-  if (mem.memory_size() == 0) {
-    return;
-  }
-
-  /* Allocate on demand. */
-  if (mem.device_pointer == 0) {
-    hip_device_->mem_alloc(mem);
-  }
-
-  /* Zero memory on device. */
-  assert(mem.device_pointer != 0);
-
-  const HIPContextScope scope(hip_device_);
-  hip_device_assert(
-      hip_device_,
-      hipMemsetD8Async((hipDeviceptr_t)mem.device_pointer, 0, mem.memory_size(), hip_stream_));
-}
-
-void HIPDeviceQueue::copy_to_device(device_memory &mem)
-{
-  assert(mem.type != MEM_GLOBAL && mem.type != MEM_TEXTURE);
-
-  if (mem.memory_size() == 0) {
-    return;
-  }
-
-  /* Allocate on demand. */
-  if (mem.device_pointer == 0) {
-    hip_device_->mem_alloc(mem);
-  }
-
-  assert(mem.device_pointer != 0);
-  assert(mem.host_pointer != nullptr);
-
-  /* Copy memory to device. */
-  const HIPContextScope scope(hip_device_);
-  hip_device_assert(
-      hip_device_,
-      hipMemcpyHtoDAsync(
-          (hipDeviceptr_t)mem.device_pointer, mem.host_pointer, mem.memory_size(), hip_stream_));
-}
-
-void HIPDeviceQueue::copy_from_device(device_memory &mem)
-{
-  assert(mem.type != MEM_GLOBAL && mem.type != MEM_TEXTURE);
-
-  if (mem.memory_size() == 0) {
-    return;
-  }
-
-  assert(mem.device_pointer != 0);
-  assert(mem.host_pointer != nullptr);
-
-  /* Copy memory from device. */
-  const HIPContextScope scope(hip_device_);
-  hip_device_assert(
-      hip_device_,
-      hipMemcpyDtoHAsync(
-          mem.host_pointer, (hipDeviceptr_t)mem.device_pointer, mem.memory_size(), hip_stream_));
-}
-
-// TODO : (Arya) Enable this after stabilizing dev branch
-unique_ptr<DeviceGraphicsInterop> HIPDeviceQueue::graphics_interop_create()
-{
-  return make_unique<HIPDeviceGraphicsInterop>(this);
-}
-
-CCL_NAMESPACE_END
-
-#endif /* WITH_HIP */

intern/cycles/device/hip/queue.h

@@ -1,68 +0,0 @@
-/*
- * Copyright 2011-2021 Blender Foundation
- *
- * Licensed under the Apache License, Version 2.0 (the "License");
- * you may not use this file except in compliance with the License.
- * You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing, software
- * distributed under the License is distributed on an "AS IS" BASIS,
- * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
- * See the License for the specific language governing permissions and
- * limitations under the License.
- */
-
-#pragma once
-
-#ifdef WITH_HIP
-
-#  include "device/device_kernel.h"
-#  include "device/device_memory.h"
-#  include "device/device_queue.h"
-
-#  include "device/hip/util.h"
-
-CCL_NAMESPACE_BEGIN
-
-class HIPDevice;
-class device_memory;
-
-/* Base class for HIP queues. */
-class HIPDeviceQueue : public DeviceQueue {
- public:
-  HIPDeviceQueue(HIPDevice *device);
-  ~HIPDeviceQueue();
-
-  virtual int num_concurrent_states(const size_t state_size) const override;
-  virtual int num_concurrent_busy_states() const override;
-
-  virtual void init_execution() override;
-
-  virtual bool kernel_available(DeviceKernel kernel) const override;
-
-  virtual bool enqueue(DeviceKernel kernel, const int work_size, void *args[]) override;
-
-  virtual bool synchronize() override;
-
-  virtual void zero_to_device(device_memory &mem) override;
-  virtual void copy_to_device(device_memory &mem) override;
-  virtual void copy_from_device(device_memory &mem) override;
-
-  virtual hipStream_t stream()
-  {
-    return hip_stream_;
-  }
-
-  // TODO : (Arya) Enable this after stabilizing the dev branch
-  virtual unique_ptr<DeviceGraphicsInterop> graphics_interop_create() override;
-
- protected:
-  HIPDevice *hip_device_;
-  hipStream_t hip_stream_;
-};
-
-CCL_NAMESPACE_END
-
-#endif /* WITH_HIP */

intern/cycles/device/hip/util.cpp

@@ -1,61 +0,0 @@
-/*
- * Copyright 2011-2021 Blender Foundation
- *
- * Licensed under the Apache License, Version 2.0 (the "License");
- * you may not use this file except in compliance with the License.
- * You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing, software
- * distributed under the License is distributed on an "AS IS" BASIS,
- * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
- * See the License for the specific language governing permissions and
- * limitations under the License.
- */
-
-#ifdef WITH_HIP
-
-#  include "device/hip/util.h"
-#  include "device/hip/device_impl.h"
-
-CCL_NAMESPACE_BEGIN
-
-HIPContextScope::HIPContextScope(HIPDevice *device) : device(device)
-{
-  hip_device_assert(device, hipCtxPushCurrent(device->hipContext));
-}
-
-HIPContextScope::~HIPContextScope()
-{
-  hip_device_assert(device, hipCtxPopCurrent(NULL));
-}
-
-#  ifndef WITH_HIP_DYNLOAD
-const char *hipewErrorString(hipError_t result)
-{
-  /* We can only give error code here without major code duplication, that
-   * should be enough since dynamic loading is only being disabled by folks
-   * who knows what they're doing anyway.
-   *
-   * NOTE: Avoid call from several threads.
-   */
-  static string error;
-  error = string_printf("%d", result);
-  return error.c_str();
-}
-
-const char *hipewCompilerPath()
-{
-  return CYCLES_HIP_HIPCC_EXECUTABLE;
-}
-
-int hipewCompilerVersion()
-{
-  return (HIP_VERSION / 100) + (HIP_VERSION % 100 / 10);
-}
-#  endif
-
-CCL_NAMESPACE_END
-
-#endif /* WITH_HIP */

intern/cycles/device/hip/util.h

@@ -1,63 +0,0 @@
-/*
- * Copyright 2011-2021 Blender Foundation
- *
- * Licensed under the Apache License, Version 2.0 (the "License");
- * you may not use this file except in compliance with the License.
- * You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing, software
- * distributed under the License is distributed on an "AS IS" BASIS,
- * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
- * See the License for the specific language governing permissions and
- * limitations under the License.
- */
-
-#pragma once
-
-#ifdef WITH_HIP
-
-#  ifdef WITH_HIP_DYNLOAD
-#    include "hipew.h"
-#  endif
-
-CCL_NAMESPACE_BEGIN
-
-class HIPDevice;
-
-/* Utility to push/pop HIP context. */
-class HIPContextScope {
- public:
-  HIPContextScope(HIPDevice *device);
-  ~HIPContextScope();
-
- private:
-  HIPDevice *device;
-};
-
-/* Utility for checking return values of HIP function calls. */
-#  define hip_device_assert(hip_device, stmt) \
-    { \
-      hipError_t result = stmt; \
-      if (result != hipSuccess) { \
-        const char *name = hipewErrorString(result); \
-        hip_device->set_error( \
-            string_printf("%s in %s (%s:%d)", name, #stmt, __FILE__, __LINE__)); \
-      } \
-    } \
-    (void)0
-
-#  define hip_assert(stmt) hip_device_assert(this, stmt)
-
-#  ifndef WITH_HIP_DYNLOAD
-/* Transparently implement some functions, so majority of the file does not need
- * to worry about difference between dynamically loaded and linked HIP at all. */
-const char *hipewErrorString(hipError_t result);
-const char *hipewCompilerPath();
-int hipewCompilerVersion();
-#  endif /* WITH_HIP_DYNLOAD */
-
-CCL_NAMESPACE_END
-
-#endif /* WITH_HIP */

intern/cycles/device/optix/device_impl.cpp

@@ -315,11 +315,6 @@ bool OptiXDevice::load_kernels(const uint kernel_features)
   group_descs[PG_HITS].kind = OPTIX_PROGRAM_GROUP_KIND_HITGROUP;
   group_descs[PG_HITS].hitgroup.moduleAH = optix_module;
   group_descs[PG_HITS].hitgroup.entryFunctionNameAH = "__anyhit__kernel_optix_shadow_all_hit";
-  group_descs[PG_HITV].kind = OPTIX_PROGRAM_GROUP_KIND_HITGROUP;
-  group_descs[PG_HITV].hitgroup.moduleCH = optix_module;
-  group_descs[PG_HITV].hitgroup.entryFunctionNameCH = "__closesthit__kernel_optix_hit";
-  group_descs[PG_HITV].hitgroup.moduleAH = optix_module;
-  group_descs[PG_HITV].hitgroup.entryFunctionNameAH = "__anyhit__kernel_optix_volume_test";
 
   if (kernel_features & KERNEL_FEATURE_HAIR) {
     if (kernel_features & KERNEL_FEATURE_HAIR_THICK) {
@@ -402,7 +397,6 @@ bool OptiXDevice::load_kernels(const uint kernel_features)
   trace_css = std::max(trace_css, stack_size[PG_HITD].cssIS + stack_size[PG_HITD].cssAH);
   trace_css = std::max(trace_css, stack_size[PG_HITS].cssIS + stack_size[PG_HITS].cssAH);
   trace_css = std::max(trace_css, stack_size[PG_HITL].cssIS + stack_size[PG_HITL].cssAH);
-  trace_css = std::max(trace_css, stack_size[PG_HITV].cssIS + stack_size[PG_HITV].cssAH);
   trace_css = std::max(trace_css,
                        stack_size[PG_HITD_MOTION].cssIS + stack_size[PG_HITD_MOTION].cssAH);
   trace_css = std::max(trace_css,
@@ -427,7 +421,6 @@ bool OptiXDevice::load_kernels(const uint kernel_features)
     pipeline_groups.push_back(groups[PG_HITD]);
     pipeline_groups.push_back(groups[PG_HITS]);
     pipeline_groups.push_back(groups[PG_HITL]);
-    pipeline_groups.push_back(groups[PG_HITV]);
     if (motion_blur) {
       pipeline_groups.push_back(groups[PG_HITD_MOTION]);
       pipeline_groups.push_back(groups[PG_HITS_MOTION]);
@@ -466,7 +459,6 @@ bool OptiXDevice::load_kernels(const uint kernel_features)
     pipeline_groups.push_back(groups[PG_HITD]);
     pipeline_groups.push_back(groups[PG_HITS]);
     pipeline_groups.push_back(groups[PG_HITL]);
-    pipeline_groups.push_back(groups[PG_HITV]);
     if (motion_blur) {
       pipeline_groups.push_back(groups[PG_HITD_MOTION]);
       pipeline_groups.push_back(groups[PG_HITS_MOTION]);
@@ -1398,33 +1390,25 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
       /* Set user instance ID to object index (but leave low bit blank). */
       instance.instanceId = ob->get_device_index() << 1;
 
-      /* Add some of the object visibility bits to the mask.
-       * __prim_visibility contains the combined visibility bits of all instances, so is not
-       * reliable if they differ between instances. But the OptiX visibility mask can only contain
-       * 8 bits, so have to trade-off here and select just a few important ones.
-       */
-      instance.visibilityMask = ob->visibility_for_tracing() & 0xFF;
-
       /* Have to have at least one bit in the mask, or else instance would always be culled. */
-      if (0 == instance.visibilityMask) {
-        instance.visibilityMask = 0xFF;
+      instance.visibilityMask = 1;
+
+      if (ob->get_geometry()->has_volume) {
+        /* Volumes have a special bit set in the visibility mask so a trace can mask only volumes.
+         */
+        instance.visibilityMask |= 2;
       }
 
-      if (ob->get_geometry()->geometry_type == Geometry::HAIR &&
-          static_cast<const Hair *>(ob->get_geometry())->curve_shape == CURVE_THICK) {
-        if (motion_blur && ob->get_geometry()->has_motion_blur()) {
+      if (ob->get_geometry()->geometry_type == Geometry::HAIR) {
+        /* Same applies to curves (so they can be skipped in local trace calls). */
+        instance.visibilityMask |= 4;
+
+        if (motion_blur && ob->get_geometry()->has_motion_blur() &&
+            static_cast<const Hair *>(ob->get_geometry())->curve_shape == CURVE_THICK) {
           /* Select between motion blur and non-motion blur built-in intersection module. */
           instance.sbtOffset = PG_HITD_MOTION - PG_HITD;
         }
       }
-      else {
-        /* Can disable __anyhit__kernel_optix_visibility_test by default (except for thick curves,
-         * since it needs to filter out endcaps there).
-         * It is enabled where necessary (visibility mask exceeds 8 bits or the other any-hit
-         * programs like __anyhit__kernel_optix_shadow_all_hit) via OPTIX_RAY_FLAG_ENFORCE_ANYHIT.
-         */
-        instance.flags = OPTIX_INSTANCE_FLAG_DISABLE_ANYHIT;
-      }
 
       /* Insert motion traversable if object has motion. */
       if (motion_blur && ob->use_motion()) {
@@ -1490,7 +1474,7 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
         delete[] reinterpret_cast<uint8_t *>(&motion_transform);
 
         /* Disable instance transform if object uses motion transform already. */
-        instance.flags |= OPTIX_INSTANCE_FLAG_DISABLE_TRANSFORM;
+        instance.flags = OPTIX_INSTANCE_FLAG_DISABLE_TRANSFORM;
 
         /* Get traversable handle to motion transform. */
         optixConvertPointerToTraversableHandle(context,
@@ -1507,7 +1491,7 @@ void OptiXDevice::build_bvh(BVH *bvh, Progress &progress, bool refit)
         }
         else {
           /* Disable instance transform if geometry already has it applied to vertex data. */
-          instance.flags |= OPTIX_INSTANCE_FLAG_DISABLE_TRANSFORM;
+          instance.flags = OPTIX_INSTANCE_FLAG_DISABLE_TRANSFORM;
           /* Non-instanced objects read ID from 'prim_object', so distinguish
            * them from instanced objects with the low bit set. */
           instance.instanceId |= 1;

intern/cycles/device/optix/device_impl.h

@@ -40,7 +40,6 @@ enum {
   PG_HITD, /* Default hit group. */
   PG_HITS, /* __SHADOW_RECORD_ALL__ hit group. */
   PG_HITL, /* __BVH_LOCAL__ hit group (only used for triangles). */
-  PG_HITV, /* __VOLUME__ hit group. */
   PG_HITD_MOTION,
   PG_HITS_MOTION,
   PG_CALL_SVM_AO,
@@ -52,7 +51,7 @@ enum {
 static const int MISS_PROGRAM_GROUP_OFFSET = PG_MISS;
 static const int NUM_MIS_PROGRAM_GROUPS = 1;
 static const int HIT_PROGAM_GROUP_OFFSET = PG_HITD;
-static const int NUM_HIT_PROGRAM_GROUPS = 6;
+static const int NUM_HIT_PROGRAM_GROUPS = 5;
 static const int CALLABLE_PROGRAM_GROUPS_BASE = PG_CALL_SVM_AO;
 static const int NUM_CALLABLE_PROGRAM_GROUPS = 3;
 

intern/cycles/integrator/path_trace.cpp

@@ -801,7 +801,7 @@ void PathTrace::tile_buffer_write_to_disk()
   }
 
   if (!tile_manager_.write_tile(*buffers)) {
-    device_->set_error("Error writing tile to file");
+    LOG(ERROR) << "Error writing tile to file.";
   }
 }
 
@@ -894,14 +894,7 @@ void PathTrace::process_full_buffer_from_disk(string_view filename)
 
   DenoiseParams denoise_params;
   if (!tile_manager_.read_full_buffer_from_disk(filename, &full_frame_buffers, &denoise_params)) {
-    const string error_message = "Error reading tiles from file";
-    if (progress_) {
-      progress_->set_error(error_message);
-      progress_->set_cancel(error_message);
-    }
-    else {
-      LOG(ERROR) << error_message;
-    }
+    LOG(ERROR) << "Error reading tiles from file.";
     return;
   }
 
@@ -1035,8 +1028,6 @@ static const char *device_type_for_description(const DeviceType type)
       return "CUDA";
     case DEVICE_OPTIX:
       return "OptiX";
-    case DEVICE_HIP:
-      return "HIP";
     case DEVICE_DUMMY:
       return "Dummy";
     case DEVICE_MULTI:

intern/cycles/integrator/path_trace_work_cpu.cpp

@@ -19,8 +19,6 @@
 #include "device/cpu/kernel.h"
 #include "device/device.h"
 
-#include "kernel/kernel_path_state.h"
-
 #include "integrator/pass_accessor_cpu.h"
 
 #include "render/buffers.h"
@@ -118,17 +116,13 @@ void PathTraceWorkCPU::render_samples_full_pipeline(KernelGlobals *kernel_global
                                                     const KernelWorkTile &work_tile,
                                                     const int samples_num)
 {
+  const bool has_shadow_catcher = device_scene_->data.integrator.has_shadow_catcher;
   const bool has_bake = device_scene_->data.bake.use;
 
-  IntegratorStateCPU integrator_states[2];
+  IntegratorStateCPU integrator_states[2] = {};
 
   IntegratorStateCPU *state = &integrator_states[0];
-  IntegratorStateCPU *shadow_catcher_state = nullptr;
-
-  if (device_scene_->data.integrator.has_shadow_catcher) {
-    shadow_catcher_state = &integrator_states[1];
-    path_state_init_queues(kernel_globals, shadow_catcher_state);
-  }
+  IntegratorStateCPU *shadow_catcher_state = &integrator_states[1];
 
   KernelWorkTile sample_work_tile = work_tile;
   float *render_buffer = buffers_->buffer.data();
@@ -153,7 +147,7 @@ void PathTraceWorkCPU::render_samples_full_pipeline(KernelGlobals *kernel_global
 
     kernels_.integrator_megakernel(kernel_globals, state, render_buffer);
 
-    if (shadow_catcher_state) {
+    if (has_shadow_catcher) {
       kernels_.integrator_megakernel(kernel_globals, shadow_catcher_state, render_buffer);
     }
 

intern/cycles/integrator/path_trace_work_gpu.cpp

@@ -95,8 +95,8 @@ void PathTraceWorkGPU::alloc_integrator_soa()
 #define KERNEL_STRUCT_END(name) \
   break; \
   }
-#define KERNEL_STRUCT_END_ARRAY(name, cpu_array_size, gpu_array_size) \
-  if (array_index == gpu_array_size - 1) { \
+#define KERNEL_STRUCT_END_ARRAY(name, array_size) \
+  if (array_index == array_size - 1) { \
     break; \
   } \
   }
@@ -738,8 +738,7 @@ void PathTraceWorkGPU::copy_to_gpu_display_naive(GPUDisplay *gpu_display,
 
   get_render_tile_film_pixels(destination, pass_mode, num_samples);
 
-  queue_->copy_from_device(gpu_display_rgba_half_);
-  queue_->synchronize();
+  gpu_display_rgba_half_.copy_from_device();
 
   gpu_display->copy_pixels_to_texture(
       gpu_display_rgba_half_.data(), texture_x, texture_y, width, height);

intern/cycles/integrator/render_scheduler.cpp

@@ -384,7 +384,7 @@ bool RenderScheduler::set_postprocess_render_work(RenderWork *render_work)
   }
 
   if (denoiser_params_.use && !state_.last_work_tile_was_denoised) {
-    render_work->tile.denoise = !tile_manager_.has_multiple_tiles();
+    render_work->tile.denoise = true;
     any_scheduled = true;
   }
 
@@ -903,12 +903,6 @@ bool RenderScheduler::work_need_denoise(bool &delayed, bool &ready_to_display)
     return false;
   }
 
-  /* When multiple tiles are used the full frame will be denoised.
-   * Avoid per-tile denoising to save up render time. */
-  if (tile_manager_.has_multiple_tiles()) {
-    return false;
-  }
-
   if (done()) {
     /* Always denoise at the last sample. */
     return true;

intern/cycles/kernel/CMakeLists.txt

@@ -35,10 +35,6 @@ set(SRC_DEVICE_CUDA
   device/cuda/kernel.cu
 )
 
-set(SRC_DEVICE_HIP
-  device/hip/kernel.cpp
-)
-
 set(SRC_DEVICE_OPTIX
   device/optix/kernel.cu
   device/optix/kernel_shader_raytrace.cu
@@ -110,12 +106,6 @@ set(SRC_DEVICE_CUDA_HEADERS
   device/cuda/globals.h
 )
 
-set(SRC_DEVICE_HIP_HEADERS
-  device/hip/compat.h
-  device/hip/config.h
-  device/hip/globals.h
-)
-
 set(SRC_DEVICE_OPTIX_HEADERS
   device/optix/compat.h
   device/optix/globals.h
@@ -468,104 +458,6 @@ if(WITH_CYCLES_CUDA_BINARIES)
   cycles_set_solution_folder(cycles_kernel_cuda)
 endif()
 
-####################################################### START
-
-# HIP module
-
-if(WITH_CYCLES_HIP_BINARIES)
-  # 64 bit only
-  set(HIP_BITS 64)
-
-  # HIP version
-  execute_process(COMMAND ${HIP_HIPCC_EXECUTABLE} "--version" OUTPUT_VARIABLE HIPCC_OUT)
-  string(REGEX REPLACE ".*release ([0-9]+)\\.([0-9]+).*" "\\1" HIP_VERSION_MAJOR "${HIPCC_OUT}")
-  string(REGEX REPLACE ".*release ([0-9]+)\\.([0-9]+).*" "\\2" HIP_VERSION_MINOR "${HIPCC_OUT}")
-  set(HIP_VERSION "${HIP_VERSION_MAJOR}${HIP_VERSION_MINOR}")
-
-
-  message(WARNING
-    "HIP version ${HIP_VERSION_MAJOR}.${HIP_VERSION_MINOR} detected")
-
-  # build for each arch
-  set(hip_sources device/hip/kernel.cpp
-    ${SRC_HEADERS}
-    ${SRC_DEVICE_HIP_HEADERS}
-    ${SRC_BVH_HEADERS}
-    ${SRC_SVM_HEADERS}
-    ${SRC_GEOM_HEADERS}
-    ${SRC_INTEGRATOR_HEADERS}
-    ${SRC_CLOSURE_HEADERS}
-    ${SRC_UTIL_HEADERS}
-  )
-  set(hip_fatbins)
-
-  macro(CYCLES_HIP_KERNEL_ADD arch prev_arch name flags sources experimental)
-    if(${arch} MATCHES "compute_.*")
-      set(format "ptx")
-    else()
-      set(format "fatbin")
-    endif()
-    set(hip_file ${name}_${arch}.${format})
-
-    set(kernel_sources ${sources})
-    if(NOT ${prev_arch} STREQUAL "none")
-      if(${prev_arch} MATCHES "compute_.*")
-        set(kernel_sources ${kernel_sources} ${name}_${prev_arch}.ptx)
-      else()
-        set(kernel_sources ${kernel_sources} ${name}_${prev_arch}.fatbin)
-      endif()
-    endif()
-
-    set(hip_kernel_src "/device/hip/${name}.cpp")
-
-    set(hip_flags ${flags}
-      -D CCL_NAMESPACE_BEGIN=
-      -D CCL_NAMESPACE_END=
-      -D HIPCC
-      -m ${HIP_BITS}
-      -I ${CMAKE_CURRENT_SOURCE_DIR}/..
-      -I ${CMAKE_CURRENT_SOURCE_DIR}/device/hip
-      --use_fast_math
-      -o ${CMAKE_CURRENT_BINARY_DIR}/${hip_file})
-
-    if(${experimental})
-      set(hip_flags ${hip_flags} -D __KERNEL_EXPERIMENTAL__)
-      set(name ${name}_experimental)
-    endif()
-
-    if(WITH_CYCLES_DEBUG)
-      set(hip_flags ${hip_flags} -D __KERNEL_DEBUG__)
-    endif()
-
-    if(WITH_NANOVDB)
-      set(hip_flags ${hip_flags}
-        -D WITH_NANOVDB
-        -I "${NANOVDB_INCLUDE_DIR}")
-    endif()
-  endmacro()
-
-  set(prev_arch "none")
-  foreach(arch ${CYCLES_HIP_BINARIES_ARCH})
-      set(hip_hipcc_executable ${HIP_HIPCC_EXECUTABLE})
-      set(hip_toolkit_root_dir ${HIP_TOOLKIT_ROOT_DIR})
-    if(DEFINED hip_hipcc_executable AND DEFINED hip_toolkit_root_dir)
-      # Compile regular kernel
-      CYCLES_HIP_KERNEL_ADD(${arch} ${prev_arch} kernel "" "${hip_sources}" FALSE)
-
-      if(WITH_CYCLES_HIP_BUILD_SERIAL)
-        set(prev_arch ${arch})
-      endif()
-
-      unset(hip_hipcc_executable)
-      unset(hip_toolkit_root_dir)
-    endif()
-  endforeach()
-
-  add_custom_target(cycles_kernel_hip ALL DEPENDS ${hip_fatbins})
-  cycles_set_solution_folder(cycles_kernel_hip)
-endif()
-
-####################################################### END
 # OptiX PTX modules
 
 if(WITH_CYCLES_DEVICE_OPTIX AND WITH_CYCLES_CUDA_BINARIES)
@@ -710,13 +602,11 @@ endif()
 cycles_add_library(cycles_kernel "${LIB}"
   ${SRC_DEVICE_CPU}
   ${SRC_DEVICE_CUDA}
-  ${SRC_DEVICE_HIP}
   ${SRC_DEVICE_OPTIX}
   ${SRC_HEADERS}
   ${SRC_DEVICE_CPU_HEADERS}
   ${SRC_DEVICE_GPU_HEADERS}
   ${SRC_DEVICE_CUDA_HEADERS}
-  ${SRC_DEVICE_HIP_HEADERS}
   ${SRC_DEVICE_OPTIX_HEADERS}
   ${SRC_BVH_HEADERS}
   ${SRC_CLOSURE_HEADERS}
@@ -731,7 +621,6 @@ source_group("geom" FILES ${SRC_GEOM_HEADERS})
 source_group("integrator" FILES ${SRC_INTEGRATOR_HEADERS})
 source_group("kernel" FILES ${SRC_HEADERS})
 source_group("device\\cpu" FILES ${SRC_DEVICE_CPU} ${SRC_DEVICE_CPU_HEADERS})
-source_group("device\\hip" FILES ${SRC_DEVICE_HIP} ${SRC_DEVICE_HIP_HEADERS})
 source_group("device\\gpu" FILES ${SRC_DEVICE_GPU_HEADERS})
 source_group("device\\cuda" FILES ${SRC_DEVICE_CUDA} ${SRC_DEVICE_CUDA_HEADERS})
 source_group("device\\optix" FILES ${SRC_DEVICE_OPTIX} ${SRC_DEVICE_OPTIX_HEADERS})
@@ -743,19 +632,14 @@ endif()
 if(WITH_CYCLES_DEVICE_OPTIX AND WITH_CYCLES_CUDA_BINARIES)
   add_dependencies(cycles_kernel cycles_kernel_optix)
 endif()
-if(WITH_CYCLES_HIP)
-  add_dependencies(cycles_kernel cycles_kernel_hip)
-endif()
 
 # Install kernel source for runtime compilation
 
 delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_DEVICE_CUDA}" ${CYCLES_INSTALL_PATH}/source/kernel/device/cuda)
-delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_DEVICE_HIP}" ${CYCLES_INSTALL_PATH}/source/kernel/device/hip)
 delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_DEVICE_OPTIX}" ${CYCLES_INSTALL_PATH}/source/kernel/device/optix)
 delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_HEADERS}" ${CYCLES_INSTALL_PATH}/source/kernel)
 delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_DEVICE_GPU_HEADERS}" ${CYCLES_INSTALL_PATH}/source/kernel/device/gpu)
 delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_DEVICE_CUDA_HEADERS}" ${CYCLES_INSTALL_PATH}/source/kernel/device/cuda)
-delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_DEVICE_HIP_HEADERS}" ${CYCLES_INSTALL_PATH}/source/kernel/device/hip)
 delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_DEVICE_OPTIX_HEADERS}" ${CYCLES_INSTALL_PATH}/source/kernel/device/optix)
 delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_BVH_HEADERS}" ${CYCLES_INSTALL_PATH}/source/kernel/bvh)
 delayed_install(${CMAKE_CURRENT_SOURCE_DIR} "${SRC_CLOSURE_HEADERS}" ${CYCLES_INSTALL_PATH}/source/kernel/closure)

intern/cycles/kernel/bvh/bvh.h

@@ -167,25 +167,15 @@ ccl_device_intersect bool scene_intersect(const KernelGlobals *kg,
   uint p4 = visibility;
   uint p5 = PRIMITIVE_NONE;
 
-  uint ray_mask = visibility & 0xFF;
-  uint ray_flags = OPTIX_RAY_FLAG_NONE;
-  if (0 == ray_mask && (visibility & ~0xFF) != 0) {
-    ray_mask = 0xFF;
-    ray_flags = OPTIX_RAY_FLAG_ENFORCE_ANYHIT;
-  }
-  else if (visibility & PATH_RAY_SHADOW_OPAQUE) {
-    ray_flags = OPTIX_RAY_FLAG_TERMINATE_ON_FIRST_HIT;
-  }
-
   optixTrace(scene_intersect_valid(ray) ? kernel_data.bvh.scene : 0,
              ray->P,
              ray->D,
              0.0f,
              ray->t,
              ray->time,
-             ray_mask,
-             ray_flags,
-             0, /* SBT offset for PG_HITD */
+             0xF,
+             OPTIX_RAY_FLAG_NONE,
+             0,  // SBT offset for PG_HITD
              0,
              0,
              p0,
@@ -261,11 +251,11 @@ ccl_device_intersect bool scene_intersect_local(const KernelGlobals *kg,
   uint p2 = ((uint64_t)local_isect) & 0xFFFFFFFF;
   uint p3 = (((uint64_t)local_isect) >> 32) & 0xFFFFFFFF;
   uint p4 = local_object;
-  /* Is set to zero on miss or if ray is aborted, so can be used as return value. */
+  // Is set to zero on miss or if ray is aborted, so can be used as return value
   uint p5 = max_hits;
 
   if (local_isect) {
-    local_isect->num_hits = 0; /* Initialize hit count to zero. */
+    local_isect->num_hits = 0;  // Initialize hit count to zero
   }
   optixTrace(scene_intersect_valid(ray) ? kernel_data.bvh.scene : 0,
              ray->P,
@@ -273,10 +263,11 @@ ccl_device_intersect bool scene_intersect_local(const KernelGlobals *kg,
              0.0f,
              ray->t,
              ray->time,
-             0xFF,
-             /* Need to always call into __anyhit__kernel_optix_local_hit. */
+             // Skip curves
+             0x3,
+             // Need to always call into __anyhit__kernel_optix_local_hit
              OPTIX_RAY_FLAG_ENFORCE_ANYHIT,
-             2, /* SBT offset for PG_HITL */
+             2,  // SBT offset for PG_HITL
              0,
              0,
              p0,
@@ -374,22 +365,17 @@ ccl_device_intersect bool scene_intersect_shadow_all(const KernelGlobals *kg,
   uint p4 = visibility;
   uint p5 = false;
 
-  uint ray_mask = visibility & 0xFF;
-  if (0 == ray_mask && (visibility & ~0xFF) != 0) {
-    ray_mask = 0xFF;
-  }
-
-  *num_hits = 0; /* Initialize hit count to zero. */
+  *num_hits = 0;  // Initialize hit count to zero
   optixTrace(scene_intersect_valid(ray) ? kernel_data.bvh.scene : 0,
              ray->P,
              ray->D,
              0.0f,
              ray->t,
              ray->time,
-             ray_mask,
-             /* Need to always call into __anyhit__kernel_optix_shadow_all_hit. */
+             0xF,
+             // Need to always call into __anyhit__kernel_optix_shadow_all_hit
              OPTIX_RAY_FLAG_ENFORCE_ANYHIT,
-             1, /* SBT offset for PG_HITS */
+             1,  // SBT offset for PG_HITS
              0,
              0,
              p0,
@@ -458,21 +444,16 @@ ccl_device_intersect bool scene_intersect_volume(const KernelGlobals *kg,
   uint p4 = visibility;
   uint p5 = PRIMITIVE_NONE;
 
-  uint ray_mask = visibility & 0xFF;
-  if (0 == ray_mask && (visibility & ~0xFF) != 0) {
-    ray_mask = 0xFF;
-  }
-
   optixTrace(scene_intersect_valid(ray) ? kernel_data.bvh.scene : 0,
              ray->P,
              ray->D,
              0.0f,
              ray->t,
              ray->time,
-             ray_mask,
-             /* Need to always call into __anyhit__kernel_optix_volume_test. */
-             OPTIX_RAY_FLAG_ENFORCE_ANYHIT,
-             3, /* SBT offset for PG_HITV */
+             // Skip everything but volumes
+             0x2,
+             OPTIX_RAY_FLAG_NONE,
+             0,  // SBT offset for PG_HITD
              0,
              0,
              p0,

intern/cycles/kernel/device/gpu/parallel_active_index.h

@@ -25,11 +25,7 @@ CCL_NAMESPACE_BEGIN
 
 #include "util/util_atomic.h"
 
-#ifdef __HIP__
-#  define GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE 1024
-#else
-#  define GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE 512
-#endif
+#define GPU_PARALLEL_ACTIVE_INDEX_DEFAULT_BLOCK_SIZE 512
 
 template<uint blocksize, typename IsActiveOp>
 __device__ void gpu_parallel_active_index_array(const uint num_states,

intern/cycles/kernel/device/gpu/parallel_prefix_sum.h

@@ -27,11 +27,7 @@ CCL_NAMESPACE_BEGIN
 
 #include "util/util_atomic.h"
 
-#ifdef __HIP__
-#  define GPU_PARALLEL_PREFIX_SUM_DEFAULT_BLOCK_SIZE 1024
-#else
-#  define GPU_PARALLEL_PREFIX_SUM_DEFAULT_BLOCK_SIZE 512
-#endif
+#define GPU_PARALLEL_PREFIX_SUM_DEFAULT_BLOCK_SIZE 512
 
 template<uint blocksize> __device__ void gpu_parallel_prefix_sum(int *values, const int num_values)
 {

intern/cycles/kernel/device/gpu/parallel_reduce.h

@@ -26,11 +26,7 @@ CCL_NAMESPACE_BEGIN
  * the overall cost of the algorithm while keeping the work complexity O(n) and
  * the step complexity O(log n). (Brent's Theorem optimization) */
 
-#ifdef __HIP__
-#  define GPU_PARALLEL_SUM_DEFAULT_BLOCK_SIZE 1024
-#else
-#  define GPU_PARALLEL_SUM_DEFAULT_BLOCK_SIZE 512
-#endif
+#define GPU_PARALLEL_SUM_DEFAULT_BLOCK_SIZE 512
 
 template<uint blocksize, typename InputT, typename OutputT, typename ConvertOp>
 __device__ void gpu_parallel_sum(

intern/cycles/kernel/device/gpu/parallel_sorted_index.h

@@ -26,11 +26,7 @@ CCL_NAMESPACE_BEGIN
 
 #include "util/util_atomic.h"
 
-#ifdef __HIP__
-#  define GPU_PARALLEL_SORTED_INDEX_DEFAULT_BLOCK_SIZE 1024
-#else
-#  define GPU_PARALLEL_SORTED_INDEX_DEFAULT_BLOCK_SIZE 512
-#endif
+#define GPU_PARALLEL_SORTED_INDEX_DEFAULT_BLOCK_SIZE 512
 #define GPU_PARALLEL_SORTED_INDEX_INACTIVE_KEY (~0)
 
 template<uint blocksize, typename GetKeyOp>

intern/cycles/kernel/device/hip/compat.h

@@ -1,121 +0,0 @@
-/*
- * Copyright 2011-2021 Blender Foundation
- *
- * Licensed under the Apache License, Version 2.0 (the "License");
- * you may not use this file except in compliance with the License.
- * You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing, software
- * distributed under the License is distributed on an "AS IS" BASIS,
- * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
- * See the License for the specific language governing permissions and
- * limitations under the License.
- */
-
-#pragma once
-
-#define __KERNEL_GPU__
-#define __KERNEL_HIP__
-#define CCL_NAMESPACE_BEGIN
-#define CCL_NAMESPACE_END
-
-#ifndef ATTR_FALLTHROUGH
-#  define ATTR_FALLTHROUGH
-#endif
-
-#ifdef __HIPCC_RTC__
-typedef unsigned int uint32_t;
-typedef unsigned long long uint64_t;
-#else
-#  include <stdint.h>
-#endif
-
-#ifdef CYCLES_HIPBIN_CC
-#  define FLT_MIN 1.175494350822287507969e-38f
-#  define FLT_MAX 340282346638528859811704183484516925440.0f
-#  define FLT_EPSILON 1.192092896e-07F
-#endif
-
-/* Qualifiers */
-
-#define ccl_device __device__ __inline__
-#define ccl_device_inline __device__ __inline__
-#define ccl_device_forceinline __device__ __forceinline__
-#define ccl_device_noinline __device__ __noinline__
-#define ccl_device_noinline_cpu ccl_device
-#define ccl_global
-#define ccl_static_constant __constant__
-#define ccl_device_constant __constant__ __device__
-#define ccl_constant const
-#define ccl_gpu_shared __shared__
-#define ccl_private
-#define ccl_may_alias
-#define ccl_addr_space
-#define ccl_restrict __restrict__
-#define ccl_loop_no_unroll
-#define ccl_align(n) __align__(n)
-#define ccl_optional_struct_init
-
-#define kernel_assert(cond)
-
-/* Types */
-#ifdef __HIP__
-#  include "hip/hip_fp16.h"
-#  include "hip/hip_runtime.h"
-#endif
-
-#ifdef _MSC_VER
-#  include <immintrin.h>
-#endif
-
-#define ccl_gpu_thread_idx_x (threadIdx.x)
-#define ccl_gpu_block_dim_x (blockDim.x)
-#define ccl_gpu_block_idx_x (blockIdx.x)
-#define ccl_gpu_grid_dim_x (gridDim.x)
-#define ccl_gpu_warp_size (warpSize)
-
-#define ccl_gpu_global_id_x() (ccl_gpu_block_idx_x * ccl_gpu_block_dim_x + ccl_gpu_thread_idx_x)
-#define ccl_gpu_global_size_x() (ccl_gpu_grid_dim_x * ccl_gpu_block_dim_x)
-
-/* GPU warp synchronization */
-
-#define ccl_gpu_syncthreads() __syncthreads()
-#define ccl_gpu_ballot(predicate) __ballot(predicate)
-#define ccl_gpu_shfl_down_sync(mask, var, detla) __shfl_down(var, detla)
-#define ccl_gpu_popc(x) __popc(x)
-
-/* GPU texture objects */
-typedef hipTextureObject_t ccl_gpu_tex_object;
-
-template<typename T>
-ccl_device_forceinline T ccl_gpu_tex_object_read_2D(const ccl_gpu_tex_object texobj,
-                                                    const float x,
-                                                    const float y)
-{
-  return tex2D<T>(texobj, x, y);
-}
-
-template<typename T>
-ccl_device_forceinline T ccl_gpu_tex_object_read_3D(const ccl_gpu_tex_object texobj,
-                                                    const float x,
-                                                    const float y,
-                                                    const float z)
-{
-  return tex3D<T>(texobj, x, y, z);
-}
-
-/* Use fast math functions */
-
-#define cosf(x) __cosf(((float)(x)))
-#define sinf(x) __sinf(((float)(x)))
-#define powf(x, y) __powf(((float)(x)), ((float)(y)))
-#define tanf(x) __tanf(((float)(x)))
-#define logf(x) __logf(((float)(x)))
-#define expf(x) __expf(((float)(x)))
-
-/* Types */
-
-#include "util/util_half.h"
-#include "util/util_types.h"

intern/cycles/kernel/device/hip/config.h

@@ -1,57 +0,0 @@
-/*
- * Copyright 2011-2021 Blender Foundation
- *
- * Licensed under the Apache License, Version 2.0 (the "License");
- * you may not use this file except in compliance with the License.
- * You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing, software
- * distributed under the License is distributed on an "AS IS" BASIS,
- * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
- * See the License for the specific language governing permissions and
- * limitations under the License.
- */
-
-/* Device data taken from HIP occupancy calculator.
- *
- * Terminology
- * - HIP GPUs have multiple streaming multiprocessors
- * - Each multiprocessor executes multiple thread blocks
- * - Each thread block contains a number of threads, also known as the block size
- * - Multiprocessors have a fixed number of registers, and the amount of registers
- *   used by each threads limits the number of threads per block.
- */
-
-/* Launch Bound Definitions */
-#define GPU_MULTIPRESSOR_MAX_REGISTERS 65536
-#define GPU_MULTIPROCESSOR_MAX_BLOCKS 64
-#define GPU_BLOCK_MAX_THREADS 1024
-#define GPU_THREAD_MAX_REGISTERS 255
-
-#define GPU_KERNEL_BLOCK_NUM_THREADS 1024
-#define GPU_KERNEL_MAX_REGISTERS 64
-
-/* Compute number of threads per block and minimum blocks per multiprocessor
- * given the maximum number of registers per thread. */
-
-#define ccl_gpu_kernel(block_num_threads, thread_num_registers) \
-  extern "C" __global__ void __launch_bounds__(block_num_threads, \
-                                               GPU_MULTIPRESSOR_MAX_REGISTERS / \
-                                                   (block_num_threads * thread_num_registers))
-
-/* sanity checks */
-
-#if GPU_KERNEL_BLOCK_NUM_THREADS > GPU_BLOCK_MAX_THREADS
-#  error "Maximum number of threads per block exceeded"
-#endif
-
-#if GPU_MULTIPRESSOR_MAX_REGISTERS / (GPU_KERNEL_BLOCK_NUM_THREADS * GPU_KERNEL_MAX_REGISTERS) > \
-    GPU_MULTIPROCESSOR_MAX_BLOCKS
-#  error "Maximum number of blocks per multiprocessor exceeded"
-#endif
-
-#if GPU_KERNEL_MAX_REGISTERS > GPU_THREAD_MAX_REGISTERS
-#  error "Maximum number of registers per thread exceeded"
-#endif

intern/cycles/kernel/device/hip/globals.h

@@ -1,49 +0,0 @@
-/*
- * Copyright 2011-2021 Blender Foundation
- *
- * Licensed under the Apache License, Version 2.0 (the "License");
- * you may not use this file except in compliance with the License.
- * You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing, software
- * distributed under the License is distributed on an "AS IS" BASIS,
- * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
- * See the License for the specific language governing permissions and
- * limitations under the License.
- */
-
-/* Constant Globals */
-
-#pragma once
-
-#include "kernel/kernel_profiling.h"
-#include "kernel/kernel_types.h"
-
-#include "kernel/integrator/integrator_state.h"
-
-CCL_NAMESPACE_BEGIN
-
-/* Not actually used, just a NULL pointer that gets passed everywhere, which we
- * hope gets optimized out by the compiler. */
-struct KernelGlobals {
-  /* NOTE: Keep the size in sync with SHADOW_STACK_MAX_HITS. */
-  int unused[1];
-};
-
-/* Global scene data and textures */
-__constant__ KernelData __data;
-#define KERNEL_TEX(type, name) __attribute__((used)) const __constant__ __device__ type *name;
-#include "kernel/kernel_textures.h"
-
-/* Integrator state */
-__constant__ IntegratorStateGPU __integrator_state;
-
-/* Abstraction macros */
-#define kernel_data __data
-#define kernel_tex_fetch(t, index) t[(index)]
-#define kernel_tex_array(t) (t)
-#define kernel_integrator_state __integrator_state
-
-CCL_NAMESPACE_END

intern/cycles/kernel/device/hip/kernel.cpp

@@ -1,28 +0,0 @@
-/*
- * Copyright 2011-2021 Blender Foundation
- *
- * Licensed under the Apache License, Version 2.0 (the "License");
- * you may not use this file except in compliance with the License.
- * You may obtain a copy of the License at
- *
- * http://www.apache.org/licenses/LICENSE-2.0
- *
- * Unless required by applicable law or agreed to in writing, software
- * distributed under the License is distributed on an "AS IS" BASIS,
- * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
- * See the License for the specific language governing permissions and
- * limitations under the License.
- */
-
-/* HIP kernel entry points */
-
-#ifdef __HIP_DEVICE_COMPILE__
-
-#  include "kernel/device/hip/compat.h"
-#  include "kernel/device/hip/config.h"
-#  include "kernel/device/hip/globals.h"
-
-#  include "kernel/device/gpu/image.h"
-#  include "kernel/device/gpu/kernel.h"
-
-#endif

intern/cycles/kernel/device/optix/kernel.cu

@@ -19,7 +19,7 @@
 #include "kernel/device/optix/compat.h"
 #include "kernel/device/optix/globals.h"
 
-#include "kernel/device/gpu/image.h"  /* Texture lookup uses normal CUDA intrinsics. */
+#include "kernel/device/gpu/image.h"  // Texture lookup uses normal CUDA intrinsics
 
 #include "kernel/integrator/integrator_state.h"
 #include "kernel/integrator/integrator_state_flow.h"
@@ -44,18 +44,18 @@ template<typename T> ccl_device_forceinline T *get_payload_ptr_2()
 template<bool always = false> ccl_device_forceinline uint get_object_id()
 {
 #ifdef __OBJECT_MOTION__
-  /* Always get the the instance ID from the TLAS.
-   * There might be a motion transform node between TLAS and BLAS which does not have one. */
+  // Always get the the instance ID from the TLAS
+  // There might be a motion transform node between TLAS and BLAS which does not have one
   uint object = optixGetInstanceIdFromHandle(optixGetTransformListHandle(0));
 #else
   uint object = optixGetInstanceId();
 #endif
-  /* Choose between always returning object ID or only for instances. */
+  // Choose between always returning object ID or only for instances
   if (always || (object & 1) == 0)
-    /* Can just remove the low bit since instance always contains object ID. */
+    // Can just remove the low bit since instance always contains object ID
     return object >> 1;
   else
-    /* Set to OBJECT_NONE if this is not an instanced object. */
+    // Set to OBJECT_NONE if this is not an instanced object
     return OBJECT_NONE;
 }
 
@@ -93,30 +93,23 @@ extern "C" __global__ void __raygen__kernel_optix_integrator_intersect_volume_st
 
 extern "C" __global__ void __miss__kernel_optix_miss()
 {
-  /* 'kernel_path_lamp_emission' checks intersection distance, so need to set it even on a miss. */
+  // 'kernel_path_lamp_emission' checks intersection distance, so need to set it even on a miss
   optixSetPayload_0(__float_as_uint(optixGetRayTmax()));
   optixSetPayload_5(PRIMITIVE_NONE);
 }
 
 extern "C" __global__ void __anyhit__kernel_optix_local_hit()
 {
-#ifdef __HAIR__
-  if (!optixIsTriangleHit()) {
-    /* Ignore curves. */
-    return optixIgnoreIntersection();
-  }
-#endif
-
 #ifdef __BVH_LOCAL__
   const uint object = get_object_id<true>();
   if (object != optixGetPayload_4() /* local_object */) {
-    /* Only intersect with matching object. */
+    // Only intersect with matching object
     return optixIgnoreIntersection();
   }
 
   const uint max_hits = optixGetPayload_5();
   if (max_hits == 0) {
-    /* Special case for when no hit information is requested, just report that something was hit */
+    // Special case for when no hit information is requested, just report that something was hit
     optixSetPayload_5(true);
     return optixTerminateRay();
   }
@@ -143,9 +136,8 @@ extern "C" __global__ void __anyhit__kernel_optix_local_hit()
   }
   else {
     if (local_isect->num_hits && optixGetRayTmax() > local_isect->hits[0].t) {
-      /* Record closest intersection only.
-       * Do not terminate ray here, since there is no guarantee about distance ordering in any-hit.
-       */
+      // Record closest intersection only
+      // Do not terminate ray here, since there is no guarantee about distance ordering in any-hit
       return optixIgnoreIntersection();
     }
 
@@ -162,14 +154,14 @@ extern "C" __global__ void __anyhit__kernel_optix_local_hit()
   isect->u = 1.0f - barycentrics.y - barycentrics.x;
   isect->v = barycentrics.x;
 
-  /* Record geometric normal. */
+  // Record geometric normal
   const uint tri_vindex = kernel_tex_fetch(__prim_tri_index, isect->prim);
   const float3 tri_a = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex + 0));
   const float3 tri_b = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex + 1));
   const float3 tri_c = float4_to_float3(kernel_tex_fetch(__prim_tri_verts, tri_vindex + 2));
   local_isect->Ng[hit] = normalize(cross(tri_b - tri_a, tri_c - tri_a));
 
-  /* Continue tracing (without this the trace call would return after the first hit). */
+  // Continue tracing (without this the trace call would return after the first hit)
   optixIgnoreIntersection();
 #endif
 }
@@ -198,7 +190,7 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit()
     u = __uint_as_float(optixGetAttribute_0());
     v = __uint_as_float(optixGetAttribute_1());
 
-    /* Filter out curve endcaps. */
+    // Filter out curve endcaps
     if (u == 0.0f || u == 1.0f) {
       ignore_intersection = true;
     }
@@ -249,10 +241,10 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit()
     isect->type = kernel_tex_fetch(__prim_type, prim);
 
 #  ifdef __TRANSPARENT_SHADOWS__
-    /* Detect if this surface has a shader with transparent shadows. */
+    // Detect if this surface has a shader with transparent shadows
     if (!shader_transparent_shadow(NULL, isect) || max_hits == 0) {
 #  endif
-      /* If no transparent shadows, all light is blocked and we can stop immediately. */
+      // If no transparent shadows, all light is blocked and we can stop immediately
       optixSetPayload_5(true);
       return optixTerminateRay();
 #  ifdef __TRANSPARENT_SHADOWS__
@@ -260,39 +252,24 @@ extern "C" __global__ void __anyhit__kernel_optix_shadow_all_hit()
 #  endif
   }
 
-  /* Continue tracing. */
+  // Continue tracing
   optixIgnoreIntersection();
 #endif
 }
 
-extern "C" __global__ void __anyhit__kernel_optix_volume_test()
+extern "C" __global__ void __anyhit__kernel_optix_visibility_test()
 {
-#ifdef __HAIR__
-  if (!optixIsTriangleHit()) {
-    /* Ignore curves. */
-    return optixIgnoreIntersection();
-  }
-#endif
-
+  uint visibility = optixGetPayload_4();
 #ifdef __VISIBILITY_FLAG__
   const uint prim = optixGetPrimitiveIndex();
-  const uint visibility = optixGetPayload_4();
   if ((kernel_tex_fetch(__prim_visibility, prim) & visibility) == 0) {
     return optixIgnoreIntersection();
   }
 #endif
 
-  const uint object = get_object_id<true>();
-  if ((kernel_tex_fetch(__object_flag, object) & SD_OBJECT_HAS_VOLUME) == 0) {
-    return optixIgnoreIntersection();
-  }
-}
-
-extern "C" __global__ void __anyhit__kernel_optix_visibility_test()
-{
 #ifdef __HAIR__
   if (!optixIsTriangleHit()) {
-    /* Filter out curve endcaps. */
+    // Filter out curve endcaps
     const float u = __uint_as_float(optixGetAttribute_0());
     if (u == 0.0f || u == 1.0f) {
       return optixIgnoreIntersection();
@@ -300,26 +277,18 @@ extern "C" __global__ void __anyhit__kernel_optix_visibility_test()
   }
 #endif
 
-#ifdef __VISIBILITY_FLAG__
-  const uint prim = optixGetPrimitiveIndex();
-  const uint visibility = optixGetPayload_4();
-  if ((kernel_tex_fetch(__prim_visibility, prim) & visibility) == 0) {
-    return optixIgnoreIntersection();
-  }
-
-  /* Shadow ray early termination. */
+  // Shadow ray early termination
   if (visibility & PATH_RAY_SHADOW_OPAQUE) {
     return optixTerminateRay();
   }
-#endif
 }
 
 extern "C" __global__ void __closesthit__kernel_optix_hit()
 {
-  optixSetPayload_0(__float_as_uint(optixGetRayTmax())); /* Intersection distance */
+  optixSetPayload_0(__float_as_uint(optixGetRayTmax()));  // Intersection distance
   optixSetPayload_3(optixGetPrimitiveIndex());
   optixSetPayload_4(get_object_id());
-  /* Can be PRIMITIVE_TRIANGLE and PRIMITIVE_MOTION_TRIANGLE or curve type and segment index. */
+  // Can be PRIMITIVE_TRIANGLE and PRIMITIVE_MOTION_TRIANGLE or curve type and segment index
   optixSetPayload_5(kernel_tex_fetch(__prim_type, optixGetPrimitiveIndex()));
 
   if (optixIsTriangleHit()) {
@@ -328,7 +297,7 @@ extern "C" __global__ void __closesthit__kernel_optix_hit()
     optixSetPayload_2(__float_as_uint(barycentrics.x));
   }
   else {
-    optixSetPayload_1(optixGetAttribute_0()); /* Same as 'optixGetCurveParameter()' */
+    optixSetPayload_1(optixGetAttribute_0());  // Same as 'optixGetCurveParameter()'
     optixSetPayload_2(optixGetAttribute_1());
   }
 }
@@ -342,7 +311,7 @@ ccl_device_inline void optix_intersection_curve(const uint prim, const uint type
   float3 P = optixGetObjectRayOrigin();
   float3 dir = optixGetObjectRayDirection();
 
-  /* The direction is not normalized by default, but the curve intersection routine expects that */
+  // The direction is not normalized by default, but the curve intersection routine expects that
   float len;
   dir = normalize_len(dir, &len);
 
@@ -354,15 +323,15 @@ ccl_device_inline void optix_intersection_curve(const uint prim, const uint type
 
   Intersection isect;
   isect.t = optixGetRayTmax();
-  /* Transform maximum distance into object space. */
+  // Transform maximum distance into object space
   if (isect.t != FLT_MAX)
     isect.t *= len;
 
   if (curve_intersect(NULL, &isect, P, dir, isect.t, visibility, object, prim, time, type)) {
     optixReportIntersection(isect.t / len,
                             type & PRIMITIVE_ALL,
-                            __float_as_int(isect.u),  /* Attribute_0 */
-                            __float_as_int(isect.v)); /* Attribute_1 */
+                            __float_as_int(isect.u),   // Attribute_0
+                            __float_as_int(isect.v));  // Attribute_1
   }
 }
 

intern/cycles/kernel/geom/geom_curve_intersect.h

@@ -713,7 +713,7 @@ ccl_device_inline void curve_shader_setup(const KernelGlobals *kg,
 
     P = transform_point(&tfm, P);
     D = transform_direction(&tfm, D * t);
-    D = safe_normalize_len(D, &t);
+    D = normalize_len(D, &t);
   }
 
   int prim = kernel_tex_fetch(__prim_index, isect_prim);
@@ -764,10 +764,8 @@ ccl_device_inline void curve_shader_setup(const KernelGlobals *kg,
     /* Thick curves, compute normal using direction from inside the curve.
      * This could be optimized by recording the normal in the intersection,
      * however for Optix this would go beyond the size of the payload. */
-    /* NOTE: It is possible that P will be the same as P_inside (precision issues, or very small
-     * radius). In this case use the view direction to approximate the normal. */
     const float3 P_inside = float4_to_float3(catmull_rom_basis_eval(P_curve, sd->u));
-    const float3 Ng = (!isequal_float3(P, P_inside)) ? normalize(P - P_inside) : -sd->I;
+    const float3 Ng = normalize(P - P_inside);
 
     sd->N = Ng;
     sd->Ng = Ng;

intern/cycles/kernel/geom/geom_motion_triangle.h

@@ -41,18 +41,7 @@ ccl_device_inline int find_attribute_motion(const KernelGlobals *kg,
   uint4 attr_map = kernel_tex_fetch(__attributes_map, attr_offset);
 
   while (attr_map.x != id) {
-    if (UNLIKELY(attr_map.x == ATTR_STD_NONE)) {
-      if (UNLIKELY(attr_map.y == 0)) {
-        return (int)ATTR_STD_NOT_FOUND;
-      }
-      else {
-        /* Chain jump to a different part of the table. */
-        attr_offset = attr_map.z;
-      }
-    }
-    else {
-      attr_offset += ATTR_PRIM_TYPES;
-    }
+    attr_offset += ATTR_PRIM_TYPES;
     attr_map = kernel_tex_fetch(__attributes_map, attr_offset);
   }
 

intern/cycles/kernel/integrator/integrator_state.h

@@ -60,15 +60,7 @@ CCL_NAMESPACE_BEGIN
  * TODO: these could be made dynamic depending on the features used in the scene. */
 
 #define INTEGRATOR_VOLUME_STACK_SIZE VOLUME_STACK_SIZE
-
-#define INTEGRATOR_SHADOW_ISECT_SIZE_CPU 1024
-#define INTEGRATOR_SHADOW_ISECT_SIZE_GPU 4
-
-#ifdef __KERNEL_CPU__
-#  define INTEGRATOR_SHADOW_ISECT_SIZE INTEGRATOR_SHADOW_ISECT_SIZE_CPU
-#else
-#  define INTEGRATOR_SHADOW_ISECT_SIZE INTEGRATOR_SHADOW_ISECT_SIZE_GPU
-#endif
+#define INTEGRATOR_SHADOW_ISECT_SIZE 4
 
 /* Data structures */
 
@@ -82,9 +74,9 @@ typedef struct IntegratorStateCPU {
 #define KERNEL_STRUCT_END(name) \
   } \
   name;
-#define KERNEL_STRUCT_END_ARRAY(name, cpu_size, gpu_size) \
+#define KERNEL_STRUCT_END_ARRAY(name, size) \
   } \
-  name[cpu_size];
+  name[size];
 #include "kernel/integrator/integrator_state_template.h"
 #undef KERNEL_STRUCT_BEGIN
 #undef KERNEL_STRUCT_MEMBER
@@ -111,9 +103,9 @@ typedef struct IntegratorStateGPU {
 #define KERNEL_STRUCT_END(name) \
   } \
   name;
-#define KERNEL_STRUCT_END_ARRAY(name, cpu_size, gpu_size) \
+#define KERNEL_STRUCT_END_ARRAY(name, size) \
   } \
-  name[gpu_size];
+  name[size];
 #include "kernel/integrator/integrator_state_template.h"
 #undef KERNEL_STRUCT_BEGIN
 #undef KERNEL_STRUCT_MEMBER

intern/cycles/kernel/integrator/integrator_state_template.h

@@ -107,7 +107,7 @@ KERNEL_STRUCT_END(subsurface)
 KERNEL_STRUCT_BEGIN(volume_stack)
 KERNEL_STRUCT_ARRAY_MEMBER(volume_stack, int, object, KERNEL_FEATURE_VOLUME)
 KERNEL_STRUCT_ARRAY_MEMBER(volume_stack, int, shader, KERNEL_FEATURE_VOLUME)
-KERNEL_STRUCT_END_ARRAY(volume_stack, INTEGRATOR_VOLUME_STACK_SIZE, INTEGRATOR_VOLUME_STACK_SIZE)
+KERNEL_STRUCT_END_ARRAY(volume_stack, INTEGRATOR_VOLUME_STACK_SIZE)
 
 /********************************* Shadow Path State **************************/
 
@@ -153,15 +153,11 @@ KERNEL_STRUCT_ARRAY_MEMBER(shadow_isect, int, object, KERNEL_FEATURE_PATH_TRACIN
 KERNEL_STRUCT_ARRAY_MEMBER(shadow_isect, int, type, KERNEL_FEATURE_PATH_TRACING)
 /* TODO: exclude for GPU. */
 KERNEL_STRUCT_ARRAY_MEMBER(shadow_isect, float3, Ng, KERNEL_FEATURE_PATH_TRACING)
-KERNEL_STRUCT_END_ARRAY(shadow_isect,
-                        INTEGRATOR_SHADOW_ISECT_SIZE_CPU,
-                        INTEGRATOR_SHADOW_ISECT_SIZE_GPU)
+KERNEL_STRUCT_END_ARRAY(shadow_isect, INTEGRATOR_SHADOW_ISECT_SIZE)
 
 /**************************** Shadow Volume Stack *****************************/
 
 KERNEL_STRUCT_BEGIN(shadow_volume_stack)
 KERNEL_STRUCT_ARRAY_MEMBER(shadow_volume_stack, int, object, KERNEL_FEATURE_VOLUME)
 KERNEL_STRUCT_ARRAY_MEMBER(shadow_volume_stack, int, shader, KERNEL_FEATURE_VOLUME)
-KERNEL_STRUCT_END_ARRAY(shadow_volume_stack,
-                        INTEGRATOR_VOLUME_STACK_SIZE,
-                        INTEGRATOR_VOLUME_STACK_SIZE)
+KERNEL_STRUCT_END_ARRAY(shadow_volume_stack, INTEGRATOR_VOLUME_STACK_SIZE)

intern/cycles/kernel/integrator/integrator_state_util.h

@@ -217,10 +217,10 @@ ccl_device_inline void integrator_state_copy_only(const IntegratorState to_state
     while (false) \
       ;
 
-#  define KERNEL_STRUCT_END_ARRAY(name, cpu_array_size, gpu_array_size) \
+#  define KERNEL_STRUCT_END_ARRAY(name, array_size) \
     ++index; \
     } \
-    while (index < gpu_array_size) \
+    while (index < array_size) \
       ;
 
 #  include "kernel/integrator/integrator_state_template.h"
@@ -264,12 +264,7 @@ ccl_device_inline void integrator_state_shadow_catcher_split(INTEGRATOR_STATE_AR
 
   IntegratorStateCPU *ccl_restrict split_state = state + 1;
 
-  /* Only copy the required subset, since shadow intersections are big and irrelevant here. */
-  split_state->path = state->path;
-  split_state->ray = state->ray;
-  split_state->isect = state->isect;
-  memcpy(split_state->volume_stack, state->volume_stack, sizeof(state->volume_stack));
-  split_state->shadow_path = state->shadow_path;
+  *split_state = *state;
 
   split_state->path.flag |= PATH_RAY_SHADOW_CATCHER_PASS;
 #endif

intern/cycles/kernel/kernel_accumulate.h

@@ -386,7 +386,7 @@ ccl_device_inline void kernel_accum_light(INTEGRATOR_STATE_CONST_ARGS,
 {
   /* The throughput for shadow paths already contains the light shader evaluation. */
   float3 contribution = INTEGRATOR_STATE(shadow_path, throughput);
-  kernel_accum_clamp(kg, &contribution, INTEGRATOR_STATE(shadow_path, bounce));
+  kernel_accum_clamp(kg, &contribution, INTEGRATOR_STATE(shadow_path, bounce) - 1);
 
   ccl_global float *buffer = kernel_accum_pixel_render_buffer(INTEGRATOR_STATE_PASS,
                                                               render_buffer);

intern/cycles/kernel/kernel_bake.h

@@ -42,16 +42,6 @@ ccl_device void kernel_displace_evaluate(const KernelGlobals *kg,
 
   object_inverse_dir_transform(kg, &sd, &D);
 
-#ifdef __KERNEL_DEBUG_NAN__
-  if (!isfinite3_safe(D)) {
-    kernel_assert(!"Cycles displacement with non-finite value detected");
-  }
-#endif
-
-  /* Ensure finite displacement, preventing BVH from becoming degenerate and avoiding possible
-   * traversal issues caused by non-finite math. */
-  D = ensure_finite3(D);
-
   /* Write output. */
   output[offset] += make_float4(D.x, D.y, D.z, 0.0f);
 }
@@ -76,16 +66,7 @@ ccl_device void kernel_background_evaluate(const KernelGlobals *kg,
   const int path_flag = PATH_RAY_EMISSION;
   shader_eval_surface<KERNEL_FEATURE_NODE_MASK_SURFACE_LIGHT>(
       INTEGRATOR_STATE_PASS_NULL, &sd, NULL, path_flag);
-  float3 color = shader_background_eval(&sd);
-
-#ifdef __KERNEL_DEBUG_NAN__
-  if (!isfinite3_safe(color)) {
-    kernel_assert(!"Cycles background with non-finite value detected");
-  }
-#endif
-
-  /* Ensure finite color, avoiding possible numerical instabilities in the path tracing kernels. */
-  color = ensure_finite3(color);
+  const float3 color = shader_background_eval(&sd);
 
   /* Write output. */
   output[offset] += make_float4(color.x, color.y, color.z, 0.0f);

intern/cycles/kernel/kernel_jitter.h

@@ -74,6 +74,10 @@ ccl_device_inline float cmj_randfloat_simple(uint i, uint p)
 
 ccl_device float pmj_sample_1D(const KernelGlobals *kg, uint sample, uint rng_hash, uint dimension)
 {
+  /* The PMJ sample sets contain a sample with (x,y) with NUM_PMJ_SAMPLES so for 1D
+   *  the x part is used as the sample (TODO(@leesonw): Add using both x and y parts
+   * independently). */
+
   /* Perform Owen shuffle of the sample number to reorder the samples. */
 #ifdef _SIMPLE_HASH_
   const uint rv = cmj_hash_simple(dimension, rng_hash);
@@ -91,10 +95,7 @@ ccl_device float pmj_sample_1D(const KernelGlobals *kg, uint sample, uint rng_ha
   const uint sample_set = s / NUM_PMJ_SAMPLES;
   const uint d = (dimension + sample_set);
   const uint dim = d % NUM_PMJ_PATTERNS;
-
-  /* The PMJ sample sets contain a sample with (x,y) with NUM_PMJ_SAMPLES so for 1D
-   *  the x part is used for even dims and the y for odd. */
-  int index = 2 * ((dim >> 1) * NUM_PMJ_SAMPLES + (s % NUM_PMJ_SAMPLES)) + (dim & 1);
+  int index = 2 * (dim * NUM_PMJ_SAMPLES + (s % NUM_PMJ_SAMPLES));
 
   float fx = kernel_tex_fetch(__sample_pattern_lut, index);
 
@@ -103,11 +104,12 @@ ccl_device float pmj_sample_1D(const KernelGlobals *kg, uint sample, uint rng_ha
 #  ifdef _SIMPLE_HASH_
   float dx = cmj_randfloat_simple(d, rng_hash);
 #  else
+  /* Only jitter within the grid interval. */
   float dx = cmj_randfloat(d, rng_hash);
 #  endif
-  /* Jitter sample locations and map back into [0 1]. */
-  fx = fx + dx;
+  fx = fx + dx * (1.0f / NUM_PMJ_SAMPLES);
   fx = fx - floorf(fx);
+
 #else
 #  warning "Not using Cranley-Patterson Rotation."
 #endif
@@ -134,7 +136,7 @@ ccl_device void pmj_sample_2D(
   /* Based on the sample number a sample pattern is selected and offset by the dimension. */
   const uint sample_set = s / NUM_PMJ_SAMPLES;
   const uint d = (dimension + sample_set);
-  uint dim = d % NUM_PMJ_PATTERNS;
+  const uint dim = d % NUM_PMJ_PATTERNS;
   int index = 2 * (dim * NUM_PMJ_SAMPLES + (s % NUM_PMJ_SAMPLES));
 
   float fx = kernel_tex_fetch(__sample_pattern_lut, index);
@@ -149,17 +151,17 @@ ccl_device void pmj_sample_2D(
   float dx = cmj_randfloat(d, rng_hash);
   float dy = cmj_randfloat(d + 1, rng_hash);
 #  endif
-  /* Jitter sample locations and map back to the unit square [0 1]x[0 1]. */
-  float sx = fx + dx;
-  float sy = fy + dy;
-  sx = sx - floorf(sx);
-  sy = sy - floorf(sy);
+  /* Only jitter within the grid cells. */
+  fx = fx + dx * (1.0f / NUM_PMJ_DIVISIONS);
+  fy = fy + dy * (1.0f / NUM_PMJ_DIVISIONS);
+  fx = fx - floorf(fx);
+  fy = fy - floorf(fy);
 #else
 #  warning "Not using Cranley Patterson Rotation."
 #endif
 
-  (*x) = sx;
-  (*y) = sy;
+  (*x) = fx;
+  (*y) = fy;
 }
 
 CCL_NAMESPACE_END

intern/cycles/render/film.cpp

@@ -434,8 +434,7 @@ void Film::update_passes(Scene *scene, bool add_sample_count_pass)
   const ObjectManager *object_manager = scene->object_manager;
   Integrator *integrator = scene->integrator;
 
-  if (!is_modified() && !object_manager->need_update() && !integrator->is_modified() &&
-      !background->is_modified()) {
+  if (!is_modified() && !object_manager->need_update() && !integrator->is_modified()) {
     return;
   }
 

intern/cycles/render/tile.cpp

@@ -436,12 +436,7 @@ bool TileManager::open_tile_output()
     return false;
   }
 
-  if (!write_state_.tile_out->open(write_state_.filename, write_state_.image_spec)) {
-    LOG(ERROR) << "Error opening tile file: " << write_state_.tile_out->geterror();
-    write_state_.tile_out = nullptr;
-    return false;
-  }
-
+  write_state_.tile_out->open(write_state_.filename, write_state_.image_spec);
   write_state_.num_tiles_written = 0;
 
   VLOG(3) << "Opened tile file " << write_state_.filename;
@@ -502,7 +497,6 @@ bool TileManager::write_tile(const RenderBuffers &tile_buffers)
                                           TypeDesc::FLOAT,
                                           pixels)) {
     LOG(ERROR) << "Error writing tile " << write_state_.tile_out->geterror();
-    return false;
   }
 
   ++write_state_.num_tiles_written;

intern/cycles/util/util_atomic.h

@@ -34,7 +34,7 @@
 
 #else /* __KERNEL_GPU__ */
 
-#  if defined(__KERNEL_CUDA__) || defined(__KERNEL_HIP__)
+#  ifdef __KERNEL_CUDA__
 
 #    define atomic_add_and_fetch_float(p, x) (atomicAdd((float *)(p), (float)(x)) + (float)(x))
 

intern/cycles/util/util_debug.cpp

@@ -59,23 +59,12 @@ DebugFlags::CUDA::CUDA() : adaptive_compile(false)
   reset();
 }
 
-DebugFlags::HIP::HIP() : adaptive_compile(false)
-{
-  reset();
-}
-
 void DebugFlags::CUDA::reset()
 {
   if (getenv("CYCLES_CUDA_ADAPTIVE_COMPILE") != NULL)
     adaptive_compile = true;
 }
 
-void DebugFlags::HIP::reset()
-{
-  if (getenv("CYCLES_HIP_ADAPTIVE_COMPILE") != NULL)
-    adaptive_compile = true;
-}
-
 DebugFlags::OptiX::OptiX()
 {
   reset();
@@ -114,10 +103,6 @@ std::ostream &operator<<(std::ostream &os, DebugFlagsConstRef debug_flags)
 
   os << "OptiX flags:\n"
      << "  Debug : " << string_from_bool(debug_flags.optix.use_debug) << "\n";
-
-  os << "HIP flags:\n"
-     << "  HIP streams : " << string_from_bool(debug_flags.hip.adaptive_compile) << "\n";
-
   return os;
 }
 

intern/cycles/util/util_debug.h

@@ -93,17 +93,6 @@ class DebugFlags {
     bool adaptive_compile;
   };
 
-  /* Descriptor of HIP feature-set to be used. */
-  struct HIP {
-    HIP();
-
-    /* Reset flags to their defaults. */
-    void reset();
-
-    /* Whether adaptive feature based runtime compile is enabled or not.*/
-    bool adaptive_compile;
-  };
-
   /* Descriptor of OptiX feature-set to be used. */
   struct OptiX {
     OptiX();
@@ -135,9 +124,6 @@ class DebugFlags {
   /* Requested OptiX flags. */
   OptiX optix;
 
-  /* Requested HIP flags. */
-  HIP hip;
-
  private:
   DebugFlags();
 

intern/cycles/util/util_half.h

@@ -29,7 +29,7 @@ CCL_NAMESPACE_BEGIN
 /* Half Floats */
 
 /* CUDA has its own half data type, no need to define then */
-#if !defined(__KERNEL_CUDA__) && !defined(__KERNEL_HIP__)
+#ifndef __KERNEL_CUDA__
 /* Implementing this as a class rather than a typedef so that the compiler can tell it apart from
  * unsigned shorts. */
 class half {
@@ -59,7 +59,7 @@ struct half4 {
   half x, y, z, w;
 };
 
-#if defined(__KERNEL_CUDA__) || defined(__KERNEL_HIP__)
+#ifdef __KERNEL_CUDA__
 
 ccl_device_inline void float4_store_half(half *h, float4 f)
 {
@@ -73,7 +73,6 @@ ccl_device_inline void float4_store_half(half *h, float4 f)
 
 ccl_device_inline void float4_store_half(half *h, float4 f)
 {
-
 #  ifndef __KERNEL_SSE2__
   for (int i = 0; i < 4; i++) {
     /* optimized float to half for pixels:
@@ -110,8 +109,6 @@ ccl_device_inline void float4_store_half(half *h, float4 f)
 #  endif
 }
 
-#  ifndef __KERNEL_HIP__
-
 ccl_device_inline float half_to_float(half h)
 {
   float f;
@@ -120,23 +117,6 @@ ccl_device_inline float half_to_float(half h)
 
   return f;
 }
-#  else
-
-ccl_device_inline float half_to_float(std::uint32_t a) noexcept
-{
-
-  std::uint32_t u = ((a << 13) + 0x70000000U) & 0x8fffe000U;
-
-  std::uint32_t v = __float_as_uint(__uint_as_float(u) *
-                                    __uint_as_float(0x77800000U) /*0x1.0p+112f*/) +
-                    0x38000000U;
-
-  u = (a & 0x7fff) != 0 ? v : u;
-
-  return __uint_as_float(u) * __uint_as_float(0x07800000U) /*0x1.0p-112f*/;
-}
-
-#  endif /* __KERNEL_HIP__ */
 
 ccl_device_inline float4 half4_to_float4(half4 h)
 {

intern/cycles/util/util_math.h

@@ -26,10 +26,6 @@
 #  include <cmath>
 #endif
 
-#ifdef __HIP__
-#  include <hip/hip_vector_types.h>
-#endif
-
 #include <float.h>
 #include <math.h>
 #include <stdio.h>
@@ -87,8 +83,7 @@ CCL_NAMESPACE_BEGIN
 
 /* Scalar */
 
-#ifndef __HIP__
-#  ifdef _WIN32
+#ifdef _WIN32
 ccl_device_inline float fmaxf(float a, float b)
 {
   return (a > b) ? a : b;
@@ -98,9 +93,7 @@ ccl_device_inline float fminf(float a, float b)
 {
   return (a < b) ? a : b;
 }
-
-#  endif /* _WIN32 */
-#endif   /* __HIP__ */
+#endif /* _WIN32 */
 
 #ifndef __KERNEL_GPU__
 using std::isfinite;
@@ -206,7 +199,6 @@ ccl_device_inline uint as_uint(float f)
   return u.i;
 }
 
-#ifndef __HIP__
 ccl_device_inline int __float_as_int(float f)
 {
   union {
@@ -246,7 +238,6 @@ ccl_device_inline float __uint_as_float(uint i)
   u.i = i;
   return u.f;
 }
-#endif
 
 ccl_device_inline int4 __float4_as_int4(float4 f)
 {
@@ -678,7 +669,7 @@ ccl_device float bits_to_01(uint bits)
 
 ccl_device_inline uint count_leading_zeros(uint x)
 {
-#if defined(__KERNEL_CUDA__) || defined(__KERNEL_OPTIX__) || defined(__KERNEL_HIP__)
+#if defined(__KERNEL_CUDA__) || defined(__KERNEL_OPTIX__)
   return __clz(x);
 #else
   assert(x != 0);
@@ -694,7 +685,7 @@ ccl_device_inline uint count_leading_zeros(uint x)
 
 ccl_device_inline uint count_trailing_zeros(uint x)
 {
-#if defined(__KERNEL_CUDA__) || defined(__KERNEL_OPTIX__) || defined(__KERNEL_HIP__)
+#if defined(__KERNEL_CUDA__) || defined(__KERNEL_OPTIX__)
   return (__ffs(x) - 1);
 #else
   assert(x != 0);
@@ -710,7 +701,7 @@ ccl_device_inline uint count_trailing_zeros(uint x)
 
 ccl_device_inline uint find_first_set(uint x)
 {
-#if defined(__KERNEL_CUDA__) || defined(__KERNEL_OPTIX__) || defined(__KERNEL_HIP__)
+#if defined(__KERNEL_CUDA__) || defined(__KERNEL_OPTIX__)
   return __ffs(x);
 #else
 #  ifdef _MSC_VER

intern/ghost/GHOST_IWindow.h

@@ -40,7 +40,7 @@
  * There are two coordinate systems:
  *
  * - The screen coordinate system. The origin of the screen is located in the
- *   upper left corner of the screen.
+ *   upper left corner of the screen.</li>
  * - The client rectangle coordinate system. The client rectangle of a window
  *   is the area that is drawable by the application (excluding title bars etc.).
  */

release/datafiles/userdef/userdef_default_theme.c

@@ -817,7 +817,7 @@ const bTheme U_theme_default = {
     },
     .shade2 = RGBA(0x7f707064),
     .grid = RGBA(0x23232300),
-    .wire = RGBA(0x232323ff),
+    .wire = RGBA(0x808080ff),
     .select = RGBA(0xed5700ff),
     .active = RGBA(0xffffffff),
     .edge_select = RGBA(0xffffffff),
@@ -1175,35 +1175,6 @@ const bTheme U_theme_default = {
       .color = RGBA(0x7a5441ff),
     },
   },
-  .strip_color = {
-    {
-      .color = RGBA(0xe2605bff),
-    },
-    {
-      .color = RGBA(0xf1a355ff),
-    },
-    {
-      .color = RGBA(0xf1dc55ff),
-    },
-    {
-      .color = RGBA(0x7bcc7bff),
-    },
-    {
-      .color = RGBA(0x5db6eaff),
-    },
-    {
-      .color = RGBA(0x8d59daff),
-    },
-    {
-      .color = RGBA(0xc673b8ff),
-    },
-    {
-      .color = RGBA(0x7a5441ff),
-    },
-    {
-      .color = RGBA(0x5f5f5fff),
-    },
-  },
 };
 
 /* clang-format on */

release/scripts/modules/bpy_extras/asset_utils.py

@@ -52,12 +52,19 @@ class AssetBrowserPanel:
     bl_space_type = 'FILE_BROWSER'
 
     @classmethod
-    def asset_browser_panel_poll(cls, context):
+    def poll(cls, context):
         return SpaceAssetInfo.is_asset_browser_poll(context)
 
+
+class AssetBrowserSpecificCategoryPanel(AssetBrowserPanel):
+    asset_categories = set()  # Set of strings like 'ANIMATIONS', see `asset_category_items` in rna_space.c
+
     @classmethod
     def poll(cls, context):
-        return cls.asset_browser_panel_poll(context)
+        return (
+            SpaceAssetInfo.is_asset_browser_poll(context)
+            and context.space_data.params.asset_category in cls.asset_categories
+        )
 
 
 class AssetMetaDataPanel:

release/scripts/modules/rna_manual_reference.py

@@ -39,8 +39,8 @@ if LANG is not None:
     url_manual_prefix = url_manual_prefix.replace("manual/en", "manual/" + LANG)
 
 url_manual_mapping = (
+	("bpy.types.cyclesworldsettings.sample_mbpy.types.cyclesworldsettings.sample_map_resolutionbpy.types.cyclesworldsettings.sample_map_resolutionap_resolution*", "render/cycles/world_settings.html#bpy-types-cyclesworldsettings-sample-mbpy-types-cyclesworldsettings-sample-map-resolutionbpy-types-cyclesworldsettings-sample-map-resolutionap-resolution"),
 	("bpy.types.movietrackingsettings.refine_intrinsics_tangential_distortion*", "movie_clip/tracking/clip/toolbar/solve.html#bpy-types-movietrackingsettings-refine-intrinsics-tangential-distortion"),
-	("bpy.types.spacesequesequencertimelineoverlaynceeditor.show_strip_offset*", "editors/video_sequencer/sequencer/display.html#bpy-types-spacesequesequencertimelineoverlaynceeditor-show-strip-offset"),
 	("bpy.types.movietrackingsettings.refine_intrinsics_radial_distortion*", "movie_clip/tracking/clip/toolbar/solve.html#bpy-types-movietrackingsettings-refine-intrinsics-radial-distortion"),
 	("bpy.types.fluiddomainsettings.sndparticle_potential_max_trappedair*", "physics/fluid/type/domain/liquid/particles.html#bpy-types-fluiddomainsettings-sndparticle-potential-max-trappedair"),
 	("bpy.types.fluiddomainsettings.sndparticle_potential_min_trappedair*", "physics/fluid/type/domain/liquid/particles.html#bpy-types-fluiddomainsettings-sndparticle-potential-min-trappedair"),
@@ -60,13 +60,11 @@ url_manual_mapping = (
 	("bpy.types.fluiddomainsettings.sndparticle_sampling_wavecrest*", "physics/fluid/type/domain/liquid/particles.html#bpy-types-fluiddomainsettings-sndparticle-sampling-wavecrest"),
 	("bpy.types.rigidbodyconstraint.use_override_solver_iterations*", "physics/rigid_body/constraints/introduction.html#bpy-types-rigidbodyconstraint-use-override-solver-iterations"),
 	("bpy.types.toolsettings.use_transform_correct_face_attributes*", "modeling/meshes/tools/tool_settings.html#bpy-types-toolsettings-use-transform-correct-face-attributes"),
-	("bpy.types.cyclesrendersettings.preview_adaptive_min_samples*", "render/cycles/render_settings/sampling.html#bpy-types-cyclesrendersettings-preview-adaptive-min-samples"),
 	("bpy.types.rendersettings.use_sequencer_override_scene_strip*", "video_editing/preview/sidebar.html#bpy-types-rendersettings-use-sequencer-override-scene-strip"),
 	("bpy.types.toolsettings.use_transform_correct_keep_connected*", "modeling/meshes/tools/tool_settings.html#bpy-types-toolsettings-use-transform-correct-keep-connected"),
-	("bpy.types.cyclesrendersettings.preview_denoising_prefilter*", "render/cycles/render_settings/sampling.html#bpy-types-cyclesrendersettings-preview-denoising-prefilter"),
 	("bpy.types.fluiddomainsettings.sndparticle_potential_radius*", "physics/fluid/type/domain/liquid/particles.html#bpy-types-fluiddomainsettings-sndparticle-potential-radius"),
 	("bpy.types.cyclesrendersettings.film_transparent_roughness*", "render/cycles/render_settings/film.html#bpy-types-cyclesrendersettings-film-transparent-roughness"),
-	("bpy.types.cyclesrendersettings.preview_adaptive_threshold*", "render/cycles/render_settings/sampling.html#bpy-types-cyclesrendersettings-preview-adaptive-threshold"),
+	("bpy.types.cyclesrendersettings.sample_all_lights_indirect*", "render/cycles/render_settings/sampling.html#bpy-types-cyclesrendersettings-sample-all-lights-indirect"),
 	("bpy.types.fluiddomainsettings.openvdb_cache_compress_type*", "physics/fluid/type/domain/cache.html#bpy-types-fluiddomainsettings-openvdb-cache-compress-type"),
 	("bpy.types.fluiddomainsettings.sndparticle_bubble_buoyancy*", "physics/fluid/type/domain/liquid/particles.html#bpy-types-fluiddomainsettings-sndparticle-bubble-buoyancy"),
 	("bpy.types.fluiddomainsettings.sndparticle_combined_export*", "physics/fluid/type/domain/liquid/particles.html#bpy-types-fluiddomainsettings-sndparticle-combined-export"),
@@ -78,13 +76,14 @@ url_manual_mapping = (
 	("bpy.types.toolsettings.annotation_stroke_placement_view3d*", "interface/annotate_tool.html#bpy-types-toolsettings-annotation-stroke-placement-view3d"),
 	("bpy.types.fluiddomainsettings.use_collision_border_front*", "physics/fluid/type/domain/settings.html#bpy-types-fluiddomainsettings-use-collision-border-front"),
 	("bpy.types.fluiddomainsettings.use_collision_border_right*", "physics/fluid/type/domain/settings.html#bpy-types-fluiddomainsettings-use-collision-border-right"),
-	("bpy.types.sequencertimelineoverlay.waveform_display_type*", "editors/video_sequencer/sequencer/display.html#bpy-types-sequencertimelineoverlay-waveform-display-type"),
 	("bpy.types.cyclesmaterialsettings.use_transparent_shadow*", "render/cycles/material_settings.html#bpy-types-cyclesmaterialsettings-use-transparent-shadow"),
 	("bpy.types.cyclesobjectsettings.shadow_terminator_offset*", "render/cycles/object_settings/object_data.html#bpy-types-cyclesobjectsettings-shadow-terminator-offset"),
 	("bpy.types.cyclesobjectsettings.use_adaptive_subdivision*", "render/cycles/object_settings/adaptive_subdiv.html#bpy-types-cyclesobjectsettings-use-adaptive-subdivision"),
 	("bpy.types.cyclesrendersettings.debug_use_spatial_splits*", "render/cycles/render_settings/performance.html#bpy-types-cyclesrendersettings-debug-use-spatial-splits"),
 	("bpy.types.cyclesrendersettings.light_sampling_threshold*", "render/cycles/render_settings/sampling.html#bpy-types-cyclesrendersettings-light-sampling-threshold"),
+	("bpy.types.cyclesrendersettings.preview_start_resolution*", "render/cycles/render_settings/performance.html#bpy-types-cyclesrendersettings-preview-start-resolution"),
 	("bpy.types.cyclesrendersettings.rolling_shutter_duration*", "render/cycles/render_settings/motion_blur.html#bpy-types-cyclesrendersettings-rolling-shutter-duration"),
+	("bpy.types.cyclesrendersettings.sample_all_lights_direct*", "render/cycles/render_settings/sampling.html#bpy-types-cyclesrendersettings-sample-all-lights-direct"),
 	("bpy.types.cyclesrendersettings.volume_preview_step_rate*", "render/cycles/render_settings/volumes.html#bpy-types-cyclesrendersettings-volume-preview-step-rate"),
 	("bpy.types.fluiddomainsettings.sndparticle_update_radius*", "physics/fluid/type/domain/liquid/particles.html#bpy-types-fluiddomainsettings-sndparticle-update-radius"),
 	("bpy.types.fluiddomainsettings.use_collision_border_back*", "physics/fluid/type/domain/settings.html#bpy-types-fluiddomainsettings-use-collision-border-back"),
@@ -98,15 +97,14 @@ url_manual_mapping = (
 	("bpy.types.gpencilsculptsettings.use_multiframe_falloff*", "grease_pencil/multiframe.html#bpy-types-gpencilsculptsettings-use-multiframe-falloff"),
 	("bpy.types.movietrackingsettings.use_keyframe_selection*", "movie_clip/tracking/clip/toolbar/solve.html#bpy-types-movietrackingsettings-use-keyframe-selection"),
 	("bpy.types.rendersettings.simplify_gpencil_antialiasing*", "render/cycles/render_settings/simplify.html#bpy-types-rendersettings-simplify-gpencil-antialiasing"),
-	("bpy.types.sequencertimelineoverlay.show_strip_duration*", "editors/video_sequencer/sequencer/display.html#bpy-types-sequencertimelineoverlay-show-strip-duration"),
 	("bpy.types.spaceoutliner.use_filter_lib_override_system*", "editors/outliner/interface.html#bpy-types-spaceoutliner-use-filter-lib-override-system"),
 	("bpy.types.toolsettings.use_transform_pivot_point_align*", "scene_layout/object/tools/tool_settings.html#bpy-types-toolsettings-use-transform-pivot-point-align"),
 	("bpy.types.brush.show_multiplane_scrape_planes_preview*", "sculpt_paint/sculpting/tools/multiplane_scrape.html#bpy-types-brush-show-multiplane-scrape-planes-preview"),
 	("bpy.types.cyclesmaterialsettings.volume_interpolation*", "render/cycles/material_settings.html#bpy-types-cyclesmaterialsettings-volume-interpolation"),
 	("bpy.types.cyclesrendersettings.debug_optix_curves_api*", "render/cycles/render_settings/debug.html#bpy-types-cyclesrendersettings-debug-optix-curves-api"),
-	("bpy.types.cyclesrendersettings.denoising_input_passes*", "render/cycles/render_settings/sampling.html#bpy-types-cyclesrendersettings-denoising-input-passes"),
 	("bpy.types.cyclesrendersettings.film_transparent_glass*", "render/cycles/render_settings/film.html#bpy-types-cyclesrendersettings-film-transparent-glass"),
 	("bpy.types.cyclesrendersettings.offscreen_dicing_scale*", "render/cycles/render_settings/subdivision.html#bpy-types-cyclesrendersettings-offscreen-dicing-scale"),
+	("bpy.types.cyclesrendersettings.use_progressive_refine*", "render/cycles/render_settings/performance.html#bpy-types-cyclesrendersettings-use-progressive-refine"),
 	("bpy.types.fluiddomainsettings.sndparticle_bubble_drag*", "physics/fluid/type/domain/liquid/particles.html#bpy-types-fluiddomainsettings-sndparticle-bubble-drag"),
 	("bpy.types.linestylegeometrymodifier_backbonestretcher*", "render/freestyle/view_layer/line_style/modifiers/geometry/backbone_stretcher.html#bpy-types-linestylegeometrymodifier-backbonestretcher"),
 	("bpy.types.linestylegeometrymodifier_sinusdisplacement*", "render/freestyle/view_layer/line_style/modifiers/geometry/sinus_displacement.html#bpy-types-linestylegeometrymodifier-sinusdisplacement"),
@@ -125,7 +123,6 @@ url_manual_mapping = (
 	("bpy.types.gpencillayer.use_annotation_onion_skinning*", "interface/annotate_tool.html#bpy-types-gpencillayer-use-annotation-onion-skinning"),
 	("bpy.types.greasepencil.use_adaptive_curve_resolution*", "grease_pencil/modes/edit/curve_editing.html#bpy-types-greasepencil-use-adaptive-curve-resolution"),
 	("bpy.types.linestylegeometrymodifier_polygonalization*", "render/freestyle/view_layer/line_style/modifiers/geometry/polygonization.html#bpy-types-linestylegeometrymodifier-polygonalization"),
-	("bpy.types.sequencertimelineoverlay.show_strip_source*", "editors/video_sequencer/sequencer/display.html#bpy-types-sequencertimelineoverlay-show-strip-source"),
 	("bpy.types.toolsettings.use_gpencil_automerge_strokes*", "grease_pencil/modes/draw/introduction.html#bpy-types-toolsettings-use-gpencil-automerge-strokes"),
 	("bpy.types.toolsettings.use_proportional_edit_objects*", "editors/3dview/controls/proportional_editing.html#bpy-types-toolsettings-use-proportional-edit-objects"),
 	("bpy.ops.view3d.edit_mesh_extrude_move_shrink_fatten*", "modeling/meshes/editing/face/extrude_faces_normal.html#bpy-ops-view3d-edit-mesh-extrude-move-shrink-fatten"),
@@ -136,7 +133,7 @@ url_manual_mapping = (
 	("bpy.types.cyclesrendersettings.motion_blur_position*", "render/cycles/render_settings/motion_blur.html#bpy-types-cyclesrendersettings-motion-blur-position"),
 	("bpy.types.cyclesrendersettings.rolling_shutter_type*", "render/cycles/render_settings/motion_blur.html#bpy-types-cyclesrendersettings-rolling-shutter-type"),
 	("bpy.types.cyclesrendersettings.transmission_bounces*", "render/cycles/render_settings/light_paths.html#bpy-types-cyclesrendersettings-transmission-bounces"),
-	("bpy.types.cyclesworldsettings.sample_map_resolution*", "render/cycles/world_settings.html#bpy-types-cyclesworldsettings-sample-map-resolution"),
+	("bpy.types.cyclesrendersettings.transmission_samples*", "render/cycles/render_settings/sampling.html#bpy-types-cyclesrendersettings-transmission-samples"),
 	("bpy.types.fluiddomainsettings.display_interpolation*", "physics/fluid/type/domain/gas/viewport_display.html#bpy-types-fluiddomainsettings-display-interpolation"),
 	("bpy.types.fluiddomainsettings.gridlines_cell_filter*", "physics/fluid/type/domain/gas/viewport_display.html#bpy-types-fluiddomainsettings-gridlines-cell-filter"),
 	("bpy.types.fluiddomainsettings.gridlines_color_field*", "physics/fluid/type/domain/gas/viewport_display.html#bpy-types-fluiddomainsettings-gridlines-color-field"),
@@ -150,13 +147,13 @@ url_manual_mapping = (
 	("bpy.types.rendersettings_simplify_gpencil_shader_fx*", "render/cycles/render_settings/simplify.html#bpy-types-rendersettings-simplify-gpencil-shader-fx"),
 	("bpy.types.rendersettings_simplify_gpencil_view_fill*", "render/cycles/render_settings/simplify.html#bpy-types-rendersettings-simplify-gpencil-view-fill"),
 	("bpy.types.sequencertoolsettings.snap_to_hold_offset*", "video_editing/sequencer/editing.html#bpy-types-sequencertoolsettings-snap-to-hold-offset"),
+	("bpy.types.spacesequenceeditor.waveform_display_type*", "editors/video_sequencer/sequencer/display.html#bpy-types-spacesequenceeditor-waveform-display-type"),
 	("bpy.types.toolsettings.use_mesh_automerge_and_split*", "modeling/meshes/tools/tool_settings.html#bpy-types-toolsettings-use-mesh-automerge-and-split"),
 	("bpy.types.brush.cloth_constraint_softbody_strength*", "sculpt_paint/sculpting/tools/cloth.html#bpy-types-brush-cloth-constraint-softbody-strength"),
 	("bpy.types.brush.elastic_deform_volume_preservation*", "sculpt_paint/sculpting/tools/elastic_deform.html#bpy-types-brush-elastic-deform-volume-preservation"),
 	("bpy.types.brushgpencilsettings.fill_simplify_level*", "grease_pencil/modes/draw/tools/fill.html#bpy-types-brushgpencilsettings-fill-simplify-level"),
 	("bpy.types.brushgpencilsettings.use_jitter_pressure*", "grease_pencil/modes/draw/tools/draw.html#bpy-types-brushgpencilsettings-use-jitter-pressure"),
 	("bpy.types.brushgpencilsettings.use_settings_random*", "grease_pencil/modes/draw/tools/draw.html#bpy-types-brushgpencilsettings-use-settings-random"),
-	("bpy.types.cyclesrendersettings.denoising_prefilter*", "render/cycles/render_settings/sampling.html#bpy-types-cyclesrendersettings-denoising-prefilter"),
 	("bpy.types.cyclesrendersettings.preview_dicing_rate*", "render/cycles/render_settings/subdivision.html#bpy-types-cyclesrendersettings-preview-dicing-rate"),
 	("bpy.types.cyclesrendersettings.sample_clamp_direct*", "render/cycles/render_settings/light_paths.html#bpy-types-cyclesrendersettings-sample-clamp-direct"),
 	("bpy.types.cyclesworldsettings.volume_interpolation*", "render/cycles/world_settings.html#bpy-types-cyclesworldsettings-volume-interpolation"),
@@ -169,7 +166,6 @@ url_manual_mapping = (
 	("bpy.types.freestylelineset.select_external_contour*", "render/freestyle/view_layer/line_set.html#bpy-types-freestylelineset-select-external-contour"),
 	("bpy.types.linestylegeometrymodifier_simplification*", "render/freestyle/view_layer/line_style/modifiers/geometry/simplification.html#bpy-types-linestylegeometrymodifier-simplification"),
 	("bpy.types.materialgpencilstyle.use_overlap_strokes*", "grease_pencil/materials/properties.html#bpy-types-materialgpencilstyle-use-overlap-strokes"),
-	("bpy.types.sequencertimelineoverlay.show_strip_name*", "editors/video_sequencer/sequencer/display.html#bpy-types-sequencertimelineoverlay-show-strip-name"),
 	("bpy.types.spacespreadsheet.geometry_component_type*", "editors/spreadsheet.html#bpy-types-spacespreadsheet-geometry-component-type"),
 	("bpy.types.toolsettings.use_gpencil_weight_data_add*", "grease_pencil/modes/draw/introduction.html#bpy-types-toolsettings-use-gpencil-weight-data-add"),
 	("bpy.types.view3doverlay.texture_paint_mode_opacity*", "editors/3dview/display/overlays.html#bpy-types-view3doverlay-texture-paint-mode-opacity"),
@@ -184,6 +180,10 @@ url_manual_mapping = (
 	("bpy.types.cyclesrendersettings.adaptive_threshold*", "render/cycles/render_settings/sampling.html#bpy-types-cyclesrendersettings-adaptive-threshold"),
 	("bpy.types.cyclesrendersettings.camera_cull_margin*", "render/cycles/render_settings/simplify.html#bpy-types-cyclesrendersettings-camera-cull-margin"),
 	("bpy.types.cyclesrendersettings.debug_use_hair_bvh*", "render/cycles/render_settings/performance.html#bpy-types-cyclesrendersettings-debug-use-hair-bvh"),
+	("bpy.types.cyclesrendersettings.mesh_light_samples*", "render/cycles/render_settings/sampling.html#bpy-types-cyclesrendersettings-mesh-light-samples"),
+	("bpy.types.cyclesrendersettings.preview_aa_samples*", "render/cycles/render_settings/sampling.html#bpy-types-cyclesrendersettings-preview-aa-samples"),
+	("bpy.types.cyclesrendersettings.subsurface_samples*", "render/cycles/render_settings/sampling.html#bpy-types-cyclesrendersettings-subsurface-samples"),
+	("bpy.types.cyclesrendersettings.use_square_samples*", "render/cycles/render_settings/sampling.html#bpy-types-cyclesrendersettings-use-square-samples"),
 	("bpy.types.fluiddomainsettings.export_manta_script*", "physics/fluid/type/domain/cache.html#bpy-types-fluiddomainsettings-export-manta-script"),
 	("bpy.types.fluiddomainsettings.fractions_threshold*", "physics/fluid/type/domain/settings.html#bpy-types-fluiddomainsettings-fractions-threshold"),
 	("bpy.types.fluiddomainsettings.particle_band_width*", "physics/fluid/type/domain/settings.html#bpy-types-fluiddomainsettings-particle-band-width"),
@@ -201,12 +201,11 @@ url_manual_mapping = (
 	("bpy.types.movietrackingsettings.use_tripod_solver*", "movie_clip/tracking/clip/toolbar/solve.html#bpy-types-movietrackingsettings-use-tripod-solver"),
 	("bpy.types.rendersettings.simplify_child_particles*", "render/cycles/render_settings/simplify.html#bpy-types-rendersettings-simplify-child-particles"),
 	("bpy.types.rendersettings.use_high_quality_normals*", "render/eevee/render_settings/performance.html#bpy-types-rendersettings-use-high-quality-normals"),
-	("bpy.types.sequencerpreviewoverlay.show_annotation*", "video_editing/preview/introduction.html#bpy-types-sequencerpreviewoverlay-show-annotation"),
-	("bpy.types.sequencerpreviewoverlay.show_safe_areas*", "video_editing/preview/introduction.html#bpy-types-sequencerpreviewoverlay-show-safe-areas"),
 	("bpy.types.sequencertoolsettings.snap_ignore_muted*", "video_editing/sequencer/editing.html#bpy-types-sequencertoolsettings-snap-ignore-muted"),
 	("bpy.types.sequencertoolsettings.snap_ignore_sound*", "video_editing/sequencer/editing.html#bpy-types-sequencertoolsettings-snap-ignore-sound"),
 	("bpy.types.spaceoutliner.use_filter_case_sensitive*", "editors/outliner/interface.html#bpy-types-spaceoutliner-use-filter-case-sensitive"),
 	("bpy.types.spaceoutliner.use_filter_object_content*", "editors/outliner/interface.html#bpy-types-spaceoutliner-use-filter-object-content"),
+	("bpy.types.spacesequenceeditor.show_strip_duration*", "editors/video_sequencer/sequencer/display.html#bpy-types-spacesequenceeditor-show-strip-duration"),
 	("bpy.types.toolsettings.use_proportional_connected*", "editors/3dview/controls/proportional_editing.html#bpy-types-toolsettings-use-proportional-connected"),
 	("bpy.types.toolsettings.use_proportional_projected*", "editors/3dview/controls/proportional_editing.html#bpy-types-toolsettings-use-proportional-projected"),
 	("bpy.types.view3doverlay.vertex_paint_mode_opacity*", "editors/3dview/display/overlays.html#bpy-types-view3doverlay-vertex-paint-mode-opacity"),
@@ -285,13 +284,13 @@ url_manual_mapping = (
 	("bpy.types.materialgpencilstyle.use_fill_holdout*", "grease_pencil/materials/properties.html#bpy-types-materialgpencilstyle-use-fill-holdout"),
 	("bpy.types.particlesettings.use_parent_particles*", "physics/particles/emitter/render.html#bpy-types-particlesettings-use-parent-particles"),
 	("bpy.types.rigidbodyconstraint.solver_iterations*", "physics/rigid_body/constraints/introduction.html#bpy-types-rigidbodyconstraint-solver-iterations"),
-	("bpy.types.sequencerpreviewoverlay.show_metadata*", "video_editing/preview/introduction.html#bpy-types-sequencerpreviewoverlay-show-metadata"),
-	("bpy.types.sequencertimelineoverlay.show_fcurves*", "editors/video_sequencer/sequencer/display.html#bpy-types-sequencertimelineoverlay-show-fcurves"),
 	("bpy.types.spaceclipeditor.use_grayscale_preview*", "editors/clip/display/clip_display.html#bpy-types-spaceclipeditor-use-grayscale-preview"),
 	("bpy.types.spaceoutliner.use_filter_lib_override*", "editors/outliner/interface.html#bpy-types-spaceoutliner-use-filter-lib-override"),
 	("bpy.types.spaceoutliner.use_filter_object_empty*", "editors/outliner/interface.html#bpy-types-spaceoutliner-use-filter-object-empty"),
 	("bpy.types.spaceoutliner.use_filter_object_light*", "editors/outliner/interface.html#bpy-types-spaceoutliner-use-filter-object-light"),
 	("bpy.types.spacesequenceeditor.proxy_render_size*", "video_editing/preview/sidebar.html#bpy-types-spacesequenceeditor-proxy-render-size"),
+	("bpy.types.spacesequenceeditor.show_strip_offset*", "editors/video_sequencer/sequencer/display.html#bpy-types-spacesequenceeditor-show-strip-offset"),
+	("bpy.types.spacesequenceeditor.show_strip_source*", "editors/video_sequencer/sequencer/display.html#bpy-types-spacesequenceeditor-show-strip-source"),
 	("bpy.types.spacespreadsheetrowfilter.column_name*", "editors/spreadsheet.html#bpy-types-spacespreadsheetrowfilter-column-name"),
 	("bpy.types.toolsettings.gpencil_stroke_placement*", "grease_pencil/modes/draw/stroke_placement.html#bpy-types-toolsettings-gpencil-stroke-placement"),
 	("bpy.types.toolsettings.use_keyframe_cycle_aware*", "editors/timeline.html#bpy-types-toolsettings-use-keyframe-cycle-aware"),
@@ -302,6 +301,7 @@ url_manual_mapping = (
 	("bpy.types.cyclesobjectsettings.use_camera_cull*", "render/cycles/object_settings/object_data.html#bpy-types-cyclesobjectsettings-use-camera-cull"),
 	("bpy.types.cyclesobjectsettings.use_motion_blur*", "render/cycles/object_settings/object_data.html#bpy-types-cyclesobjectsettings-use-motion-blur"),
 	("bpy.types.cyclesrendersettings.diffuse_bounces*", "render/cycles/render_settings/light_paths.html#bpy-types-cyclesrendersettings-diffuse-bounces"),
+	("bpy.types.cyclesrendersettings.diffuse_samples*", "render/cycles/render_settings/sampling.html#bpy-types-cyclesrendersettings-diffuse-samples"),
 	("bpy.types.cyclesrendersettings.preview_samples*", "render/cycles/render_settings/sampling.html#bpy-types-cyclesrendersettings-preview-samples"),
 	("bpy.types.cyclesrendersettings.use_camera_cull*", "render/cycles/render_settings/simplify.html#bpy-types-cyclesrendersettings-use-camera-cull"),
 	("bpy.types.cyclesworldsettings.volume_step_size*", "render/cycles/world_settings.html#bpy-types-cyclesworldsettings-volume-step-size"),
@@ -340,7 +340,9 @@ url_manual_mapping = (
 	("bpy.types.cyclesmaterialsettings.displacement*", "render/cycles/material_settings.html#bpy-types-cyclesmaterialsettings-displacement"),
 	("bpy.types.cyclesrendersettings.debug_bvh_type*", "render/cycles/render_settings/debug.html#bpy-types-cyclesrendersettings-debug-bvh-type"),
 	("bpy.types.cyclesrendersettings.glossy_bounces*", "render/cycles/render_settings/light_paths.html#bpy-types-cyclesrendersettings-glossy-bounces"),
+	("bpy.types.cyclesrendersettings.glossy_samples*", "render/cycles/render_settings/sampling.html#bpy-types-cyclesrendersettings-glossy-samples"),
 	("bpy.types.cyclesrendersettings.volume_bounces*", "render/cycles/render_settings/light_paths.html#bpy-types-cyclesrendersettings-volume-bounces"),
+	("bpy.types.cyclesrendersettings.volume_samples*", "render/cycles/render_settings/sampling.html#bpy-types-cyclesrendersettings-volume-samples"),
 	("bpy.types.cyclesworldsettings.sampling_method*", "render/cycles/world_settings.html#bpy-types-cyclesworldsettings-sampling-method"),
 	("bpy.types.cyclesworldsettings.volume_sampling*", "render/cycles/world_settings.html#bpy-types-cyclesworldsettings-volume-sampling"),
 	("bpy.types.editbone.bbone_handle_use_scale_end*", "animation/armatures/bones/properties/bendy_bones.html#bpy-types-editbone-bbone-handle-use-scale-end"),
@@ -368,12 +370,16 @@ url_manual_mapping = (
 	("bpy.types.spacegrapheditor.show_extrapolation*", "editors/graph_editor/introduction.html#bpy-types-spacegrapheditor-show-extrapolation"),
 	("bpy.types.spaceoutliner.use_filter_collection*", "editors/outliner/interface.html#bpy-types-spaceoutliner-use-filter-collection"),
 	("bpy.types.spacesequenceeditor.display_channel*", "video_editing/preview/sidebar.html#bpy-types-spacesequenceeditor-display-channel"),
+	("bpy.types.spacesequenceeditor.show_annotation*", "video_editing/preview/introduction.html#bpy-types-spacesequenceeditor-show-annotation"),
 	("bpy.types.spacesequenceeditor.show_region_hud*", "video_editing/sequencer/navigating.html#bpy-types-spacesequenceeditor-show-region-hud"),
+	("bpy.types.spacesequenceeditor.show_safe_areas*", "video_editing/preview/introduction.html#bpy-types-spacesequenceeditor-show-safe-areas"),
+	("bpy.types.spacesequenceeditor.show_strip_name*", "editors/video_sequencer/sequencer/display.html#bpy-types-spacesequenceeditor-show-strip-name"),
 	("bpy.types.spacespreadsheet.show_only_selected*", "editors/spreadsheet.html#bpy-types-spacespreadsheet-show-only-selected"),
 	("bpy.types.spacespreadsheetrowfilter.operation*", "editors/spreadsheet.html#bpy-types-spacespreadsheetrowfilter-operation"),
 	("bpy.types.spacespreadsheetrowfilter.threshold*", "editors/spreadsheet.html#bpy-types-spacespreadsheetrowfilter-threshold"),
 	("bpy.types.toolsettings.use_snap_grid_absolute*", "editors/3dview/controls/snapping.html#bpy-types-toolsettings-use-snap-grid-absolute"),
 	("bpy.types.view3doverlay.show_face_orientation*", "editors/3dview/display/overlays.html#bpy-types-view3doverlay-show-face-orientation"),
+	("bpy.types.worldlighting.use_ambient_occlusion*", "render/cycles/world_settings.html#bpy-types-worldlighting-use-ambient-occlusion"),
 	("bpy.ops.object.blenderkit_material_thumbnail*", "addons/3d_view/blenderkit.html#bpy-ops-object-blenderkit-material-thumbnail"),
 	("bpy.ops.object.multires_higher_levels_delete*", "modeling/modifiers/generate/multiresolution.html#bpy-ops-object-multires-higher-levels-delete"),
 	("bpy.ops.object.vertex_group_copy_to_selected*", "modeling/meshes/properties/vertex_groups/vertex_groups.html#bpy-ops-object-vertex-group-copy-to-selected"),
@@ -454,9 +460,9 @@ url_manual_mapping = (
 	("bpy.types.sculpt.constant_detail_resolution*", "sculpt_paint/sculpting/tool_settings/dyntopo.html#bpy-types-sculpt-constant-detail-resolution"),
 	("bpy.types.spaceclipeditor.annotation_source*", "movie_clip/tracking/clip/sidebar/view.html#bpy-types-spaceclipeditor-annotation-source"),
 	("bpy.types.spaceclipeditor.show_blue_channel*", "editors/clip/display/clip_display.html#bpy-types-spaceclipeditor-show-blue-channel"),
-	("bpy.types.spacefilebrowser.system_bookmarks*", "editors/file_browser.html#bpy-types-spacefilebrowser-system-bookmarks"),
 	("bpy.types.spaceoutliner.use_filter_children*", "editors/outliner/interface.html#bpy-types-spaceoutliner-use-filter-children"),
 	("bpy.types.spaceoutliner.use_filter_complete*", "editors/outliner/interface.html#bpy-types-spaceoutliner-use-filter-complete"),
+	("bpy.types.spacesequenceeditor.show_metadata*", "video_editing/preview/introduction.html#bpy-types-spacesequenceeditor-show-metadata"),
 	("bpy.types.spacespreadsheet.attribute_domain*", "editors/spreadsheet.html#bpy-types-spacespreadsheet-attribute-domain"),
 	("bpy.types.spacespreadsheetrowfilter.enabled*", "editors/spreadsheet.html#bpy-types-spacespreadsheetrowfilter-enabled"),
 	("bpy.types.spaceview3d.transform_orientation*", "editors/3dview/controls/orientation.html#bpy-types-spaceview3d-transform-orientation"),
@@ -478,10 +484,10 @@ url_manual_mapping = (
 	("bpy.types.cyclesrendersettings.blur_glossy*", "render/cycles/render_settings/light_paths.html#bpy-types-cyclesrendersettings-blur-glossy"),
 	("bpy.types.cyclesrendersettings.dicing_rate*", "render/cycles/render_settings/subdivision.html#bpy-types-cyclesrendersettings-dicing-rate"),
 	("bpy.types.cyclesrendersettings.max_bounces*", "render/cycles/render_settings/light_paths.html#bpy-types-cyclesrendersettings-max-bounces"),
+	("bpy.types.cyclesrendersettings.progressive*", "render/cycles/render_settings/sampling.html#bpy-types-cyclesrendersettings-progressive"),
 	("bpy.types.cyclesrendersettings.use_fast_gi*", "render/cycles/render_settings/light_paths.html#bpy-types-cyclesrendersettings-use-fast-gi"),
 	("bpy.types.editbone.bbone_custom_handle_end*", "animation/armatures/bones/properties/bendy_bones.html#bpy-types-editbone-bbone-custom-handle-end"),
 	("bpy.types.editbone.bbone_handle_type_start*", "animation/armatures/bones/properties/bendy_bones.html#bpy-types-editbone-bbone-handle-type-start"),
-	("bpy.types.fileselectparams.recursion_level*", "editors/file_browser.html#bpy-types-fileselectparams-recursion-level"),
 	("bpy.types.fluiddomainsettings.adapt_margin*", "physics/fluid/type/domain/gas/adaptive_domain.html#bpy-types-fluiddomainsettings-adapt-margin"),
 	("bpy.types.fluiddomainsettings.burning_rate*", "physics/fluid/type/domain/settings.html#bpy-types-fluiddomainsettings-burning-rate"),
 	("bpy.types.fluiddomainsettings.guide_parent*", "physics/fluid/type/domain/guides.html#bpy-types-fluiddomainsettings-guide-parent"),
@@ -510,6 +516,7 @@ url_manual_mapping = (
 	("bpy.types.spaceclipeditor.show_red_channel*", "editors/clip/display/clip_display.html#bpy-types-spaceclipeditor-show-red-channel"),
 	("bpy.types.spaceclipeditor.use_mute_footage*", "editors/clip/display/clip_display.html#bpy-types-spaceclipeditor-use-mute-footage"),
 	("bpy.types.spacesequenceeditor.overlay_type*", "video_editing/preview/sidebar.html#bpy-types-spacesequenceeditor-overlay-type"),
+	("bpy.types.spacesequenceeditor.show_fcurves*", "editors/video_sequencer/sequencer/display.html#bpy-types-spacesequenceeditor-show-fcurves"),
 	("bpy.types.spaceuveditor.sticky_select_mode*", "editors/uv/selecting.html#bpy-types-spaceuveditor-sticky-select-mode"),
 	("bpy.types.spaceview3d.show_object_viewport*", "editors/3dview/display/visibility.html#bpy-types-spaceview3d-show-object-viewport"),
 	("bpy.types.view3doverlay.show_fade_inactive*", "editors/3dview/display/overlays.html#bpy-types-view3doverlay-show-fade-inactive"),
@@ -529,8 +536,10 @@ url_manual_mapping = (
 	("bpy.types.brush.surface_smooth_iterations*", "sculpt_paint/sculpting/tools/smooth.html#bpy-types-brush-surface-smooth-iterations"),
 	("bpy.types.brushgpencilsettings.pen_jitter*", "grease_pencil/modes/draw/tools/draw.html#bpy-types-brushgpencilsettings-pen-jitter"),
 	("bpy.types.cyclescurverendersettings.shape*", "render/cycles/render_settings/hair.html#bpy-types-cyclescurverendersettings-shape"),
+	("bpy.types.cyclesrendersettings.aa_samples*", "render/cycles/render_settings/sampling.html#bpy-types-cyclesrendersettings-aa-samples"),
 	("bpy.types.cyclesrendersettings.ao_bounces*", "render/cycles/render_settings/light_paths.html#bpy-types-cyclesrendersettings-ao-bounces"),
-	("bpy.types.cyclesrendersettings.time_limit*", "render/cycles/render_settings/sampling.html#bpy-types-cyclesrendersettings-time-limit"),
+	("bpy.types.cyclesrendersettings.ao_samples*", "render/cycles/render_settings/sampling.html#bpy-types-cyclesrendersettings-ao-samples"),
+	("bpy.types.cyclesrendersettings.tile_order*", "render/cycles/render_settings/performance.html#bpy-types-cyclesrendersettings-tile-order"),
 	("bpy.types.cyclesvisibilitysettings.camera*", "render/cycles/world_settings.html#bpy-types-cyclesvisibilitysettings-camera"),
 	("bpy.types.cyclesworldsettings.max_bounces*", "render/cycles/world_settings.html#bpy-types-cyclesworldsettings-max-bounces"),
 	("bpy.types.fluiddomainsettings.domain_type*", "physics/fluid/type/domain/settings.html#bpy-types-fluiddomainsettings-domain-type"),
@@ -566,8 +575,6 @@ url_manual_mapping = (
 	("bpy.types.rendersettings.use_single_layer*", "render/layers/view_layer.html#bpy-types-rendersettings-use-single-layer"),
 	("bpy.types.sceneeevee.use_taa_reprojection*", "render/eevee/render_settings/sampling.html#bpy-types-sceneeevee-use-taa-reprojection"),
 	("bpy.types.sequenceeditor.use_overlay_lock*", "video_editing/preview/sidebar.html#bpy-types-sequenceeditor-use-overlay-lock"),
-	("bpy.types.spacefilebrowser.recent_folders*", "editors/file_browser.html#bpy-types-spacefilebrowser-recent-folders"),
-	("bpy.types.spacefilebrowser.system_folders*", "editors/file_browser.html#bpy-types-spacefilebrowser-system-folders"),
 	("bpy.types.spaceoutliner.use_filter_object*", "editors/outliner/interface.html#bpy-types-spaceoutliner-use-filter-object"),
 	("bpy.types.spacesequenceeditor.use_proxies*", "video_editing/preview/sidebar.html#bpy-types-spacesequenceeditor-use-proxies"),
 	("bpy.types.spaceuveditor.show_pixel_coords*", "editors/uv/sidebar.html#bpy-types-spaceuveditor-show-pixel-coords"),
@@ -599,7 +606,6 @@ url_manual_mapping = (
 	("bpy.types.cyclescamerasettings.longitude*", "render/cycles/object_settings/cameras.html#bpy-types-cyclescamerasettings-longitude"),
 	("bpy.types.editbone.bbone_handle_type_end*", "animation/armatures/bones/properties/bendy_bones.html#bpy-types-editbone-bbone-handle-type-end"),
 	("bpy.types.editbone.use_endroll_as_inroll*", "animation/armatures/bones/properties/bendy_bones.html#bpy-types-editbone-use-endroll-as-inroll"),
-	("bpy.types.fileselectparams.filter_search*", "editors/file_browser.html#bpy-types-fileselectparams-filter-search"),
 	("bpy.types.fluiddomainsettings.cache_type*", "physics/fluid/type/domain/cache.html#bpy-types-fluiddomainsettings-cache-type"),
 	("bpy.types.fluiddomainsettings.flip_ratio*", "physics/fluid/type/domain/settings.html#bpy-types-fluiddomainsettings-flip-ratio"),
 	("bpy.types.fluiddomainsettings.guide_beta*", "physics/fluid/type/domain/guides.html#bpy-types-fluiddomainsettings-guide-beta"),
@@ -668,8 +674,6 @@ url_manual_mapping = (
 	("bpy.types.cyclesrendersettings.caustics*", "render/cycles/render_settings/light_paths.html#bpy-types-cyclesrendersettings-caustics"),
 	("bpy.types.cyclesrendersettings.denoiser*", "render/cycles/render_settings/sampling.html#bpy-types-cyclesrendersettings-denoiser"),
 	("bpy.types.editbone.use_inherit_rotation*", "animation/armatures/bones/properties/relations.html#bpy-types-editbone-use-inherit-rotation"),
-	("bpy.types.fileselectparams.display_size*", "editors/file_browser.html#bpy-types-fileselectparams-display-size"),
-	("bpy.types.fileselectparams.display_type*", "editors/file_browser.html#bpy-types-fileselectparams-display-type"),
 	("bpy.types.fluiddomainsettings.use_guide*", "physics/fluid/type/domain/guides.html#bpy-types-fluiddomainsettings-use-guide"),
 	("bpy.types.fluiddomainsettings.use_noise*", "physics/fluid/type/domain/gas/noise.html#bpy-types-fluiddomainsettings-use-noise"),
 	("bpy.types.fluiddomainsettings.use_slice*", "physics/fluid/type/domain/gas/viewport_display.html#bpy-types-fluiddomainsettings-use-slice"),
@@ -732,8 +736,6 @@ url_manual_mapping = (
 	("bpy.types.compositornodebrightcontrast*", "compositing/types/color/bright_contrast.html#bpy-types-compositornodebrightcontrast"),
 	("bpy.types.compositornodedoubleedgemask*", "compositing/types/matte/double_edge_mask.html#bpy-types-compositornodedoubleedgemask"),
 	("bpy.types.cyclesrendersettings.samples*", "render/cycles/render_settings/sampling.html#bpy-types-cyclesrendersettings-samples"),
-	("bpy.types.fileselectparams.show_hidden*", "editors/file_browser.html#bpy-types-fileselectparams-show-hidden"),
-	("bpy.types.fileselectparams.sort_method*", "editors/file_browser.html#bpy-types-fileselectparams-sort-method"),
 	("bpy.types.fluiddomainsettings.clipping*", "physics/fluid/type/domain/settings.html#bpy-types-fluiddomainsettings-clipping"),
 	("bpy.types.fluiddomainsettings.use_mesh*", "physics/fluid/type/domain/liquid/mesh.html#bpy-types-fluiddomainsettings-use-mesh"),
 	("bpy.types.freestylelinestyle.angle_max*", "render/freestyle/view_layer/line_style/strokes.html#bpy-types-freestylelinestyle-angle-max"),
@@ -750,7 +752,6 @@ url_manual_mapping = (
 	("bpy.types.object.use_empty_image_alpha*", "modeling/empties.html#bpy-types-object-use-empty-image-alpha"),
 	("bpy.types.rendersettings.frame_map_new*", "render/output/properties/frame_range.html#bpy-types-rendersettings-frame-map-new"),
 	("bpy.types.rendersettings.frame_map_old*", "render/output/properties/frame_range.html#bpy-types-rendersettings-frame-map-old"),
-	("bpy.types.rendersettings.use_auto_tile*", "render/cycles/render_settings/performance.html#bpy-types-rendersettings-use-auto-tile"),
 	("bpy.types.rendersettings.use_overwrite*", "render/output/properties/output.html#bpy-types-rendersettings-use-overwrite"),
 	("bpy.types.rendersettings.use_sequencer*", "render/output/properties/post_processing.html#bpy-types-rendersettings-use-sequencer"),
 	("bpy.types.sceneeevee.volumetric_shadow*", "render/eevee/render_settings/volumetrics.html#bpy-types-sceneeevee-volumetric-shadow"),
@@ -796,7 +797,6 @@ url_manual_mapping = (
 	("bpy.types.compositornodesetalpha.mode*", "compositing/types/converter/set_alpha.html#bpy-types-compositornodesetalpha-mode"),
 	("bpy.types.dopesheet.use_filter_invert*", "editors/graph_editor/channels.html#bpy-types-dopesheet-use-filter-invert"),
 	("bpy.types.editbone.use_local_location*", "animation/armatures/bones/properties/relations.html#bpy-types-editbone-use-local-location"),
-	("bpy.types.fileselectparams.use_filter*", "editors/file_browser.html#bpy-types-fileselectparams-use-filter"),
 	("bpy.types.fluiddomainsettings.gravity*", "physics/fluid/type/domain/settings.html#bpy-types-fluiddomainsettings-gravity"),
 	("bpy.types.fluidflowsettings.flow_type*", "physics/fluid/type/flow.html#bpy-types-fluidflowsettings-flow-type"),
 	("bpy.types.fluidflowsettings.subframes*", "physics/fluid/type/flow.html#bpy-types-fluidflowsettings-subframes"),
@@ -868,7 +868,6 @@ url_manual_mapping = (
 	("bpy.types.compositornodecolorbalance*", "compositing/types/color/color_balance.html#bpy-types-compositornodecolorbalance"),
 	("bpy.types.compositornodekeyingscreen*", "compositing/types/matte/keying_screen.html#bpy-types-compositornodekeyingscreen"),
 	("bpy.types.dynamicpaintcanvassettings*", "physics/dynamic_paint/canvas.html#bpy-types-dynamicpaintcanvassettings"),
-	("bpy.types.fileselectparams.directory*", "editors/file_browser.html#bpy-types-fileselectparams-directory"),
 	("bpy.types.fluidflowsettings.use_flow*", "physics/fluid/type/flow.html#bpy-types-fluidflowsettings-use-flow"),
 	("bpy.types.fmodifierfunctiongenerator*", "editors/graph_editor/fcurves/modifiers.html#bpy-types-fmodifierfunctiongenerator"),
 	("bpy.types.geometrynodeattributeclamp*", "modeling/geometry_nodes/attribute/attribute_clamp.html#bpy-types-geometrynodeattributeclamp"),
@@ -892,7 +891,6 @@ url_manual_mapping = (
 	("bpy.types.shadernodeambientocclusion*", "render/shader_nodes/input/ao.html#bpy-types-shadernodeambientocclusion"),
 	("bpy.types.shadernodevolumeabsorption*", "render/shader_nodes/shader/volume_absorption.html#bpy-types-shadernodevolumeabsorption"),
 	("bpy.types.shadernodevolumeprincipled*", "render/shader_nodes/shader/volume_principled.html#bpy-types-shadernodevolumeprincipled"),
-	("bpy.types.spacefilebrowser.bookmarks*", "editors/file_browser.html#bpy-types-spacefilebrowser-bookmarks"),
 	("bpy.types.spaceoutliner.display_mode*", "editors/outliner/interface.html#bpy-types-spaceoutliner-display-mode"),
 	("bpy.types.spaceoutliner.filter_state*", "editors/outliner/interface.html#bpy-types-spaceoutliner-filter-state"),
 	("bpy.types.toolsettings.keyframe_type*", "editors/timeline.html#bpy-types-toolsettings-keyframe-type"),
@@ -942,7 +940,6 @@ url_manual_mapping = (
 	("bpy.types.cyclesrendersettings.seed*", "render/cycles/render_settings/sampling.html#bpy-types-cyclesrendersettings-seed"),
 	("bpy.types.dynamicpaintbrushsettings*", "physics/dynamic_paint/brush.html#bpy-types-dynamicpaintbrushsettings"),
 	("bpy.types.editbone.use_scale_easing*", "animation/armatures/bones/properties/bendy_bones.html#bpy-types-editbone-use-scale-easing"),
-	("bpy.types.fileselectparams.filename*", "editors/file_browser.html#bpy-types-fileselectparams-filename"),
 	("bpy.types.fluiddomainsettings.alpha*", "physics/fluid/type/domain/settings.html#bpy-types-fluiddomainsettings-alpha"),
 	("bpy.types.fluidflowsettings.density*", "physics/fluid/type/flow.html#bpy-types-fluidflowsettings-density"),
 	("bpy.types.freestylelineset.qi_start*", "render/freestyle/view_layer/line_set.html#bpy-types-freestylelineset-qi-start"),
@@ -1051,8 +1048,6 @@ url_manual_mapping = (
 	("bpy.types.nodesocketinterface.name*", "interface/controls/nodes/groups.html#bpy-types-nodesocketinterface-name"),
 	("bpy.types.object.is_shadow_catcher*", "render/cycles/object_settings/object_data.html#bpy-types-object-is-shadow-catcher"),
 	("bpy.types.particleinstancemodifier*", "modeling/modifiers/physics/particle_instance.html#bpy-types-particleinstancemodifier"),
-	("bpy.types.rendersettings.tile_size*", "render/cycles/render_settings/performance.html#bpy-types-rendersettings-tile-size"),
-	("bpy.types.sequencertimelineoverlay*", "editors/video_sequencer/sequencer/display.html#bpy-types-sequencertimelineoverlay"),
 	("bpy.types.sequencetransform.offset*", "video_editing/sequencer/sidebar/strip.html#bpy-types-sequencetransform-offset"),
 	("bpy.types.shadernodebrightcontrast*", "render/shader_nodes/color/bright_contrast.html#bpy-types-shadernodebrightcontrast"),
 	("bpy.types.shadernodebsdfprincipled*", "render/shader_nodes/shader/principled.html#bpy-types-shadernodebsdfprincipled"),
@@ -1148,7 +1143,6 @@ url_manual_mapping = (
 	("bpy.types.rendersettings.fps_base*", "render/output/properties/format.html#bpy-types-rendersettings-fps-base"),
 	("bpy.types.rigidbodyobject.enabled*", "physics/rigid_body/properties/settings.html#bpy-types-rigidbodyobject-enabled"),
 	("bpy.types.sceneeevee.use_overscan*", "render/eevee/render_settings/film.html#bpy-types-sceneeevee-use-overscan"),
-	("bpy.types.sequencerpreviewoverlay*", "video_editing/preview/introduction.html#bpy-types-sequencerpreviewoverlay"),
 	("bpy.types.sequencetransform.scale*", "video_editing/sequencer/sidebar/strip.html#bpy-types-sequencetransform-scale"),
 	("bpy.types.shadernodeeeveespecular*", "render/shader_nodes/shader/specular_bsdf.html#bpy-types-shadernodeeeveespecular"),
 	("bpy.types.shadernodehuesaturation*", "render/shader_nodes/color/hue_saturation.html#bpy-types-shadernodehuesaturation"),
@@ -1159,7 +1153,7 @@ url_manual_mapping = (
 	("bpy.types.vertexweightmixmodifier*", "modeling/modifiers/modify/weight_mix.html#bpy-types-vertexweightmixmodifier"),
 	("bpy.types.viewlayer.use_freestyle*", "render/freestyle/view_layer/freestyle.html#bpy-types-viewlayer-use-freestyle"),
 	("bpy.types.volumedisplay.use_slice*", "modeling/volumes/properties.html#bpy-types-volumedisplay-use-slice"),
-	("bpy.types.worldlighting.ao_factor*", "render/cycles/render_settings/light_paths.html#bpy-types-worldlighting-ao-factor"),
+	("bpy.types.worldlighting.ao_factor*", "render/cycles/world_settings.html#bpy-types-worldlighting-ao-factor"),
 	("bpy.types.worldmistsettings.depth*", "render/cycles/world_settings.html#bpy-types-worldmistsettings-depth"),
 	("bpy.types.worldmistsettings.start*", "render/cycles/world_settings.html#bpy-types-worldmistsettings-start"),
 	("bpy.ops.armature.armature_layers*", "animation/armatures/bones/editing/change_layers.html#bpy-ops-armature-armature-layers"),
@@ -1261,11 +1255,11 @@ url_manual_mapping = (
 	("bpy.types.shadernodevectorrotate*", "render/shader_nodes/vector/vector_rotate.html#bpy-types-shadernodevectorrotate"),
 	("bpy.types.sound.use_memory_cache*", "video_editing/sequencer/sidebar/strip.html#bpy-types-sound-use-memory-cache"),
 	("bpy.types.spaceview3d.show_gizmo*", "editors/3dview/display/gizmo.html#bpy-types-spaceview3d-show-gizmo"),
-	("bpy.types.texturegpencilmodifier*", "grease_pencil/modifiers/modify/texture_mapping.html#bpy-types-texturegpencilmodifier"),
+	("bpy.types.texturegpencilmodifier*", "grease_pencil/modifiers/color/texture_mapping.html#bpy-types-texturegpencilmodifier"),
 	("bpy.types.volumedisplacemodifier*", "modeling/modifiers/deform/volume_displace.html#bpy-types-volumedisplacemodifier"),
 	("bpy.types.volumerender.step_size*", "modeling/volumes/properties.html#bpy-types-volumerender-step-size"),
 	("bpy.types.weightednormalmodifier*", "modeling/modifiers/modify/weighted_normal.html#bpy-types-weightednormalmodifier"),
-	("bpy.types.worldlighting.distance*", "render/cycles/render_settings/light_paths.html#bpy-types-worldlighting-distance"),
+	("bpy.types.worldlighting.distance*", "render/cycles/world_settings.html#bpy-types-worldlighting-distance"),
 	("bpy.ops.armature.autoside_names*", "animation/armatures/bones/editing/naming.html#bpy-ops-armature-autoside-names"),
 	("bpy.ops.armature.calculate_roll*", "animation/armatures/bones/editing/bone_roll.html#bpy-ops-armature-calculate-roll"),
 	("bpy.ops.armature.duplicate_move*", "animation/armatures/bones/editing/duplicate.html#bpy-ops-armature-duplicate-move"),
@@ -1366,6 +1360,8 @@ url_manual_mapping = (
 	("bpy.types.object.visible_shadow*", "render/cycles/object_settings/object_data.html#bpy-types-object-visible-shadow"),
 	("bpy.types.offsetgpencilmodifier*", "grease_pencil/modifiers/deform/offset.html#bpy-types-offsetgpencilmodifier"),
 	("bpy.types.posebone.custom_shape*", "animation/armatures/bones/properties/display.html#bpy-types-posebone-custom-shape"),
+	("bpy.types.rendersettings.tile_x*", "render/cycles/render_settings/performance.html#bpy-types-rendersettings-tile-x"),
+	("bpy.types.rendersettings.tile_y*", "render/cycles/render_settings/performance.html#bpy-types-rendersettings-tile-y"),
 	("bpy.types.rigifyselectioncolors*", "addons/rigging/rigify/metarigs.html#bpy-types-rigifyselectioncolors"),
 	("bpy.types.sceneeevee.volumetric*", "render/eevee/render_settings/volumetrics.html#bpy-types-sceneeevee-volumetric"),
 	("bpy.types.screen.show_statusbar*", "interface/window_system/topbar.html#bpy-types-screen-show-statusbar"),
@@ -1577,7 +1573,7 @@ url_manual_mapping = (
 	("bpy.types.stretchtoconstraint*", "animation/constraints/tracking/stretch_to.html#bpy-types-stretchtoconstraint"),
 	("bpy.types.texturenodecurvergb*", "editors/texture_node/types/color/rgb_curves.html#bpy-types-texturenodecurvergb"),
 	("bpy.types.texturenodevaltorgb*", "editors/texture_node/types/converter/rgb_to_bw.html#bpy-types-texturenodevaltorgb"),
-	("bpy.types.timegpencilmodifier*", "grease_pencil/modifiers/modify/time_offset.html#bpy-types-timegpencilmodifier"),
+	("bpy.types.timegpencilmodifier*", "grease_pencil/modifiers/deform/time_offset.html#bpy-types-timegpencilmodifier"),
 	("bpy.types.tintgpencilmodifier*", "grease_pencil/modifiers/color/tint.html#bpy-types-tintgpencilmodifier"),
 	("bpy.types.transformconstraint*", "animation/constraints/transform/transformation.html#bpy-types-transformconstraint"),
 	("bpy.types.triangulatemodifier*", "modeling/modifiers/generate/triangulate.html#bpy-types-triangulatemodifier"),
@@ -1656,7 +1652,6 @@ url_manual_mapping = (
 	("bpy.types.curve.twist_smooth*", "modeling/curves/properties/shape.html#bpy-types-curve-twist-smooth"),
 	("bpy.types.curvepaintsettings*", "modeling/curves/tools/draw.html#bpy-types-curvepaintsettings"),
 	("bpy.types.fcurve.array_index*", "editors/graph_editor/fcurves/properties.html#bpy-types-fcurve-array-index"),
-	("bpy.types.fileselectidfilter*", "editors/file_browser.html#bpy-types-fileselectidfilter"),
 	("bpy.types.fmodifiergenerator*", "editors/graph_editor/fcurves/modifiers.html#bpy-types-fmodifiergenerator"),
 	("bpy.types.freestylelinestyle*", "render/freestyle/view_layer/line_style/index.html#bpy-types-freestylelinestyle"),
 	("bpy.types.gammacrosssequence*", "video_editing/sequencer/strips/transitions/gamma_cross.html#bpy-types-gammacrosssequence"),
@@ -1825,7 +1820,6 @@ url_manual_mapping = (
 	("bpy.ops.clip.track_markers*", "movie_clip/tracking/clip/editing/track.html#bpy-ops-clip-track-markers"),
 	("bpy.ops.curve.extrude_move*", "modeling/curves/editing/control_points.html#bpy-ops-curve-extrude-move"),
 	("bpy.ops.curve.make_segment*", "modeling/curves/editing/control_points.html#bpy-ops-curve-make-segment"),
-	("bpy.ops.file.directory_new*", "editors/file_browser.html#bpy-ops-file-directory-new"),
 	("bpy.ops.graph.euler_filter*", "editors/graph_editor/fcurves/editing.html#bpy-ops-graph-euler-filter"),
 	("bpy.ops.marker.camera_bind*", "animation/markers.html#bpy-ops-marker-camera-bind"),
 	("bpy.ops.mask.select_circle*", "movie_clip/masking/selecting.html#bpy-ops-mask-select-circle"),
@@ -1880,7 +1874,6 @@ url_manual_mapping = (
 	("bpy.types.editbone.bbone_x*", "animation/armatures/bones/properties/bendy_bones.html#bpy-types-editbone-bbone-x"),
 	("bpy.types.editbone.bbone_z*", "animation/armatures/bones/properties/bendy_bones.html#bpy-types-editbone-bbone-z"),
 	("bpy.types.fcurve.data_path*", "editors/graph_editor/fcurves/properties.html#bpy-types-fcurve-data-path"),
-	("bpy.types.fileselectparams*", "editors/file_browser.html#bpy-types-fileselectparams"),
 	("bpy.types.fmodifierstepped*", "editors/graph_editor/fcurves/modifiers.html#bpy-types-fmodifierstepped"),
 	("bpy.types.freestylelineset*", "render/freestyle/view_layer/line_set.html#bpy-types-freestylelineset"),
 	("bpy.types.mask.frame_start*", "movie_clip/masking/sidebar.html#bpy-types-mask-frame-start"),
@@ -1911,7 +1904,6 @@ url_manual_mapping = (
 	("bpy.types.softbodymodifier*", "physics/soft_body/index.html#bpy-types-softbodymodifier"),
 	("bpy.types.softbodysettings*", "physics/soft_body/settings/index.html#bpy-types-softbodysettings"),
 	("bpy.types.solidifymodifier*", "modeling/modifiers/generate/solidify.html#bpy-types-solidifymodifier"),
-	("bpy.types.spacefilebrowser*", "editors/file_browser.html#bpy-types-spacefilebrowser"),
 	("bpy.types.spacegrapheditor*", "editors/graph_editor/index.html#bpy-types-spacegrapheditor"),
 	("bpy.types.spacepreferences*", "editors/preferences/index.html#bpy-types-spacepreferences"),
 	("bpy.types.spacespreadsheet*", "editors/spreadsheet.html#bpy-types-spacespreadsheet"),
@@ -1931,7 +1923,6 @@ url_manual_mapping = (
 	("bpy.ops.clip.solve_camera*", "movie_clip/tracking/clip/editing/track.html#bpy-ops-clip-solve-camera"),
 	("bpy.ops.constraint.delete*", "animation/constraints/interface/header.html#bpy-ops-constraint-delete"),
 	("bpy.ops.curve.smooth_tilt*", "modeling/curves/editing/control_points.html#bpy-ops-curve-smooth-tilt"),
-	("bpy.ops.file.reset_recent*", "editors/file_browser.html#bpy-ops-file-reset-recent"),
 	("bpy.ops.fluid.bake_guides*", "physics/fluid/type/domain/guides.html#bpy-ops-fluid-bake-guides"),
 	("bpy.ops.fluid.free_guides*", "physics/fluid/type/domain/guides.html#bpy-ops-fluid-free-guides"),
 	("bpy.ops.font.style_toggle*", "modeling/texts/editing.html#bpy-ops-font-style-toggle"),
@@ -1953,7 +1944,6 @@ url_manual_mapping = (
 	("bpy.ops.object.origin_set*", "scene_layout/object/origin.html#bpy-ops-object-origin-set"),
 	("bpy.ops.object.parent_set*", "scene_layout/object/editing/parent.html#bpy-ops-object-parent-set"),
 	("bpy.ops.object.pointcloud*", "modeling/point_cloud.html#bpy-ops-object-pointcloud"),
-	("bpy.ops.object.proxy_make*", "files/linked_libraries/library_proxies.html#bpy-ops-object-proxy-make"),
 	("bpy.ops.object.select_all*", "scene_layout/object/selecting.html#bpy-ops-object-select-all"),
 	("bpy.ops.object.shade_flat*", "scene_layout/object/editing/shading.html#bpy-ops-object-shade-flat"),
 	("bpy.ops.pose.group_assign*", "animation/armatures/properties/bone_groups.html#bpy-ops-pose-group-assign"),
@@ -2261,7 +2251,6 @@ url_manual_mapping = (
 	("bpy.ops.clip.prefetch*", "movie_clip/tracking/clip/editing/clip.html#bpy-ops-clip-prefetch"),
 	("bpy.ops.clip.set_axis*", "movie_clip/tracking/clip/editing/reconstruction.html#bpy-ops-clip-set-axis"),
 	("bpy.ops.file.pack_all*", "files/blend/packed_data.html#bpy-ops-file-pack-all"),
-	("bpy.ops.file.previous*", "editors/file_browser.html#bpy-ops-file-previous"),
 	("bpy.ops.gpencil.paste*", "grease_pencil/modes/edit/grease_pencil_menu.html#bpy-ops-gpencil-paste"),
 	("bpy.ops.image.project*", "sculpt_paint/texture_paint/tool_settings/options.html#bpy-ops-image-project"),
 	("bpy.ops.image.replace*", "editors/image/editing.html#bpy-ops-image-replace"),
@@ -2310,8 +2299,6 @@ url_manual_mapping = (
 	("bpy.ops.curve.delete*", "modeling/curves/editing/curve.html#bpy-ops-curve-delete"),
 	("bpy.ops.curve.reveal*", "modeling/curves/editing/curve.html#bpy-ops-curve-reveal"),
 	("bpy.ops.curve.smooth*", "modeling/curves/editing/control_points.html#bpy-ops-curve-smooth"),
-	("bpy.ops.file.execute*", "editors/file_browser.html#bpy-ops-file-execute"),
-	("bpy.ops.file.refresh*", "editors/file_browser.html#bpy-ops-file-refresh"),
 	("bpy.ops.fluid.preset*", "physics/fluid/type/domain/liquid/diffusion.html#bpy-ops-fluid-preset"),
 	("bpy.ops.gpencil.copy*", "grease_pencil/modes/edit/grease_pencil_menu.html#bpy-ops-gpencil-copy"),
 	("bpy.ops.graph.delete*", "editors/graph_editor/fcurves/editing.html#bpy-ops-graph-delete"),
@@ -2353,8 +2340,6 @@ url_manual_mapping = (
 	("bpy.types.vectorfont*", "modeling/texts/index.html#bpy-types-vectorfont"),
 	("bpy.ops.clip.reload*", "movie_clip/tracking/clip/editing/clip.html#bpy-ops-clip-reload"),
 	("bpy.ops.curve.split*", "modeling/curves/editing/curve.html#bpy-ops-curve-split"),
-	("bpy.ops.file.cancel*", "editors/file_browser.html#bpy-ops-file-cancel"),
-	("bpy.ops.file.parent*", "editors/file_browser.html#bpy-ops-file-parent"),
 	("bpy.ops.graph.clean*", "editors/graph_editor/fcurves/editing.html#bpy-ops-graph-clean"),
 	("bpy.ops.graph.paste*", "editors/graph_editor/fcurves/editing.html#bpy-ops-graph-paste"),
 	("bpy.ops.mesh.bisect*", "modeling/meshes/editing/mesh/bisect.html#bpy-ops-mesh-bisect"),
@@ -2434,7 +2419,6 @@ url_manual_mapping = (
 	("bpy.types.spacenla*", "editors/nla/index.html#bpy-types-spacenla"),
 	("bpy.types.sunlight*", "render/lights/light_object.html#bpy-types-sunlight"),
 	("bpy.ops.clip.open*", "movie_clip/tracking/clip/editing/clip.html#bpy-ops-clip-open"),
-	("bpy.ops.file.next*", "editors/file_browser.html#bpy-ops-file-next"),
 	("bpy.ops.image.new*", "editors/image/editing.html#bpy-ops-image-new"),
 	("bpy.ops.mesh.fill*", "modeling/meshes/editing/face/fill.html#bpy-ops-mesh-fill"),
 	("bpy.ops.mesh.poke*", "modeling/meshes/editing/face/poke_faces.html#bpy-ops-mesh-poke"),
@@ -2545,7 +2529,6 @@ url_manual_mapping = (
 	("bpy.ops.anim*", "animation/index.html#bpy-ops-anim"),
 	("bpy.ops.boid*", "physics/particles/emitter/physics/boids.html#bpy-ops-boid"),
 	("bpy.ops.clip*", "movie_clip/index.html#bpy-ops-clip"),
-	("bpy.ops.file*", "editors/file_browser.html#bpy-ops-file"),
 	("bpy.ops.font*", "modeling/texts/index.html#bpy-ops-font"),
 	("bpy.ops.mask*", "movie_clip/masking/index.html#bpy-ops-mask"),
 	("bpy.ops.mesh*", "modeling/meshes/index.html#bpy-ops-mesh"),

release/scripts/presets/keyconfig/Blender.py

@@ -217,15 +217,6 @@ class Prefs(bpy.types.KeyConfigPreferences):
         update=update_fn,
     )
 
-    use_file_single_click: BoolProperty(
-        name="Open Folders on Single Click",
-        description=(
-            "Navigate into folders by clicking on them once instead of twice"
-        ),
-        default=False,
-        update=update_fn,
-    )
-
     def draw(self, layout):
         from bpy import context
 
@@ -282,10 +273,6 @@ class Prefs(bpy.types.KeyConfigPreferences):
         sub.prop(self, "use_pie_click_drag")
         sub.prop(self, "use_v3d_shade_ex_pie")
 
-        # File Browser settings.
-        col = layout.column()
-        col.label(text="File Browser")
-        col.row().prop(self, "use_file_single_click")
 
 blender_default = bpy.utils.execfile(os.path.join(DIRNAME, "keymap_data", "blender_default.py"))
 
@@ -325,7 +312,6 @@ def load():
             ),
             use_alt_click_leader=kc_prefs.use_alt_click_leader,
             use_pie_click_drag=kc_prefs.use_pie_click_drag,
-            use_file_single_click=kc_prefs.use_file_single_click,
         ),
     )
 

release/scripts/presets/keyconfig/keymap_data/blender_default.py

@@ -66,7 +66,6 @@ class Params:
         # Alt-MMB axis switching 'RELATIVE' or 'ABSOLUTE' axis switching.
         "v3d_alt_mmb_drag_action",
 
-        "use_file_single_click",
         # Convenience variables:
         # (derived from other settings).
         #
@@ -77,14 +76,7 @@ class Params:
         "select_mouse_value_fallback",
         # Shorthand for: `('CLICK_DRAG' if params.use_pie_click_drag else 'PRESS')`
         "pie_value",
-        # Shorthand for: `{"type": params.tool_tweak, "value": 'ANY'}`.
-        "tool_tweak_event",
         # Shorthand for: `{"type": params.tool_maybe_tweak, "value": params.tool_maybe_tweak_value}`.
-        #
-        # NOTE: This is typically used for active tool key-map items however it should never
-        # be used for selection tools (the default box-select tool for example).
-        # Since this means with RMB select enabled in edit-mode for e.g.
-        # `Ctrl-LMB` would be caught by box-select instead of add/extrude.
         "tool_maybe_tweak_event",
     )
 
@@ -107,7 +99,6 @@ class Params:
             use_alt_tool_or_cursor=False,
             use_alt_click_leader=False,
             use_pie_click_drag=False,
-            use_file_single_click=False,
             v3d_tilde_action='VIEW',
             v3d_alt_mmb_drag_action='RELATIVE',
     ):
@@ -192,13 +183,10 @@ class Params:
         self.use_alt_click_leader = use_alt_click_leader
         self.use_pie_click_drag = use_pie_click_drag
 
-        self.use_file_single_click = use_file_single_click
-
         # Convenience variables:
         self.use_fallback_tool_rmb = self.use_fallback_tool if self.select_mouse == 'RIGHT' else False
         self.select_mouse_value_fallback = 'CLICK' if self.use_fallback_tool_rmb else self.select_mouse_value
         self.pie_value = 'CLICK_DRAG' if use_pie_click_drag else 'PRESS'
-        self.tool_tweak_event = {"type": self.tool_tweak, "value": 'ANY'}
         self.tool_maybe_tweak_event = {"type": self.tool_maybe_tweak, "value": self.tool_maybe_tweak_value}
 
 
@@ -2155,20 +2143,16 @@ def km_file_browser_main(params):
         {"items": items},
     )
 
-    if not params.use_file_single_click:
-        items.extend([
-            ("file.select", {"type": 'LEFTMOUSE', "value": 'DOUBLE_CLICK'},
-            {"properties": [("open", True), ("deselect_all", not params.legacy)]}),
-        ])
-
     items.extend([
         ("file.mouse_execute", {"type": 'LEFTMOUSE', "value": 'DOUBLE_CLICK'}, None),
         # Both .execute and .select are needed here. The former only works if
         # there's a file operator (i.e. not in regular editor mode) but is
         # needed to load files. The latter makes selection work if there's no
         # operator (i.e. in regular editor mode).
+        ("file.select", {"type": 'LEFTMOUSE', "value": 'DOUBLE_CLICK'},
+         {"properties": [("open", True), ("deselect_all", not params.legacy)]}),
         ("file.select", {"type": 'LEFTMOUSE', "value": 'PRESS'},
-         {"properties": [("open", params.use_file_single_click), ("deselect_all", not params.legacy)]}),
+         {"properties": [("open", False), ("deselect_all", not params.legacy)]}),
         ("file.select", {"type": 'LEFTMOUSE', "value": 'CLICK', "ctrl": True},
          {"properties": [("extend", True), ("open", False)]}),
         ("file.select", {"type": 'LEFTMOUSE', "value": 'CLICK', "shift": True},
@@ -5685,9 +5669,6 @@ def km_knife_tool_modal_map(_params):
         ("IGNORE_SNAP_OFF", {"type": 'LEFT_CTRL', "value": 'RELEASE', "any": True}, None),
         ("IGNORE_SNAP_ON", {"type": 'RIGHT_CTRL', "value": 'PRESS', "any": True}, None),
         ("IGNORE_SNAP_OFF", {"type": 'RIGHT_CTRL', "value": 'RELEASE', "any": True}, None),
-        ("X_AXIS", {"type": 'X', "value": 'PRESS'}, None),
-        ("Y_AXIS", {"type": 'Y', "value": 'PRESS'}, None),
-        ("Z_AXIS", {"type": 'Z', "value": 'PRESS'}, None),
         ("ANGLE_SNAP_TOGGLE", {"type": 'A', "value": 'PRESS'}, None),
         ("CYCLE_ANGLE_SNAP_EDGE", {"type": 'R', "value": 'PRESS'}, None),
         ("CUT_THROUGH_TOGGLE", {"type": 'C', "value": 'PRESS'}, None),
@@ -6144,9 +6125,12 @@ def km_image_editor_tool_uv_cursor(params):
         "Image Editor Tool: Uv, Cursor",
         {"space_type": 'IMAGE_EDITOR', "region_type": 'WINDOW'},
         {"items": [
-            ("uv.cursor_set", {"type": params.tool_mouse, "value": 'PRESS'}, None),
+            ("uv.cursor_set",
+             {"type": params.tool_mouse, "value": 'PRESS'},
+             None),
             # Don't use `tool_maybe_tweak_event` since it conflicts with `PRESS` that places the cursor.
-            ("transform.translate", params.tool_tweak_event,
+            ("transform.translate",
+             {"type": params.tool_tweak, "value": 'ANY'},
              {"properties": [("release_confirm", True), ("cursor_transform", True)]}),
         ]},
     )
@@ -6171,8 +6155,7 @@ def km_image_editor_tool_uv_select_box(params, *, fallback):
         {"items": [
             *([] if (fallback and not params.use_fallback_tool) else _template_items_tool_select_actions_simple(
                 "uv.select_box",
-                # Don't use `tool_maybe_tweak_event`, see comment for this slot.
-                **({"type": params.select_tweak, "value": 'ANY'} if fallback else params.tool_tweak_event))),
+                **({"type": params.select_tweak, "value": 'ANY'} if fallback else params.tool_maybe_tweak_event))),
             *_template_uv_select_for_fallback(params, fallback),
         ]},
     )
@@ -6201,7 +6184,7 @@ def km_image_editor_tool_uv_select_lasso(params, *, fallback):
         {"items": [
             *([] if (fallback and not params.use_fallback_tool) else _template_items_tool_select_actions_simple(
                 "uv.select_lasso",
-                **({"type": params.select_tweak, "value": 'ANY'} if fallback else params.tool_tweak_event))),
+                **({"type": params.select_tweak, "value": 'ANY'} if fallback else params.tool_maybe_tweak_event))),
             *_template_uv_select_for_fallback(params, fallback),
         ]},
     )
@@ -6338,7 +6321,7 @@ def km_3d_view_tool_cursor(params):
         {"items": [
             ("view3d.cursor3d", {"type": params.tool_mouse, "value": 'PRESS'}, None),
             # Don't use `tool_maybe_tweak_event` since it conflicts with `PRESS` that places the cursor.
-            ("transform.translate", params.tool_tweak_event,
+            ("transform.translate", {"type": params.tool_tweak, "value": 'ANY'},
              {"properties": [("release_confirm", True), ("cursor_transform", True)]}),
         ]},
     )
@@ -6364,8 +6347,7 @@ def km_3d_view_tool_select_box(params, *, fallback):
         {"items": [
             *([] if (fallback and not params.use_fallback_tool) else _template_items_tool_select_actions(
                 "view3d.select_box",
-                # Don't use `tool_maybe_tweak_event`, see comment for this slot.
-                **({"type": params.select_tweak, "value": 'ANY'} if fallback else params.tool_tweak_event))),
+                **({"type": params.select_tweak, "value": 'ANY'} if fallback else params.tool_maybe_tweak_event))),
             *_template_view3d_select_for_fallback(params, fallback),
         ]},
     )
@@ -6395,7 +6377,7 @@ def km_3d_view_tool_select_lasso(params, *, fallback):
         {"items": [
             *([] if (fallback and not params.use_fallback_tool) else _template_items_tool_select_actions(
                 "view3d.select_lasso",
-                **({"type": params.select_tweak, "value": 'ANY'} if fallback else params.tool_tweak_event))),
+                **({"type": params.select_tweak, "value": 'ANY'} if fallback else params.tool_maybe_tweak_event))),
             *_template_view3d_select_for_fallback(params, fallback),
         ]}
     )
@@ -7050,8 +7032,10 @@ def km_3d_view_tool_sculpt_mask_by_color(params):
         "3D View Tool: Sculpt, Mask by Color",
         {"space_type": 'VIEW_3D', "region_type": 'WINDOW'},
         {"items": [
-            ("sculpt.mask_by_color", {"type": params.tool_mouse, "value": 'ANY'}, None),
-            ("sculpt.mask_by_color", params.tool_tweak_event, None),
+            ("sculpt.mask_by_color", {"type": params.tool_mouse, "value": 'ANY'},
+             None),
+            ("sculpt.mask_by_color", {"type": params.tool_tweak, "value": 'ANY'},
+             None),
         ]},
     )
 
@@ -7061,7 +7045,8 @@ def km_3d_view_tool_sculpt_face_set_edit(params):
         "3D View Tool: Sculpt, Face Set Edit",
         {"space_type": 'VIEW_3D', "region_type": 'WINDOW'},
         {"items": [
-            ("sculpt.face_set_edit", {"type": params.tool_mouse, "value": 'PRESS'}, None),
+            ("sculpt.face_set_edit", {"type": params.tool_mouse, "value": 'PRESS'},
+             None),
         ]},
     )
 
@@ -7250,8 +7235,7 @@ def km_3d_view_tool_edit_gpencil_select_box(params, *, fallback):
         {"items": [
             *([] if (fallback and not params.use_fallback_tool) else _template_items_tool_select_actions(
                 "gpencil.select_box",
-                # Don't use `tool_maybe_tweak_event`, see comment for this slot.
-                **({"type": params.select_tweak, "value": 'ANY'} if fallback else params.tool_tweak_event))),
+                **({"type": params.select_tweak, "value": 'ANY'} if fallback else params.tool_maybe_tweak_event))),
             *_template_view3d_gpencil_select_for_fallback(params, fallback),
         ]},
     )
@@ -7281,7 +7265,7 @@ def km_3d_view_tool_edit_gpencil_select_lasso(params, *, fallback):
         {"items": [
             *([] if (fallback and not params.use_fallback_tool) else _template_items_tool_select_actions(
                 "gpencil.select_lasso",
-                **({"type": params.select_tweak, "value": 'ANY'} if fallback else params.tool_tweak_event))),
+                **({"type": params.select_tweak, "value": 'ANY'} if fallback else params.tool_maybe_tweak_event))),
             *_template_view3d_gpencil_select_for_fallback(params, fallback),
         ]}
     )
@@ -7380,7 +7364,7 @@ def km_3d_view_tool_sculpt_gpencil_select_box(params):
     return (
         "3D View Tool: Sculpt Gpencil, Select Box",
         {"space_type": 'VIEW_3D', "region_type": 'WINDOW'},
-        {"items": _template_items_tool_select_actions("gpencil.select_box", **params.tool_tweak_event)},
+        {"items": _template_items_tool_select_actions("gpencil.select_box", **params.tool_maybe_tweak_event)},
     )
 
 
@@ -7399,7 +7383,7 @@ def km_3d_view_tool_sculpt_gpencil_select_lasso(params):
     return (
         "3D View Tool: Sculpt Gpencil, Select Lasso",
         {"space_type": 'VIEW_3D', "region_type": 'WINDOW'},
-        {"items": _template_items_tool_select_actions("gpencil.select_lasso", **params.tool_tweak_event)},
+        {"items": _template_items_tool_select_actions("gpencil.select_lasso", **params.tool_maybe_tweak_event)},
     )
 
 
@@ -7421,9 +7405,8 @@ def km_sequencer_editor_tool_select_box(params, *, fallback):
         _fallback_id("Sequencer Tool: Select Box", fallback),
         {"space_type": 'SEQUENCE_EDITOR', "region_type": 'WINDOW'},
         {"items": [
-            # Don't use `tool_maybe_tweak_event`, see comment for this slot.
             *_template_items_tool_select_actions_simple(
-                "sequencer.select_box", **params.tool_tweak_event,
+                "sequencer.select_box", **params.tool_maybe_tweak_event,
                 properties=[("tweak", params.select_mouse == 'LEFTMOUSE')],
             ),
             # RMB select can already set the frame, match the tweak tool.

release/scripts/presets/keyconfig/keymap_data/industry_compatible_data.py

@@ -3876,9 +3876,6 @@ def km_knife_tool_modal_map(_params):
         ("IGNORE_SNAP_OFF", {"type": 'LEFT_SHIFT', "value": 'RELEASE', "any": True}, None),
         ("IGNORE_SNAP_ON", {"type": 'RIGHT_SHIFT', "value": 'PRESS', "any": True}, None),
         ("IGNORE_SNAP_OFF", {"type": 'RIGHT_SHIFT', "value": 'RELEASE', "any": True}, None),
-        ("X_AXIS", {"type": 'X', "value": 'PRESS'}, None),
-        ("Y_AXIS", {"type": 'Y', "value": 'PRESS'}, None),
-        ("Z_AXIS", {"type": 'Z', "value": 'PRESS'}, None),
         ("ANGLE_SNAP_TOGGLE", {"type": 'A', "value": 'PRESS'}, None),
         ("CYCLE_ANGLE_SNAP_EDGE", {"type": 'R', "value": 'PRESS'}, None),
         ("CUT_THROUGH_TOGGLE", {"type": 'C', "value": 'PRESS'}, None),

release/scripts/startup/bl_ui/properties_data_mesh.py

@@ -639,6 +639,7 @@ class DATA_PT_mesh_attributes(MeshButtonsPanel, Panel):
 
         add_attributes(mesh.attributes)
         add_attributes(mesh.uv_layers)
+        add_attributes(mesh.vertex_colors)
         add_attributes(ob.vertex_groups)
 
         colliding_names = [name for name, layers in attributes_by_name.items() if len(layers) >= 2]

release/scripts/startup/bl_ui/space_filebrowser.py

@@ -648,6 +648,30 @@ class ASSETBROWSER_MT_select(AssetBrowserMenu, Menu):
         layout.operator("file.select_box")
 
 
+class ASSETBROWSER_PT_navigation_bar(asset_utils.AssetBrowserPanel, Panel):
+    bl_label = "Asset Navigation"
+    bl_region_type = 'TOOLS'
+    bl_options = {'HIDE_HEADER'}
+
+    @classmethod
+    def poll(cls, context):
+        return (
+            asset_utils.AssetBrowserPanel.poll(context) and
+            context.preferences.experimental.use_extended_asset_browser
+        )
+
+    def draw(self, context):
+        layout = self.layout
+
+        space_file = context.space_data
+
+        col = layout.column()
+
+        col.scale_x = 1.3
+        col.scale_y = 1.3
+        col.prop(space_file.params, "asset_category", expand=True)
+
+
 class ASSETBROWSER_PT_metadata(asset_utils.AssetBrowserPanel, Panel):
     bl_region_type = 'TOOL_PROPS'
     bl_label = "Asset Metadata"
@@ -762,10 +786,9 @@ class ASSETBROWSER_MT_context_menu(AssetBrowserMenu, Menu):
 
         layout.separator()
 
-        sub = layout.column()
+        sub = layout.row()
         sub.operator_context = 'EXEC_DEFAULT'
-        sub.operator("asset.clear", text="Clear Asset").set_fake_user = False
-        sub.operator("asset.clear", text="Clear Asset (Set Fake User)").set_fake_user = True
+        sub.operator("asset.clear", text="Clear Asset")
 
         layout.separator()
 
@@ -797,6 +820,7 @@ classes = (
     ASSETBROWSER_MT_editor_menus,
     ASSETBROWSER_MT_view,
     ASSETBROWSER_MT_select,
+    ASSETBROWSER_PT_navigation_bar,
     ASSETBROWSER_PT_metadata,
     ASSETBROWSER_PT_metadata_preview,
     ASSETBROWSER_PT_metadata_details,

release/scripts/startup/bl_ui/space_image.py

@@ -598,10 +598,16 @@ class IMAGE_HT_tool_header(Header):
     def draw(self, context):
         layout = self.layout
 
+        layout.template_header()
+
         self.draw_tool_settings(context)
 
         layout.separator_spacer()
 
+        IMAGE_HT_header.draw_xform_template(layout, context)
+
+        layout.separator_spacer()
+
         self.draw_mode_settings(context)
 
     def draw_tool_settings(self, context):
@@ -756,7 +762,8 @@ class IMAGE_HT_header(Header):
         show_uvedit = sima.show_uvedit
         show_maskedit = sima.show_maskedit
 
-        layout.template_header()
+        if not show_region_tool_header:
+            layout.template_header()
 
         if sima.mode != 'UV':
             layout.prop(sima, "ui_mode", text="")
@@ -777,7 +784,8 @@ class IMAGE_HT_header(Header):
 
         layout.separator_spacer()
 
-        IMAGE_HT_header.draw_xform_template(layout, context)
+        if not show_region_tool_header:
+            IMAGE_HT_header.draw_xform_template(layout, context)
 
         layout.template_ID(sima, "image", new="image.new", open="image.open")
 
@@ -926,10 +934,6 @@ class IMAGE_PT_snapping(Panel):
             row = col.row(align=True)
             row.prop(tool_settings, "snap_target", expand=True)
 
-        col.separator()
-        if 'INCREMENT' in tool_settings.snap_uv_element:
-            col.prop(tool_settings, "use_snap_uv_grid_absolute")
-
         col.label(text="Affect")
         row = col.row(align=True)
         row.prop(tool_settings, "use_snap_translate", text="Move", toggle=True)
@@ -1463,33 +1467,6 @@ class IMAGE_PT_udim_grid(Panel):
         col = layout.column()
         col.prop(uvedit, "tile_grid_shape", text="Grid Shape")
 
-class IMAGE_PT_custom_grid(Panel):
-    bl_space_type = 'IMAGE_EDITOR'
-    bl_region_type = 'UI'
-    bl_category = "View"
-    bl_label = "Custom Grid"
-
-    @classmethod
-    def poll(cls, context):
-        sima = context.space_data
-        return sima.show_uvedit
-
-    def draw_header(self, context):
-        sima = context.space_data
-        uvedit = sima.uv_editor
-        self.layout.prop(uvedit, "use_custom_grid", text="")
-
-    def draw(self, context):
-        layout = self.layout
-
-        sima = context.space_data
-        uvedit = sima.uv_editor
-
-        layout.use_property_split = True
-        layout.use_property_decorate = False
-
-        col = layout.column()
-        col.prop(uvedit, "custom_grid_subdivisions", text="Subdivisions")
 
 class IMAGE_PT_overlay(Panel):
     bl_space_type = 'IMAGE_EDITOR'
@@ -1675,7 +1652,6 @@ classes = (
     IMAGE_PT_uv_cursor,
     IMAGE_PT_annotation,
     IMAGE_PT_udim_grid,
-    IMAGE_PT_custom_grid,
     IMAGE_PT_overlay,
     IMAGE_PT_overlay_uv_edit,
     IMAGE_PT_overlay_uv_edit_geometry,

release/scripts/startup/bl_ui/space_node.py

@@ -754,11 +754,6 @@ class NodeTreeInterfacePanel:
             # Display descriptions only for Geometry Nodes, since it's only used in the modifier panel.
             if tree.type == 'GEOMETRY':
                 layout.prop(active_socket, "description")
-                field_socket_prefixes = {
-                    "NodeSocketInt", "NodeSocketColor", "NodeSocketVector", "NodeSocketBool", "NodeSocketFloat"}
-                is_field_type = any(active_socket.bl_socket_idname.startswith(prefix) for prefix in field_socket_prefixes)
-                if in_out == "OUT" and is_field_type:
-                    layout.prop(active_socket, "attribute_domain")
             active_socket.draw(context, layout)
 
 

release/scripts/startup/bl_ui/space_outliner.py

@@ -323,8 +323,7 @@ class OUTLINER_MT_asset(Menu):
         space = context.space_data
 
         layout.operator("asset.mark")
-        layout.operator("asset.clear", text="Clear Asset").set_fake_user = False
-        layout.operator("asset.clear", text="Clear Asset (Set Fake User)").set_fake_user = True
+        layout.operator("asset.clear")
 
 
 class OUTLINER_PT_filter(Panel):

release/scripts/startup/bl_ui/space_sequencer.py

@@ -99,6 +99,8 @@ class SEQUENCER_HT_tool_header(Header):
     def draw(self, context):
         layout = self.layout
 
+        layout.template_header()
+
         self.draw_tool_settings(context)
 
         # TODO: options popover.
@@ -130,7 +132,8 @@ class SEQUENCER_HT_header(Header):
 
         show_region_tool_header = st.show_region_tool_header
 
-        layout.template_header()
+        if not show_region_tool_header:
+            layout.template_header()
 
         layout.prop(st, "view_type", text="")
 
@@ -149,8 +152,6 @@ class SEQUENCER_HT_header(Header):
             row = layout.row(align=True)
             row.prop(sequencer_tool_settings, "overlap_mode", text="")
             row = layout.row(align=True)
-            row.prop(sequencer_tool_settings, "pivot_point", text="", icon_only=True)
-            row = layout.row(align=True)
             row.prop(tool_settings, "use_snap_sequencer", text="")
             sub = row.row(align=True)
             sub.popover(panel="SEQUENCER_PT_snapping")
@@ -241,7 +242,6 @@ class SEQUENCER_PT_sequencer_overlay(Panel):
         layout.prop(overlay_settings, "show_strip_name", text="Name")
         layout.prop(overlay_settings, "show_strip_source", text="Source")
         layout.prop(overlay_settings, "show_strip_duration", text="Duration")
-        layout.prop(overlay_settings, "show_strip_tag_color", text="Color Tags")
 
         layout.separator()
 
@@ -375,9 +375,9 @@ class SEQUENCER_MT_view(Menu):
                 layout.operator("view2d.zoom_border", text="Zoom")
                 layout.menu("SEQUENCER_MT_preview_zoom")
 
-            layout.prop(st, "use_zoom_to_fit")
-
-            if st.display_mode == 'WAVEFORM':
+            if st.display_mode == 'IMAGE':
+                layout.prop(st, "use_zoom_to_fit")
+            elif st.display_mode == 'WAVEFORM':
                 layout.separator()
                 layout.prop(st, "show_separate_color", text="Show Separate Color Channels")
 
@@ -869,9 +869,6 @@ class SEQUENCER_MT_strip(Menu):
                 layout.operator("sequencer.meta_make")
                 layout.operator("sequencer.meta_toggle", text="Toggle Meta")
 
-        layout.separator()
-        layout.menu("SEQUENCER_MT_color_tag_picker")
-
         layout.separator()
         layout.menu("SEQUENCER_MT_strip_lock_mute")
 
@@ -968,9 +965,6 @@ class SEQUENCER_MT_context_menu(Menu):
                 layout.operator("sequencer.meta_make")
                 layout.operator("sequencer.meta_toggle", text="Toggle Meta")
 
-        layout.separator()
-        layout.menu("SEQUENCER_MT_color_tag_picker")
-
         layout.separator()
 
         layout.menu("SEQUENCER_MT_strip_lock_mute")
@@ -1003,41 +997,6 @@ class SequencerButtonsPanel_Output:
         return cls.has_preview(context)
 
 
-class SEQUENCER_PT_color_tag_picker(Panel):
-    bl_label = "Color Tag"
-    bl_space_type = 'SEQUENCE_EDITOR'
-    bl_region_type = 'UI'
-    bl_category = "Strip"
-    bl_options = {'HIDE_HEADER', 'INSTANCED'}
-
-    @classmethod
-    def poll(cls, context):
-        return context.active_sequence_strip is not None
-
-    def draw(self, context):
-        layout = self.layout
-
-        row = layout.row(align=True)
-        row.operator("sequencer.strip_color_tag_set", icon="X").color = 'NONE'
-        for i in range(1, 10):
-            icon = 'SEQUENCE_COLOR_%02d' % i
-            row.operator("sequencer.strip_color_tag_set", icon=icon).color = 'COLOR_%02d' % i
-
-
-class SEQUENCER_MT_color_tag_picker(Menu):
-    bl_label = "Set Color Tag"
-
-    @classmethod
-    def poll(cls, context):
-        return context.active_sequence_strip is not None
-
-    def draw(self, context):
-        layout = self.layout
-
-        row = layout.row(align=True)
-        row.operator_enum("sequencer.strip_color_tag_set", "color", icon_only=True)
-
-
 class SEQUENCER_PT_strip(SequencerButtonsPanel, Panel):
     bl_label = ""
     bl_options = {'HIDE_HEADER'}
@@ -1081,20 +1040,9 @@ class SEQUENCER_PT_strip(SequencerButtonsPanel, Panel):
         else:
             icon_header = 'SEQ_SEQUENCER'
 
-        row = layout.row(align=True)
-        row.use_property_decorate = False
+        row = layout.row()
         row.label(text="", icon=icon_header)
-        row.separator()
         row.prop(strip, "name", text="")
-
-        sub = row.row(align=True)
-        if strip.color_tag == 'NONE':
-            sub.popover(panel="SEQUENCER_PT_color_tag_picker", text="", icon='COLOR')
-        else:
-            icon = 'SEQUENCE_' + strip.color_tag
-            sub.popover(panel="SEQUENCER_PT_color_tag_picker", text="", icon=icon)
-
-        row.separator()
         row.prop(strip, "mute", toggle=True, icon_only=True, emboss=False)
 
 
@@ -2380,11 +2328,8 @@ classes = (
     SEQUENCER_MT_strip_transform,
     SEQUENCER_MT_strip_input,
     SEQUENCER_MT_strip_lock_mute,
-    SEQUENCER_MT_color_tag_picker,
     SEQUENCER_MT_context_menu,
 
-    SEQUENCER_PT_color_tag_picker,
-
     SEQUENCER_PT_active_tool,
     SEQUENCER_PT_strip,
 

release/scripts/startup/bl_ui/space_toolsystem_common.py

@@ -190,7 +190,7 @@ class ToolActivePanelHelper:
         ToolSelectPanelHelper.draw_active_tool_header(
             context,
             layout.column(),
-            show_tool_icon=True,
+            show_tool_name=True,
             tool_key=ToolSelectPanelHelper._tool_key_from_context(context, space_type=self.bl_space_type),
         )
 
@@ -766,7 +766,7 @@ class ToolSelectPanelHelper:
     def draw_active_tool_header(
             context, layout,
             *,
-            show_tool_icon=False,
+            show_tool_name=False,
             tool_key=None,
     ):
         if tool_key is None:
@@ -783,12 +783,9 @@ class ToolSelectPanelHelper:
             return None
         # Note: we could show 'item.text' here but it makes the layout jitter when switching tools.
         # Add some spacing since the icon is currently assuming regular small icon size.
-        if show_tool_icon:
-            layout.label(text="    " + item.label, icon_value=icon_value)
+        layout.label(text="    " + item.label if show_tool_name else " ", icon_value=icon_value)
+        if show_tool_name:
             layout.separator()
-        else:
-            layout.label(text=item.label)
-
         draw_settings = item.draw_settings
         if draw_settings is not None:
             draw_settings(context, layout, tool)

release/scripts/startup/bl_ui/space_toolsystem_toolbar.py

@@ -497,14 +497,18 @@ class _defs_view3d_add:
         props = tool.operator_properties("view3d.interactive_add")
         if not extra:
             row = layout.row()
+            row.scale_x = 0.8
             row.label(text="Depth:")
             row = layout.row()
+            row.scale_x = 0.9
             row.prop(props, "plane_depth", text="")
             row = layout.row()
+            row.scale_x = 0.8
             row.label(text="Orientation:")
             row = layout.row()
             row.prop(props, "plane_orientation", text="")
             row = layout.row()
+            row.scale_x = 0.8
             row.prop(props, "snap_target")
 
             region_is_header = bpy.context.region.type == 'TOOL_HEADER'

release/scripts/startup/bl_ui/space_userpref.py

@@ -1073,25 +1073,6 @@ class USERPREF_PT_theme_collection_colors(ThemePanel, CenterAlignMixIn, Panel):
             flow.prop(ui, "color", text=iface_("Color %d") % i, translate=False)
 
 
-class USERPREF_PT_theme_strip_colors(ThemePanel, CenterAlignMixIn, Panel):
-    bl_label = "Strip Colors"
-    bl_options = {'DEFAULT_CLOSED'}
-
-    def draw_header(self, _context):
-        layout = self.layout
-
-        layout.label(icon='SEQ_STRIP_DUPLICATE')
-
-    def draw_centered(self, context, layout):
-        theme = context.preferences.themes[0]
-
-        layout.use_property_split = True
-
-        flow = layout.grid_flow(row_major=False, columns=0, even_columns=True, even_rows=False, align=False)
-        for i, ui in enumerate(theme.strip_color, 1):
-            flow.prop(ui, "color", text=iface_("Color %d") % i, translate=False)
-
-
 # Base class for dynamically defined theme-space panels.
 # This is not registered.
 class PreferenceThemeSpacePanel:
@@ -2270,6 +2251,7 @@ class USERPREF_PT_experimental_new_features(ExperimentalPanel, Panel):
                 ({"property": "use_sculpt_tools_tilt"}, "T82877"),
                 ({"property": "use_extended_asset_browser"}, ("project/view/130/", "Project Page")),
                 ({"property": "use_override_templates"}, ("T73318", "Milestone 4")),
+                ({"property": "use_geometry_nodes_fields"}, "T91274"),
             ),
         )
 
@@ -2303,7 +2285,6 @@ class USERPREF_PT_experimental_debugging(ExperimentalPanel, Panel):
                 ({"property": "override_auto_resync"}, "T83811"),
                 ({"property": "proxy_to_override_auto_conversion"}, "T91671"),
                 ({"property": "use_cycles_debug"}, None),
-                ({"property": "use_geometry_nodes_legacy"}, "T91274"),
             ),
         )
 
@@ -2367,7 +2348,6 @@ classes = (
     USERPREF_PT_theme_text_style,
     USERPREF_PT_theme_bone_color_sets,
     USERPREF_PT_theme_collection_colors,
-    USERPREF_PT_theme_strip_colors,
 
     USERPREF_PT_file_paths_data,
     USERPREF_PT_file_paths_render,

release/scripts/startup/bl_ui/space_view3d.py

@@ -46,10 +46,16 @@ class VIEW3D_HT_tool_header(Header):
     def draw(self, context):
         layout = self.layout
 
+        layout.row(align=True).template_header()
+
         self.draw_tool_settings(context)
 
         layout.separator_spacer()
 
+        VIEW3D_HT_header.draw_xform_template(layout, context)
+
+        layout.separator_spacer()
+
         self.draw_mode_settings(context)
 
     def draw_tool_settings(self, context):
@@ -598,8 +604,10 @@ class VIEW3D_HT_header(Header):
         tool_settings = context.tool_settings
         view = context.space_data
         shading = view.shading
+        show_region_tool_header = view.show_region_tool_header
 
-        layout.row(align=True).template_header()
+        if not show_region_tool_header:
+            layout.row(align=True).template_header()
 
         row = layout.row(align=True)
         obj = context.active_object
@@ -746,7 +754,7 @@ class VIEW3D_HT_header(Header):
                     )
 
             layout.separator_spacer()
-        else:
+        elif not show_region_tool_header:
             # Transform settings depending on tool header visibility
             VIEW3D_HT_header.draw_xform_template(layout, context)
 

release/scripts/startup/nodeitems_builtins.py

@@ -180,8 +180,11 @@ def object_eevee_cycles_shader_nodes_poll(context):
             eevee_cycles_shader_nodes_poll(context))
 
 
-def geometry_nodes_legacy_poll(context):
-    return context.preferences.experimental.use_geometry_nodes_legacy
+def geometry_nodes_fields_poll(context):
+    return context.preferences.experimental.use_geometry_nodes_fields
+
+def geometry_nodes_fields_legacy_poll(context):
+    return not context.preferences.experimental.use_geometry_nodes_fields
 
 
 # All standard node categories currently used in nodes.
@@ -480,27 +483,27 @@ texture_node_categories = [
 geometry_node_categories = [
     # Geometry Nodes
     GeometryNodeCategory("GEO_ATTRIBUTE", "Attribute", items=[
-        NodeItem("GeometryNodeLegacyAttributeRandomize", poll=geometry_nodes_legacy_poll),
-        NodeItem("GeometryNodeLegacyAttributeMath", poll=geometry_nodes_legacy_poll),
-        NodeItem("GeometryNodeLegacyAttributeClamp", poll=geometry_nodes_legacy_poll),
-        NodeItem("GeometryNodeLegacyAttributeCompare", poll=geometry_nodes_legacy_poll),
-        NodeItem("GeometryNodeLegacyAttributeConvert", poll=geometry_nodes_legacy_poll),
-        NodeItem("GeometryNodeLegacyAttributeCurveMap", poll=geometry_nodes_legacy_poll),
-        NodeItem("GeometryNodeLegacyAttributeFill", poll=geometry_nodes_legacy_poll),
-        NodeItem("GeometryNodeLegacyAttributeMix", poll=geometry_nodes_legacy_poll),
-        NodeItem("GeometryNodeLegacyAttributeProximity", poll=geometry_nodes_legacy_poll),
-        NodeItem("GeometryNodeLegacyAttributeColorRamp", poll=geometry_nodes_legacy_poll),
-        NodeItem("GeometryNodeLegacyAttributeVectorMath", poll=geometry_nodes_legacy_poll),
-        NodeItem("GeometryNodeLegacyAttributeVectorRotate", poll=geometry_nodes_legacy_poll),
-        NodeItem("GeometryNodeLegacyAttributeSampleTexture", poll=geometry_nodes_legacy_poll),
-        NodeItem("GeometryNodeLegacyAttributeCombineXYZ", poll=geometry_nodes_legacy_poll),
-        NodeItem("GeometryNodeLegacyAttributeSeparateXYZ", poll=geometry_nodes_legacy_poll),
-        NodeItem("GeometryNodeLegacyAttributeMapRange", poll=geometry_nodes_legacy_poll),
-        NodeItem("GeometryNodeLegacyAttributeTransfer", poll=geometry_nodes_legacy_poll),
-        NodeItem("GeometryNodeAttributeRemove", poll=geometry_nodes_legacy_poll),
+        NodeItem("GeometryNodeLegacyAttributeRandomize", poll=geometry_nodes_fields_legacy_poll),
+        NodeItem("GeometryNodeLegacyAttributeMath", poll=geometry_nodes_fields_legacy_poll),
+        NodeItem("GeometryNodeLegacyAttributeClamp", poll=geometry_nodes_fields_legacy_poll),
+        NodeItem("GeometryNodeLegacyAttributeCompare", poll=geometry_nodes_fields_legacy_poll),
+        NodeItem("GeometryNodeLegacyAttributeConvert", poll=geometry_nodes_fields_legacy_poll),
+        NodeItem("GeometryNodeLegacyAttributeCurveMap", poll=geometry_nodes_fields_legacy_poll),
+        NodeItem("GeometryNodeLegacyAttributeFill", poll=geometry_nodes_fields_legacy_poll),
+        NodeItem("GeometryNodeLegacyAttributeMix", poll=geometry_nodes_fields_legacy_poll),
+        NodeItem("GeometryNodeLegacyAttributeProximity", poll=geometry_nodes_fields_legacy_poll),
+        NodeItem("GeometryNodeLegacyAttributeColorRamp", poll=geometry_nodes_fields_legacy_poll),
+        NodeItem("GeometryNodeLegacyAttributeVectorMath", poll=geometry_nodes_fields_legacy_poll),
+        NodeItem("GeometryNodeLegacyAttributeVectorRotate", poll=geometry_nodes_fields_legacy_poll),
+        NodeItem("GeometryNodeLegacyAttributeSampleTexture", poll=geometry_nodes_fields_legacy_poll),
+        NodeItem("GeometryNodeLegacyAttributeCombineXYZ", poll=geometry_nodes_fields_legacy_poll),
+        NodeItem("GeometryNodeLegacyAttributeSeparateXYZ", poll=geometry_nodes_fields_legacy_poll),
+        NodeItem("GeometryNodeLegacyAttributeMapRange", poll=geometry_nodes_fields_legacy_poll),
+        NodeItem("GeometryNodeLegacyAttributeTransfer", poll=geometry_nodes_fields_legacy_poll),
+        NodeItem("GeometryNodeAttributeRemove", poll=geometry_nodes_fields_legacy_poll),
 
-        NodeItem("GeometryNodeAttributeCapture"),
-        NodeItem("GeometryNodeAttributeStatistic"),
+        NodeItem("GeometryNodeAttributeCapture", poll=geometry_nodes_fields_poll),
+        NodeItem("GeometryNodeAttributeStatistic", poll=geometry_nodes_fields_poll),
     ]),
     GeometryNodeCategory("GEO_COLOR", "Color", items=[
         NodeItem("ShaderNodeMixRGB"),
@@ -510,25 +513,24 @@ geometry_node_categories = [
         NodeItem("ShaderNodeCombineRGB"),
     ]),
     GeometryNodeCategory("GEO_CURVE", "Curve", items=[
-        NodeItem("GeometryNodeLegacyCurveSubdivide", poll=geometry_nodes_legacy_poll),
-        NodeItem("GeometryNodeLegacyCurveReverse", poll=geometry_nodes_legacy_poll),
-        NodeItem("GeometryNodeLegacyCurveSplineType", poll=geometry_nodes_legacy_poll),
-        NodeItem("GeometryNodeLegacyCurveSetHandles", poll=geometry_nodes_legacy_poll),
-        NodeItem("GeometryNodeLegacyCurveSelectHandles", poll=geometry_nodes_legacy_poll),
-        NodeItem("GeometryNodeLegacyMeshToCurve", poll=geometry_nodes_legacy_poll),
-        NodeItem("GeometryNodeLegacyCurveToPoints", poll=geometry_nodes_legacy_poll),
-        NodeItem("GeometryNodeLegacyCurveEndpoints", poll=geometry_nodes_legacy_poll),
+        NodeItem("GeometryNodeLegacyCurveSubdivide", poll=geometry_nodes_fields_legacy_poll),
+        NodeItem("GeometryNodeLegacyCurveReverse", poll=geometry_nodes_fields_legacy_poll),
+        NodeItem("GeometryNodeLegacyCurveSplineType", poll=geometry_nodes_fields_legacy_poll),
+        NodeItem("GeometryNodeLegacyCurveSetHandles", poll=geometry_nodes_fields_legacy_poll),
+        NodeItem("GeometryNodeLegacyCurveSelectHandles", poll=geometry_nodes_fields_legacy_poll),
+        NodeItem("GeometryNodeLegacyMeshToCurve", poll=geometry_nodes_fields_legacy_poll),
 
         NodeItem("GeometryNodeCurveToMesh"),
         NodeItem("GeometryNodeCurveResample"),
+        NodeItem("GeometryNodeCurveToPoints"),
+        NodeItem("GeometryNodeCurveEndpoints"),
         NodeItem("GeometryNodeCurveFill"),
         NodeItem("GeometryNodeCurveTrim"),
         NodeItem("GeometryNodeCurveLength"),
-        NodeItem("GeometryNodeCurveParameter"),
-        NodeItem("GeometryNodeInputTangent"),
-        NodeItem("GeometryNodeCurveSample"),
-        NodeItem("GeometryNodeCurveFillet"),
-        NodeItem("GeometryNodeCurveReverse"),
+        NodeItem("GeometryNodeCurveParameter", poll=geometry_nodes_fields_poll),
+        NodeItem("GeometryNodeInputTangent", poll=geometry_nodes_fields_poll),
+        NodeItem("GeometryNodeCurveSample", poll=geometry_nodes_fields_poll),
+        NodeItem("GeometryNodeCurveFillet", poll=geometry_nodes_fields_poll),
     ]),
     GeometryNodeCategory("GEO_PRIMITIVES_CURVE", "Curve Primitives", items=[
         NodeItem("GeometryNodeCurvePrimitiveLine"),
@@ -540,21 +542,20 @@ geometry_node_categories = [
         NodeItem("GeometryNodeCurvePrimitiveBezierSegment"),
     ]),
     GeometryNodeCategory("GEO_GEOMETRY", "Geometry", items=[
-        NodeItem("GeometryNodeLegacyDeleteGeometry", poll=geometry_nodes_legacy_poll),
-        NodeItem("GeometryNodeLegacyRaycast", poll=geometry_nodes_legacy_poll),
+        NodeItem("GeometryNodeLegacyDeleteGeometry", poll=geometry_nodes_fields_legacy_poll),
+        NodeItem("GeometryNodeLegacyRaycast", poll=geometry_nodes_fields_legacy_poll),
 
-        NodeItem("GeometryNodeProximity"),
         NodeItem("GeometryNodeBoundBox"),
         NodeItem("GeometryNodeConvexHull"),
         NodeItem("GeometryNodeTransform"),
         NodeItem("GeometryNodeJoinGeometry"),
         NodeItem("GeometryNodeSeparateComponents"),
-        NodeItem("GeometryNodeSetPosition"),
-        NodeItem("GeometryNodeRealizeInstances"),
+        NodeItem("GeometryNodeSetPosition", poll=geometry_nodes_fields_poll),
+        NodeItem("GeometryNodeRealizeInstances", poll=geometry_nodes_fields_poll),
     ]),
     GeometryNodeCategory("GEO_INPUT", "Input", items=[
-        NodeItem("FunctionNodeLegacyRandomFloat", poll=geometry_nodes_legacy_poll),
-
+        NodeItem("FunctionNodeLegacyRandomFloat", poll=geometry_nodes_fields_legacy_poll),
+        
         NodeItem("GeometryNodeObjectInfo"),
         NodeItem("GeometryNodeCollectionInfo"),
         NodeItem("ShaderNodeValue"),
@@ -562,26 +563,24 @@ geometry_node_categories = [
         NodeItem("FunctionNodeInputVector"),
         NodeItem("GeometryNodeInputMaterial"),
         NodeItem("GeometryNodeIsViewport"),
-        NodeItem("GeometryNodeInputPosition"),
-        NodeItem("GeometryNodeInputIndex"),
-        NodeItem("GeometryNodeInputNormal"),
+        NodeItem("GeometryNodeInputPosition", poll=geometry_nodes_fields_poll),
+        NodeItem("GeometryNodeInputIndex", poll=geometry_nodes_fields_poll),
+        NodeItem("GeometryNodeInputNormal", poll=geometry_nodes_fields_poll),
     ]),
     GeometryNodeCategory("GEO_MATERIAL", "Material", items=[
-        NodeItem("GeometryNodeLegacyMaterialAssign", poll=geometry_nodes_legacy_poll),
-        NodeItem("GeometryNodeLegacySelectByMaterial", poll=geometry_nodes_legacy_poll),
+        NodeItem("GeometryNodeLegacyMaterialAssign", poll=geometry_nodes_fields_legacy_poll),
+        NodeItem("GeometryNodeLegacySelectByMaterial", poll=geometry_nodes_fields_legacy_poll),
 
-        NodeItem("GeometryNodeMaterialAssign"),
-        NodeItem("GeometryNodeMaterialSelection"),
+        NodeItem("GeometryNodeMaterialAssign", poll=geometry_nodes_fields_poll),
+        NodeItem("GeometryNodeMaterialSelection", poll=geometry_nodes_fields_poll),
         NodeItem("GeometryNodeMaterialReplace"),
     ]),
     GeometryNodeCategory("GEO_MESH", "Mesh", items=[
-        NodeItem("GeometryNodeLegacyEdgeSplit", poll=geometry_nodes_legacy_poll),
-        NodeItem("GeometryNodeLegacySubdivisionSurface", poll=geometry_nodes_legacy_poll),
-
         NodeItem("GeometryNodeBoolean"),
         NodeItem("GeometryNodeTriangulate"),
+        NodeItem("GeometryNodeEdgeSplit"),
+        NodeItem("GeometryNodeSubdivisionSurface"),
         NodeItem("GeometryNodeMeshSubdivide"),
-        NodeItem("GeometryNodePointsToVertices"),
     ]),
     GeometryNodeCategory("GEO_PRIMITIVES_MESH", "Mesh Primitives", items=[
         NodeItem("GeometryNodeMeshCircle"),
@@ -594,16 +593,14 @@ geometry_node_categories = [
         NodeItem("GeometryNodeMeshUVSphere"),
     ]),
     GeometryNodeCategory("GEO_POINT", "Point", items=[
-        NodeItem("GeometryNodeMeshToPoints"),
-        NodeItem("GeometryNodeInstanceOnPoints"),
-        NodeItem("GeometryNodeDistributePointsOnFaces"),
-        NodeItem("GeometryNodeLegacyPointDistribute", poll=geometry_nodes_legacy_poll),
-        NodeItem("GeometryNodeLegacyPointInstance", poll=geometry_nodes_legacy_poll),
-        NodeItem("GeometryNodeLegacyPointSeparate", poll=geometry_nodes_legacy_poll),
-        NodeItem("GeometryNodeLegacyPointScale", poll=geometry_nodes_legacy_poll),
-        NodeItem("GeometryNodeLegacyPointTranslate", poll=geometry_nodes_legacy_poll),
-        NodeItem("GeometryNodeLegacyRotatePoints", poll=geometry_nodes_legacy_poll),
-        NodeItem("GeometryNodeLegacyAlignRotationToVector", poll=geometry_nodes_legacy_poll),
+        NodeItem("GeometryNodeDistributePointsOnFaces", poll=geometry_nodes_fields_poll),
+        NodeItem("GeometryNodeLegacyPointDistribute", poll=geometry_nodes_fields_legacy_poll),
+        NodeItem("GeometryNodeLegacyPointInstance", poll=geometry_nodes_fields_legacy_poll),
+        NodeItem("GeometryNodeLegacyPointSeparate", poll=geometry_nodes_fields_legacy_poll),
+        NodeItem("GeometryNodeLegacyPointScale", poll=geometry_nodes_fields_legacy_poll),
+        NodeItem("GeometryNodeLegacyPointTranslate", poll=geometry_nodes_fields_legacy_poll),
+        NodeItem("GeometryNodeLegacyRotatePoints", poll=geometry_nodes_fields_legacy_poll),
+        NodeItem("GeometryNodeLegacyAlignRotationToVector", poll=geometry_nodes_fields_legacy_poll),
     ]),
     GeometryNodeCategory("GEO_TEXT", "Text", items=[
         NodeItem("FunctionNodeStringLength"),
@@ -621,10 +618,10 @@ geometry_node_categories = [
         NodeItem("FunctionNodeFloatCompare"),
         NodeItem("FunctionNodeFloatToInt"),
         NodeItem("GeometryNodeSwitch"),
-        NodeItem("FunctionNodeRandomValue"),
+        NodeItem("FunctionNodeRandomValue", poll=geometry_nodes_fields_poll),
     ]),
     GeometryNodeCategory("GEO_TEXTURE", "Texture", items=[
-        NodeItem("ShaderNodeTexNoise"),
+        NodeItem("ShaderNodeTexNoise", poll=geometry_nodes_fields_poll),
     ]),
     GeometryNodeCategory("GEO_VECTOR", "Vector", items=[
         NodeItem("ShaderNodeVectorCurve"),
@@ -637,7 +634,7 @@ geometry_node_categories = [
         NodeItem("GeometryNodeViewer"),
     ]),
     GeometryNodeCategory("GEO_VOLUME", "Volume", items=[
-        NodeItem("GeometryNodeLegacyPointsToVolume", poll=geometry_nodes_legacy_poll),
+        NodeItem("GeometryNodeLegacyPointsToVolume", poll=geometry_nodes_fields_legacy_poll),
 
         NodeItem("GeometryNodeVolumeToMesh"),
     ]),

source/blender/blenkernel/BKE_asset_catalog.hh

@@ -65,23 +65,13 @@ class AssetCatalogService {
   void load_from_disk(const CatalogFilePath &file_or_directory_path);
 
   /**
-   * Write the catalog definitions to disk in response to the blend file being saved.
+   * Write the catalog definitions to disk.
+   * The provided directory path is only used when there is no CDF loaded from disk yet but assets
+   * still have to be saved.
    *
-   * The location where the catalogs are saved is variable, and depends on the location of the
-   * blend file. The first matching rule wins:
-   *
-   * - Already loaded a CDF from disk?
-   *    -> Always write to that file.
-   * - The directory containing the blend file has a blender_assets.cats.txt file?
-   *    -> Merge with & write to that file.
-   * - The directory containing the blend file is part of an asset library, as per
-   *   the user's preferences?
-   *    -> Merge with & write to ${ASSET_LIBRARY_ROOT}/blender_assets.cats.txt
-   * - Create a new file blender_assets.cats.txt next to the blend file.
-   *
-   * Return true on success, which either means there were no in-memory categories to save,
-   * or the save was successful. */
-  bool write_to_disk_on_blendfile_save(const CatalogFilePath &blend_file_path);
+   * Return true on success, which either means there were no in-memory categories to save, or the
+   * save was successful. */
+  bool write_to_disk(const CatalogFilePath &directory_for_new_files);
 
   /**
    * Merge on-disk changes into the in-memory asset catalogs.
@@ -96,10 +86,6 @@ class AssetCatalogService {
   /** Return catalog with the given ID. Return nullptr if not found. */
   AssetCatalog *find_catalog(CatalogID catalog_id);
 
-  /** Return first catalog with the given path. Return nullptr if not found. This is not an
-   * efficient call as it's just a linear search over the catalogs. */
-  AssetCatalog *find_catalog_by_path(const CatalogPath &path) const;
-
   /** Create a catalog with some sensible auto-generated catalog ID.
    * The catalog will be saved to the default catalog file.*/
   AssetCatalog *create_catalog(const CatalogPath &catalog_path);
@@ -109,11 +95,6 @@ class AssetCatalogService {
    * written. */
   void delete_catalog(CatalogID catalog_id);
 
-  /**
-   * Update the catalog path, also updating the catalog path of all sub-catalogs.
-   */
-  void update_catalog_path(CatalogID catalog_id, const CatalogPath &new_catalog_path);
-
   AssetCatalogTree *get_catalog_tree();
 
   /** Return true only if there are no catalogs known. */
@@ -139,87 +120,52 @@ class AssetCatalogService {
   std::unique_ptr<AssetCatalogDefinitionFile> construct_cdf_in_memory(
       const CatalogFilePath &file_path);
 
-  /**
-   * Find a suitable path to write a CDF to.
-   *
-   * This depends on the location of the blend file, and on whether a CDF already exists next to it
-   * or whether the blend file is saved inside an asset library.
-   */
-  static CatalogFilePath find_suitable_cdf_path_for_writing(
-      const CatalogFilePath &blend_file_path);
-
   std::unique_ptr<AssetCatalogTree> read_into_tree();
   void rebuild_tree();
 };
 
-/**
- * Representation of a catalog path in the #AssetCatalogTree.
- */
 class AssetCatalogTreeItem {
-  friend class AssetCatalogTree;
+  friend class AssetCatalogService;
 
  public:
-  /** Container for child items. Uses a #std::map to keep items ordered by their name (i.e. their
-   * last catalog component). */
   using ChildMap = std::map<std::string, AssetCatalogTreeItem>;
-  using ItemIterFn = FunctionRef<void(AssetCatalogTreeItem &)>;
+  using ItemIterFn = FunctionRef<void(const AssetCatalogTreeItem &)>;
 
-  AssetCatalogTreeItem(StringRef name,
-                       CatalogID catalog_id,
-                       const AssetCatalogTreeItem *parent = nullptr);
+  AssetCatalogTreeItem(StringRef name, const AssetCatalogTreeItem *parent = nullptr);
 
-  CatalogID get_catalog_id() const;
   StringRef get_name() const;
   /** Return the full catalog path, defined as the name of this catalog prefixed by the full
    * catalog path of its parent and a separator. */
   CatalogPath catalog_path() const;
   int count_parents() const;
-  bool has_children() const;
 
-  /** Iterate over children calling \a callback for each of them, but do not recurse into their
-   * children. */
-  void foreach_child(const ItemIterFn callback);
+  static void foreach_item_recursive(const ChildMap &children_, const ItemIterFn callback);
 
  protected:
   /** Child tree items, ordered by their names. */
   ChildMap children_;
   /** The user visible name of this component. */
   CatalogPathComponent name_;
-  CatalogID catalog_id_;
 
   /** Pointer back to the parent item. Used to reconstruct the hierarchy from an item (e.g. to
    * build a path). */
   const AssetCatalogTreeItem *parent_ = nullptr;
-
- private:
-  static void foreach_item_recursive(ChildMap &children_, ItemIterFn callback);
 };
 
 /**
  * A representation of the catalog paths as tree structure. Each component of the catalog tree is
- * represented by an #AssetCatalogTreeItem. The last path component of an item is used as its name,
- * which may also be shown to the user.
- * An item can not have multiple children with the same name. That means the name uniquely
- * identifies an item within its parent.
- *
+ * represented by a #AssetCatalogTreeItem.
  * There is no single root tree element, the #AssetCatalogTree instance itself represents the root.
  */
 class AssetCatalogTree {
-  using ChildMap = AssetCatalogTreeItem::ChildMap;
-  using ItemIterFn = AssetCatalogTreeItem::ItemIterFn;
+  friend class AssetCatalogService;
 
  public:
-  /** Ensure an item representing \a path is in the tree, adding it if necessary. */
-  void insert_item(const AssetCatalog &catalog);
-
-  void foreach_item(const AssetCatalogTreeItem::ItemIterFn callback);
-  /** Iterate over root items calling \a callback for each of them, but do not recurse into their
-   * children. */
-  void foreach_root_item(const ItemIterFn callback);
+  void foreach_item(const AssetCatalogTreeItem::ItemIterFn callback) const;
 
  protected:
   /** Child tree items, ordered by their names. */
-  ChildMap root_items_;
+  AssetCatalogTreeItem::ChildMap children_;
 };
 
 /** Keeps track of which catalogs are defined in a certain file on disk.
@@ -231,7 +177,6 @@ class AssetCatalogDefinitionFile {
    * Later versioning code may be added to handle older files. */
   const static int SUPPORTED_VERSION;
   const static std::string VERSION_MARKER;
-  const static std::string HEADER;
 
   CatalogFilePath file_path;
 
@@ -297,15 +242,6 @@ class AssetCatalog {
     bool is_deleted = false;
   } flags;
 
-  /**
-   * \return true only if this catalog's path is contained within the given path.
-   * When this catalog's path is equal to the given path, return true as well.
-   *
-   * Note that non-normalized paths (so for example starting or ending with a slash) are not
-   * supported, and result in undefined behavior.
-   */
-  bool is_contained_in(const CatalogPath &other_path) const;
-
   /**
    * Create a new Catalog with the given path, auto-generating a sensible catalog simple-name.
    *

source/blender/blenkernel/BKE_asset_library.h

@@ -20,8 +20,6 @@
 
 #pragma once
 
-struct Main;
-
 #ifdef __cplusplus
 extern "C" {
 #endif
@@ -33,41 +31,6 @@ typedef struct AssetLibrary AssetLibrary;
 struct AssetLibrary *BKE_asset_library_load(const char *library_path);
 void BKE_asset_library_free(struct AssetLibrary *asset_library);
 
-/**
- * Try to find an appropriate location for an asset library root from a file or directory path.
- * Does not check if \a input_path exists.
- *
- * The design is made to find an appropriate asset library path from a .blend file path, but
- * technically works with any file or directory as \a input_path.
- * Design is:
- * * If \a input_path lies within a known asset library path (i.e. an asset library registered in
- *   the Preferences), return the asset library path.
- * * Otherwise, if \a input_path has a parent path, return the parent path (e.g. to use the
- *   directory a .blend file is in as asset library root).
- * * If \a input_path is empty or doesn't have a parent path (e.g. because a .blend wasn't saved
- *   yet), there is no suitable path. The caller has to decide how to handle this case.
- *
- * \param r_library_path: The returned asset library path with a trailing slash, or an empty string
- *                        if no suitable path is found. Assumed to be a buffer of at least
- *                        #FILE_MAXDIR bytes.
- *
- * \return True if the function could find a valid, that is, a non-empty path to return in \a
- *         r_library_path.
- */
-bool BKE_asset_library_find_suitable_root_path_from_path(
-    const char *input_path, char r_library_path[768 /* FILE_MAXDIR */]);
-/**
- * Uses the current location on disk of the file represented by \a bmain as input to
- * #BKE_asset_library_find_suitable_root_path_from_path(). Refer to it for a design
- * description.
- *
- * \return True if the function could find a valid, that is, a non-empty path to return in \a
- *         r_library_path. If \a bmain wasn't saved into a file yet, the return value will be
- *         false.
- */
-bool BKE_asset_library_find_suitable_root_path_from_main(
-    const struct Main *bmain, char r_library_path[768 /* FILE_MAXDIR */]);
-
 #ifdef __cplusplus
 }
 #endif

source/blender/blenkernel/BKE_asset_library.hh

@@ -27,7 +27,6 @@
 #include "BKE_asset_library.h"
 
 #include "BKE_asset_catalog.hh"
-#include "BKE_callbacks.h"
 
 #include <memory>
 
@@ -37,18 +36,6 @@ struct AssetLibrary {
   std::unique_ptr<AssetCatalogService> catalog_service;
 
   void load(StringRefNull library_root_directory);
-
-  void on_save_handler_register();
-  void on_save_handler_unregister();
-
-  void on_save_post(struct Main *, struct PointerRNA **pointers, const int num_pointers);
-
- private:
-  bCallbackFuncStore on_save_callback_store_;
 };
 
 }  // namespace blender::bke
-
-blender::bke::AssetCatalogService *BKE_asset_library_get_catalog_service(
-    const ::AssetLibrary *library);
-blender::bke::AssetCatalogTree *BKE_asset_library_get_catalog_tree(const ::AssetLibrary *library);

source/blender/blenkernel/BKE_blender_version.h

@@ -39,13 +39,13 @@ extern "C" {
 
 /* Blender file format version. */
 #define BLENDER_FILE_VERSION BLENDER_VERSION
-#define BLENDER_FILE_SUBVERSION 31
+#define BLENDER_FILE_SUBVERSION 25
 
 /* Minimum Blender version that supports reading file written with the current
  * version. Older Blender versions will test this and show a warning if the file
  * was written with too new a version. */
 #define BLENDER_FILE_MIN_VERSION 300
-#define BLENDER_FILE_MIN_SUBVERSION 26
+#define BLENDER_FILE_MIN_SUBVERSION 11
 
 /** User readable version string. */
 const char *BKE_blender_version_string(void);

source/blender/blenkernel/BKE_callbacks.h

@@ -130,7 +130,6 @@ void BKE_callback_exec_id_depsgraph(struct Main *bmain,
                                     struct Depsgraph *depsgraph,
                                     eCbEvent evt);
 void BKE_callback_add(bCallbackFuncStore *funcstore, eCbEvent evt);
-void BKE_callback_remove(bCallbackFuncStore *funcstore, eCbEvent evt);
 
 void BKE_callback_global_init(void);
 void BKE_callback_global_finalize(void);

source/blender/blenkernel/BKE_geometry_set.hh

@@ -281,8 +281,6 @@ struct GeometrySet {
     return this->remove(Component::static_type);
   }
 
-  void keep_only(const blender::Span<GeometryComponentType> component_types);
-
   void add(const GeometryComponent &component);
 
   blender::Vector<const GeometryComponent *> get_components_for_read() const;
@@ -311,10 +309,6 @@ struct GeometrySet {
       bool include_instances,
       blender::Map<blender::bke::AttributeIDRef, AttributeKind> &r_attributes) const;
 
-  using ForeachSubGeometryCallback = blender::FunctionRef<void(GeometrySet &geometry_set)>;
-
-  void modify_geometry_sets(ForeachSubGeometryCallback callback);
-
   /* Utility methods for creation. */
   static GeometrySet create_with_mesh(
       Mesh *mesh, GeometryOwnershipType ownership = GeometryOwnershipType::Owned);
@@ -329,8 +323,6 @@ struct GeometrySet {
   bool has_instances() const;
   bool has_volume() const;
   bool has_curve() const;
-  bool has_realized_data() const;
-  bool is_empty() const;
 
   const Mesh *get_mesh_for_read() const;
   const PointCloud *get_pointcloud_for_read() const;
@@ -505,40 +497,13 @@ class InstanceReference {
   {
   }
 
-  InstanceReference(const InstanceReference &other) : type_(other.type_), data_(other.data_)
+  InstanceReference(const InstanceReference &src) : type_(src.type_), data_(src.data_)
   {
-    if (other.geometry_set_) {
-      geometry_set_ = std::make_unique<GeometrySet>(*other.geometry_set_);
+    if (src.type_ == Type::GeometrySet) {
+      geometry_set_ = std::make_unique<GeometrySet>(*src.geometry_set_);
     }
   }
 
-  InstanceReference(InstanceReference &&other)
-      : type_(other.type_), data_(other.data_), geometry_set_(std::move(other.geometry_set_))
-  {
-    other.type_ = Type::None;
-    other.data_ = nullptr;
-  }
-
-  InstanceReference &operator=(const InstanceReference &other)
-  {
-    if (this == &other) {
-      return *this;
-    }
-    this->~InstanceReference();
-    new (this) InstanceReference(other);
-    return *this;
-  }
-
-  InstanceReference &operator=(InstanceReference &&other)
-  {
-    if (this == &other) {
-      return *this;
-    }
-    this->~InstanceReference();
-    new (this) InstanceReference(std::move(other));
-    return *this;
-  }
-
   Type type() const
   {
     return type_;
@@ -630,7 +595,6 @@ class InstancesComponent : public GeometryComponent {
   void add_instance(int instance_handle, const blender::float4x4 &transform, const int id = -1);
 
   blender::Span<InstanceReference> references() const;
-  void remove_unused_references();
 
   void ensure_geometry_instances();
   GeometrySet &geometry_set_from_reference(const int reference_index);
@@ -722,13 +686,6 @@ class AttributeFieldInput : public fn::FieldInput {
   {
   }
 
-  template<typename T> static fn::Field<T> Create(std::string name)
-  {
-    const CPPType &type = CPPType::get<T>();
-    auto field_input = std::make_shared<AttributeFieldInput>(std::move(name), type);
-    return fn::Field<T>{field_input};
-  }
-
   StringRefNull attribute_name() const
   {
     return name_;

source/blender/blenkernel/BKE_lib_override.h

@@ -173,8 +173,6 @@ void BKE_lib_override_library_main_unused_cleanup(struct Main *bmain);
 void BKE_lib_override_library_update(struct Main *bmain, struct ID *local);
 void BKE_lib_override_library_main_update(struct Main *bmain);
 
-bool BKE_lib_override_library_id_is_user_deletable(struct Main *bmain, struct ID *id);
-
 /* Storage (.blend file writing) part. */
 
 /* For now, we just use a temp main list. */