Mesh: Replace MLoop struct with generic attributes #104424
|
@ -172,6 +172,7 @@ if(UNIX AND NOT APPLE)
|
|||
include(cmake/wayland_protocols.cmake)
|
||||
# Can be removed when the build-bot upgrades to v1.20.x or newer.
|
||||
include(cmake/wayland.cmake)
|
||||
include(cmake/wayland_libdecor.cmake)
|
||||
endif()
|
||||
|
||||
include(cmake/harvest.cmake)
|
||||
|
|
|
@ -133,6 +133,7 @@ download_source(NASM)
|
|||
download_source(XR_OPENXR_SDK)
|
||||
download_source(WL_PROTOCOLS)
|
||||
download_source(WAYLAND)
|
||||
download_source(WAYLAND_LIBDECOR)
|
||||
download_source(ISPC)
|
||||
download_source(GMP)
|
||||
download_source(POTRACE)
|
||||
|
|
|
@ -118,6 +118,8 @@ else()
|
|||
|
||||
harvest(wayland-protocols/share/wayland-protocols wayland-protocols/share/wayland-protocols/ "*.xml")
|
||||
harvest(wayland/bin wayland/bin "wayland-scanner")
|
||||
harvest(wayland/include wayland/include "*.h")
|
||||
harvest(wayland_libdecor/include wayland_libdecor/include "*.h")
|
||||
else()
|
||||
harvest(blosc/lib openvdb/lib "*.a")
|
||||
harvest(xml2/lib opencollada/lib "*.a")
|
||||
|
|
|
@ -464,6 +464,12 @@ set(WAYLAND_URI https://gitlab.freedesktop.org/wayland/wayland/-/releases/1.21.0
|
|||
set(WAYLAND_HASH f2653a2293bcd882d756c6a83d278903)
|
||||
set(WAYLAND_HASH_TYPE MD5)
|
||||
|
||||
set(WAYLAND_LIBDECOR_VERSION 0.1.0)
|
||||
set(WAYLAND_LIBDECOR_FILE libdecor-${WAYLAND_LIBDECOR_VERSION}.tar.xz)
|
||||
set(WAYLAND_LIBDECOR_URI https://gitlab.gnome.org/jadahl/libdecor/uploads/81adf91d27620e20bcc5f6b9b312d768/libdecor-${WAYLAND_LIBDECOR_VERSION}.tar.xz )
|
||||
set(WAYLAND_LIBDECOR_HASH 47b59eba76faa3787f0878bf8700e912)
|
||||
set(WAYLAND_LIBDECOR_HASH_TYPE MD5)
|
||||
|
||||
set(ISPC_VERSION v1.17.0)
|
||||
set(ISPC_URI https://github.com/ispc/ispc/archive/${ISPC_VERSION}.tar.gz)
|
||||
set(ISPC_HASH 4f476a3109332a77fe839a9014c60ca9)
|
||||
|
|
|
@ -6,9 +6,11 @@ ExternalProject_Add(external_wayland
|
|||
URL_HASH ${WAYLAND_HASH_TYPE}=${WAYLAND_HASH}
|
||||
PREFIX ${BUILD_DIR}/wayland
|
||||
PATCH_COMMAND ${PATCH_CMD} -d ${BUILD_DIR}/wayland/src/external_wayland < ${PATCH_DIR}/wayland.diff
|
||||
# Use `-E` so the `PKG_CONFIG_PATH` can be defined to link against our own LIBEXPAT.
|
||||
CONFIGURE_COMMAND ${CMAKE_COMMAND} -E env PKG_CONFIG_PATH=${LIBDIR}/expat/lib/pkgconfig:${LIBDIR}/xml2/lib/pkgconfig:$PKG_CONFIG_PATH
|
||||
meson --prefix ${LIBDIR}/wayland -Ddocumentation=false -Dtests=false -Dlibraries=false . ../external_wayland
|
||||
# Use `-E` so the `PKG_CONFIG_PATH` can be defined to link against our own LIBEXPAT & LIBXML2.
|
||||
# Note that passing link args "ffi/lib" should not be needed, but
|
||||
# `pkgconfig` would incorrectly look in "ffi/lib/../lib64" otherwise.
|
||||
CONFIGURE_COMMAND ${CMAKE_COMMAND} -E env PKG_CONFIG_PATH=${LIBDIR}/expat/lib/pkgconfig:${LIBDIR}/xml2/lib/pkgconfig:${LIBDIR}/ffi/lib/pkgconfig:$PKG_CONFIG_PATH
|
||||
meson --prefix ${LIBDIR}/wayland -Ddocumentation=false -Dtests=false -Dc_link_args=-L${LIBDIR}/ffi/lib . ../external_wayland
|
||||
BUILD_COMMAND ninja
|
||||
INSTALL_COMMAND ninja install
|
||||
)
|
||||
|
@ -17,4 +19,5 @@ add_dependencies(
|
|||
external_wayland
|
||||
external_expat
|
||||
external_xml2
|
||||
external_ffi
|
||||
)
|
||||
|
|
|
@ -0,0 +1,15 @@
|
|||
# SPDX-License-Identifier: GPL-2.0-or-later
|
||||
|
||||
# NOTE: currently only the header file is extracted, no compilation is needed
|
||||
# as the library is dynamically loaded when found on the system.
|
||||
|
||||
ExternalProject_Add(external_wayland_libdecor
|
||||
URL file://${PACKAGE_DIR}/${WAYLAND_LIBDECOR_FILE}
|
||||
DOWNLOAD_DIR ${DOWNLOAD_DIR}
|
||||
URL_HASH ${WAYLAND_LIBDECOR_HASH_TYPE}=${WAYLAND_LIBDECOR_HASH}
|
||||
PREFIX ${BUILD_DIR}/wayland_libdecor
|
||||
BUILD_COMMAND echo .
|
||||
CONFIGURE_COMMAND echo .
|
||||
INSTALL_COMMAND cp ../external_wayland_libdecor/src/libdecor.h ${LIBDIR}/wayland_libdecor/include/libdecor-0/libdecor.h
|
||||
INSTALL_DIR ${LIBDIR}/wayland_libdecor/include/libdecor-0
|
||||
)
|
|
@ -62,6 +62,13 @@ FIND_LIBRARY(SYCL_LIBRARY
|
|||
lib64 lib
|
||||
)
|
||||
|
||||
if(WIN32)
|
||||
string(REPLACE ".lib" "d.lib" SYCL_LIBRARY_DEBUG ${SYCL_LIBRARY})
|
||||
set(SYCL_LIBRARY_DEBUG ${SYCL_LIBRARY_DEBUG} CACHE FILEPATH "Path to SYCL debug library")
|
||||
else()
|
||||
set(SYCL_LIBRARY_DEBUG ${SYCL_LIBRARY} CACHE FILEPATH "Path to SYCL debug library")
|
||||
endif()
|
||||
|
||||
FIND_PATH(SYCL_INCLUDE_DIR
|
||||
NAMES
|
||||
CL/sycl.hpp
|
||||
|
@ -85,4 +92,5 @@ ENDIF()
|
|||
|
||||
MARK_AS_ADVANCED(
|
||||
_SYCL_INCLUDE_PARENT_DIR
|
||||
SYCL_LIBRARY_DEBUG
|
||||
)
|
||||
|
|
|
@ -335,10 +335,18 @@ if(WITH_CYCLES_DEVICE_ONEAPI)
|
|||
set(LEVEL_ZERO_ROOT_DIR ${CYCLES_LEVEL_ZERO})
|
||||
endif()
|
||||
|
||||
set(CYCLES_SYCL ${LIBDIR}/dpcpp CACHE PATH "Path to DPC++ and SYCL installation")
|
||||
set(CYCLES_SYCL ${LIBDIR}/dpcpp CACHE PATH "Path to oneAPI DPC++ compiler")
|
||||
if(EXISTS ${CYCLES_SYCL} AND NOT SYCL_ROOT_DIR)
|
||||
set(SYCL_ROOT_DIR ${CYCLES_SYCL})
|
||||
endif()
|
||||
file(GLOB _sycl_runtime_libraries
|
||||
${SYCL_ROOT_DIR}/lib/libsycl.so
|
||||
${SYCL_ROOT_DIR}/lib/libsycl.so.[0-9]
|
||||
${SYCL_ROOT_DIR}/lib/libsycl.so.[0-9].[0-9].[0-9]-[0-9]
|
||||
${SYCL_ROOT_DIR}/lib/libpi_level_zero.so
|
||||
)
|
||||
list(APPEND PLATFORM_BUNDLED_LIBRARIES ${_sycl_runtime_libraries})
|
||||
unset(_sycl_runtime_libraries)
|
||||
endif()
|
||||
|
||||
if(WITH_OPENVDB)
|
||||
|
@ -691,14 +699,23 @@ endif()
|
|||
|
||||
if(WITH_GHOST_WAYLAND)
|
||||
find_package(PkgConfig)
|
||||
pkg_check_modules(wayland-client wayland-client>=1.12)
|
||||
pkg_check_modules(wayland-egl wayland-egl)
|
||||
pkg_check_modules(wayland-scanner wayland-scanner)
|
||||
pkg_check_modules(xkbcommon xkbcommon)
|
||||
pkg_check_modules(wayland-cursor wayland-cursor)
|
||||
pkg_check_modules(wayland-protocols wayland-protocols>=1.15)
|
||||
|
||||
if(${wayland-protocols_FOUND})
|
||||
# When dynamically linked WAYLAND is used and `${LIBDIR}/wayland` is present,
|
||||
# there is no need to search for the libraries as they are not needed for building.
|
||||
# Only the headers are needed which can reference the known paths.
|
||||
if(EXISTS "${LIBDIR}/wayland" AND WITH_GHOST_WAYLAND_DYNLOAD)
|
||||
set(_use_system_wayland OFF)
|
||||
else()
|
||||
set(_use_system_wayland ON)
|
||||
endif()
|
||||
|
||||
if(_use_system_wayland)
|
||||
pkg_check_modules(wayland-client wayland-client>=1.12)
|
||||
pkg_check_modules(wayland-egl wayland-egl)
|
||||
pkg_check_modules(wayland-scanner wayland-scanner)
|
||||
pkg_check_modules(wayland-cursor wayland-cursor)
|
||||
pkg_check_modules(wayland-protocols wayland-protocols>=1.15)
|
||||
pkg_get_variable(WAYLAND_PROTOCOLS_DIR wayland-protocols pkgdatadir)
|
||||
else()
|
||||
# CentOS 7 packages have too old a version, a newer version exist in the
|
||||
|
@ -712,6 +729,15 @@ if(WITH_GHOST_WAYLAND)
|
|||
if(EXISTS ${WAYLAND_PROTOCOLS_DIR})
|
||||
set(wayland-protocols_FOUND ON)
|
||||
endif()
|
||||
|
||||
set(wayland-client_INCLUDE_DIRS "${LIBDIR}/wayland/include")
|
||||
set(wayland-egl_INCLUDE_DIRS "${LIBDIR}/wayland/include")
|
||||
set(wayland-cursor_INCLUDE_DIRS "${LIBDIR}/wayland/include")
|
||||
|
||||
set(wayland-client_FOUND ON)
|
||||
set(wayland-egl_FOUND ON)
|
||||
set(wayland-scanner_FOUND ON)
|
||||
set(wayland-cursor_FOUND ON)
|
||||
endif()
|
||||
|
||||
if (NOT ${wayland-client_FOUND})
|
||||
|
@ -745,7 +771,11 @@ if(WITH_GHOST_WAYLAND)
|
|||
endif()
|
||||
|
||||
if(WITH_GHOST_WAYLAND_LIBDECOR)
|
||||
pkg_check_modules(libdecor REQUIRED libdecor-0>=0.1)
|
||||
if(_use_system_wayland)
|
||||
pkg_check_modules(libdecor REQUIRED libdecor-0>=0.1)
|
||||
else()
|
||||
set(libdecor_INCLUDE_DIRS "${LIBDIR}/wayland_libdecor/include/libdecor-0")
|
||||
endif()
|
||||
endif()
|
||||
|
||||
list(APPEND PLATFORM_LINKLIBS
|
||||
|
@ -815,6 +845,8 @@ if(WITH_GHOST_WAYLAND)
|
|||
# End wayland-scanner version check.
|
||||
|
||||
endif()
|
||||
|
||||
unset(_use_system_wayland)
|
||||
endif()
|
||||
|
||||
if(WITH_GHOST_X11)
|
||||
|
|
|
@ -952,5 +952,23 @@ endif()
|
|||
set(ZSTD_INCLUDE_DIRS ${LIBDIR}/zstd/include)
|
||||
set(ZSTD_LIBRARIES ${LIBDIR}/zstd/lib/zstd_static.lib)
|
||||
|
||||
set(LEVEL_ZERO_ROOT_DIR ${LIBDIR}/level_zero)
|
||||
set(SYCL_ROOT_DIR ${LIBDIR}/dpcpp)
|
||||
if(WITH_CYCLES_DEVICE_ONEAPI)
|
||||
set(LEVEL_ZERO_ROOT_DIR ${LIBDIR}/level_zero)
|
||||
set(CYCLES_SYCL ${LIBDIR}/dpcpp CACHE PATH "Path to oneAPI DPC++ compiler")
|
||||
if(EXISTS ${CYCLES_SYCL} AND NOT SYCL_ROOT_DIR)
|
||||
set(SYCL_ROOT_DIR ${CYCLES_SYCL})
|
||||
endif()
|
||||
file(GLOB _sycl_runtime_libraries_glob
|
||||
${SYCL_ROOT_DIR}/bin/sycl.dll
|
||||
${SYCL_ROOT_DIR}/bin/sycl[0-9].dll
|
||||
)
|
||||
foreach(sycl_runtime_library IN LISTS _sycl_runtime_libraries_glob)
|
||||
string(REPLACE ".dll" "$<$<CONFIG:Debug>:d>.dll" sycl_runtime_library ${sycl_runtime_library})
|
||||
list(APPEND _sycl_runtime_libraries ${sycl_runtime_library})
|
||||
endforeach()
|
||||
unset(_sycl_runtime_libraries_glob)
|
||||
|
||||
list(APPEND _sycl_runtime_libraries ${SYCL_ROOT_DIR}/bin/pi_level_zero.dll)
|
||||
list(APPEND PLATFORM_BUNDLED_LIBRARIES ${_sycl_runtime_libraries})
|
||||
unset(_sycl_runtime_libraries)
|
||||
endif()
|
||||
|
|
|
@ -525,8 +525,8 @@ class CyclesRenderSettings(bpy.types.PropertyGroup):
|
|||
|
||||
use_deterministic_guiding: BoolProperty(
|
||||
name="Deterministic",
|
||||
description="Makes path guiding deterministic which means renderings will be"
|
||||
"reproducible with the same pixel values every time. This feature slows down"
|
||||
description="Makes path guiding deterministic which means renderings will be "
|
||||
"reproducible with the same pixel values every time. This feature slows down "
|
||||
"training",
|
||||
default=True,
|
||||
)
|
||||
|
@ -562,7 +562,7 @@ class CyclesRenderSettings(bpy.types.PropertyGroup):
|
|||
description="The maximum number of samples used for training path guiding. "
|
||||
"Higher samples lead to more accurate guiding, however may also unnecessarily slow "
|
||||
"down rendering once guiding is accurate enough. "
|
||||
"A value 0 will continue training until the last sample",
|
||||
"A value of 0 will continue training until the last sample",
|
||||
min=0,
|
||||
soft_min=1,
|
||||
default=128,
|
||||
|
|
|
@ -142,7 +142,6 @@ set(SRC
|
|||
${SRC_DUMMY}
|
||||
${SRC_MULTI}
|
||||
${SRC_OPTIX}
|
||||
${SRC_ONEAPI}
|
||||
${SRC_HEADERS}
|
||||
)
|
||||
|
||||
|
@ -188,7 +187,25 @@ if(WITH_CYCLES_DEVICE_METAL)
|
|||
)
|
||||
endif()
|
||||
if (WITH_CYCLES_DEVICE_ONEAPI)
|
||||
if(WIN32)
|
||||
set(cycles_kernel_oneapi_lib ${CMAKE_CURRENT_BINARY_DIR}/../kernel/cycles_kernel_oneapi.lib)
|
||||
else()
|
||||
set(cycles_kernel_oneapi_lib ${CMAKE_CURRENT_BINARY_DIR}/../kernel/libcycles_kernel_oneapi.so)
|
||||
endif()
|
||||
list(APPEND LIB
|
||||
${cycles_kernel_oneapi_lib}
|
||||
"$<$<CONFIG:Debug>:${SYCL_LIBRARY_DEBUG}>"
|
||||
"$<$<CONFIG:Release>:${SYCL_LIBRARY}>"
|
||||
"$<$<CONFIG:RelWithDebInfo>:${SYCL_LIBRARY}>"
|
||||
"$<$<CONFIG:MinSizeRel>:${SYCL_LIBRARY}>"
|
||||
)
|
||||
add_definitions(-DWITH_ONEAPI)
|
||||
list(APPEND SRC
|
||||
${SRC_ONEAPI}
|
||||
)
|
||||
list(APPEND INC_SYS
|
||||
${SYCL_INCLUDE_DIR}
|
||||
)
|
||||
endif()
|
||||
|
||||
if(WITH_OPENIMAGEDENOISE)
|
||||
|
|
|
@ -19,62 +19,12 @@
|
|||
|
||||
CCL_NAMESPACE_BEGIN
|
||||
|
||||
#ifdef WITH_ONEAPI
|
||||
static OneAPIDLLInterface oneapi_dll;
|
||||
#endif
|
||||
|
||||
#ifdef _WIN32
|
||||
# define LOAD_ONEAPI_SHARED_LIBRARY(path) (void *)(LoadLibrary(path))
|
||||
# define LOAD_ONEAPI_SHARED_LIBRARY_ERROR() GetLastError()
|
||||
# define FREE_SHARED_LIBRARY(handle) FreeLibrary((HMODULE)handle)
|
||||
# define GET_SHARED_LIBRARY_SYMBOL(handle, name) GetProcAddress((HMODULE)handle, name)
|
||||
#elif __linux__
|
||||
# define LOAD_ONEAPI_SHARED_LIBRARY(path) dlopen(path, RTLD_NOW)
|
||||
# define LOAD_ONEAPI_SHARED_LIBRARY_ERROR() dlerror()
|
||||
# define FREE_SHARED_LIBRARY(handle) dlclose(handle)
|
||||
# define GET_SHARED_LIBRARY_SYMBOL(handle, name) dlsym(handle, name)
|
||||
#endif
|
||||
|
||||
bool device_oneapi_init()
|
||||
{
|
||||
#if !defined(WITH_ONEAPI)
|
||||
return false;
|
||||
#else
|
||||
|
||||
string lib_path = path_get("lib");
|
||||
# ifdef _WIN32
|
||||
lib_path = path_join(lib_path, "cycles_kernel_oneapi.dll");
|
||||
# else
|
||||
lib_path = path_join(lib_path, "cycles_kernel_oneapi.so");
|
||||
# endif
|
||||
void *lib_handle = LOAD_ONEAPI_SHARED_LIBRARY(lib_path.c_str());
|
||||
|
||||
/* This shouldn't happen, but it still makes sense to have a branch for this. */
|
||||
if (lib_handle == NULL) {
|
||||
LOG(ERROR) << "oneAPI kernel shared library cannot be loaded: "
|
||||
<< LOAD_ONEAPI_SHARED_LIBRARY_ERROR();
|
||||
return false;
|
||||
}
|
||||
|
||||
# define DLL_INTERFACE_CALL(function, return_type, ...) \
|
||||
(oneapi_dll.function) = reinterpret_cast<decltype(oneapi_dll.function)>( \
|
||||
GET_SHARED_LIBRARY_SYMBOL(lib_handle, #function)); \
|
||||
if (oneapi_dll.function == NULL) { \
|
||||
LOG(ERROR) << "oneAPI shared library function \"" << #function \
|
||||
<< "\" has not been loaded from kernel shared - disable oneAPI " \
|
||||
"library disable oneAPI implementation due to this"; \
|
||||
FREE_SHARED_LIBRARY(lib_handle); \
|
||||
return false; \
|
||||
}
|
||||
# include "kernel/device/oneapi/dll_interface_template.h"
|
||||
# undef DLL_INTERFACE_CALL
|
||||
|
||||
VLOG_INFO << "oneAPI kernel shared library has been loaded successfully";
|
||||
|
||||
/* We need to have this oneapi kernel shared library during all life-span of the Blender.
|
||||
* So it is not unloaded because of this.
|
||||
* FREE_SHARED_LIBRARY(lib_handle); */
|
||||
|
||||
/* NOTE(@nsirgien): we need to enable JIT cache from here and
|
||||
* right now this cache policy is controlled by env. variables. */
|
||||
/* NOTE(hallade) we also disable use of copy engine as it
|
||||
|
@ -109,17 +59,10 @@ bool device_oneapi_init()
|
|||
#endif
|
||||
}
|
||||
|
||||
#if defined(_WIN32) || defined(__linux__)
|
||||
# undef LOAD_SYCL_SHARED_LIBRARY
|
||||
# undef LOAD_ONEAPI_SHARED_LIBRARY
|
||||
# undef FREE_SHARED_LIBRARY
|
||||
# undef GET_SHARED_LIBRARY_SYMBOL
|
||||
#endif
|
||||
|
||||
Device *device_oneapi_create(const DeviceInfo &info, Stats &stats, Profiler &profiler)
|
||||
{
|
||||
#ifdef WITH_ONEAPI
|
||||
return new OneapiDevice(info, oneapi_dll, stats, profiler);
|
||||
return new OneapiDevice(info, stats, profiler);
|
||||
#else
|
||||
(void)info;
|
||||
(void)stats;
|
||||
|
@ -165,7 +108,7 @@ static void device_iterator_cb(const char *id, const char *name, int num, void *
|
|||
void device_oneapi_info(vector<DeviceInfo> &devices)
|
||||
{
|
||||
#ifdef WITH_ONEAPI
|
||||
(oneapi_dll.oneapi_iterate_devices)(device_iterator_cb, &devices);
|
||||
OneapiDevice::iterate_devices(device_iterator_cb, &devices);
|
||||
#else /* WITH_ONEAPI */
|
||||
(void)devices;
|
||||
#endif /* WITH_ONEAPI */
|
||||
|
@ -175,10 +118,10 @@ string device_oneapi_capabilities()
|
|||
{
|
||||
string capabilities;
|
||||
#ifdef WITH_ONEAPI
|
||||
char *c_capabilities = (oneapi_dll.oneapi_device_capabilities)();
|
||||
char *c_capabilities = OneapiDevice::device_capabilities();
|
||||
if (c_capabilities) {
|
||||
capabilities = c_capabilities;
|
||||
(oneapi_dll.oneapi_free)(c_capabilities);
|
||||
free(c_capabilities);
|
||||
}
|
||||
#endif
|
||||
return capabilities;
|
||||
|
|
|
@ -8,7 +8,7 @@
|
|||
# include "util/debug.h"
|
||||
# include "util/log.h"
|
||||
|
||||
# include "kernel/device/oneapi/kernel.h"
|
||||
# include "kernel/device/oneapi/globals.h"
|
||||
|
||||
CCL_NAMESPACE_BEGIN
|
||||
|
||||
|
@ -19,26 +19,19 @@ static void queue_error_cb(const char *message, void *user_ptr)
|
|||
}
|
||||
}
|
||||
|
||||
OneapiDevice::OneapiDevice(const DeviceInfo &info,
|
||||
OneAPIDLLInterface &oneapi_dll_object,
|
||||
Stats &stats,
|
||||
Profiler &profiler)
|
||||
OneapiDevice::OneapiDevice(const DeviceInfo &info, Stats &stats, Profiler &profiler)
|
||||
: Device(info, stats, profiler),
|
||||
device_queue_(nullptr),
|
||||
texture_info_(this, "texture_info", MEM_GLOBAL),
|
||||
kg_memory_(nullptr),
|
||||
kg_memory_device_(nullptr),
|
||||
kg_memory_size_(0),
|
||||
oneapi_dll_(oneapi_dll_object)
|
||||
kg_memory_size_(0)
|
||||
{
|
||||
need_texture_info_ = false;
|
||||
|
||||
oneapi_dll_.oneapi_set_error_cb(queue_error_cb, &oneapi_error_string_);
|
||||
oneapi_set_error_cb(queue_error_cb, &oneapi_error_string_);
|
||||
|
||||
/* OneAPI calls should be initialized on this moment. */
|
||||
assert(oneapi_dll_.oneapi_create_queue != nullptr);
|
||||
|
||||
bool is_finished_ok = oneapi_dll_.oneapi_create_queue(device_queue_, info.num);
|
||||
bool is_finished_ok = create_queue(device_queue_, info.num);
|
||||
if (is_finished_ok == false) {
|
||||
set_error("oneAPI queue initialization error: got runtime exception \"" +
|
||||
oneapi_error_string_ + "\"");
|
||||
|
@ -50,7 +43,7 @@ OneapiDevice::OneapiDevice(const DeviceInfo &info,
|
|||
}
|
||||
|
||||
size_t globals_segment_size;
|
||||
is_finished_ok = oneapi_dll_.oneapi_kernel_globals_size(device_queue_, globals_segment_size);
|
||||
is_finished_ok = kernel_globals_size(device_queue_, globals_segment_size);
|
||||
if (is_finished_ok == false) {
|
||||
set_error("oneAPI constant memory initialization got runtime exception \"" +
|
||||
oneapi_error_string_ + "\"");
|
||||
|
@ -59,27 +52,27 @@ OneapiDevice::OneapiDevice(const DeviceInfo &info,
|
|||
VLOG_DEBUG << "Successfully created global/constant memory segment (kernel globals object)";
|
||||
}
|
||||
|
||||
kg_memory_ = oneapi_dll_.oneapi_usm_aligned_alloc_host(device_queue_, globals_segment_size, 16);
|
||||
oneapi_dll_.oneapi_usm_memset(device_queue_, kg_memory_, 0, globals_segment_size);
|
||||
kg_memory_ = usm_aligned_alloc_host(device_queue_, globals_segment_size, 16);
|
||||
usm_memset(device_queue_, kg_memory_, 0, globals_segment_size);
|
||||
|
||||
kg_memory_device_ = oneapi_dll_.oneapi_usm_alloc_device(device_queue_, globals_segment_size);
|
||||
kg_memory_device_ = usm_alloc_device(device_queue_, globals_segment_size);
|
||||
|
||||
kg_memory_size_ = globals_segment_size;
|
||||
|
||||
max_memory_on_device_ = oneapi_dll_.oneapi_get_memcapacity(device_queue_);
|
||||
max_memory_on_device_ = get_memcapacity();
|
||||
}
|
||||
|
||||
OneapiDevice::~OneapiDevice()
|
||||
{
|
||||
texture_info_.free();
|
||||
oneapi_dll_.oneapi_usm_free(device_queue_, kg_memory_);
|
||||
oneapi_dll_.oneapi_usm_free(device_queue_, kg_memory_device_);
|
||||
usm_free(device_queue_, kg_memory_);
|
||||
usm_free(device_queue_, kg_memory_device_);
|
||||
|
||||
for (ConstMemMap::iterator mt = const_mem_map_.begin(); mt != const_mem_map_.end(); mt++)
|
||||
delete mt->second;
|
||||
|
||||
if (device_queue_)
|
||||
oneapi_dll_.oneapi_free_queue(device_queue_);
|
||||
free_queue(device_queue_);
|
||||
}
|
||||
|
||||
bool OneapiDevice::check_peer_access(Device * /*peer_device*/)
|
||||
|
@ -99,7 +92,7 @@ bool OneapiDevice::load_kernels(const uint requested_features)
|
|||
* with specialization constants, but it hasn't been implemented yet. */
|
||||
(void)requested_features;
|
||||
|
||||
bool is_finished_ok = oneapi_dll_.oneapi_run_test_kernel(device_queue_);
|
||||
bool is_finished_ok = oneapi_run_test_kernel(device_queue_);
|
||||
if (is_finished_ok == false) {
|
||||
set_error("oneAPI kernel load: got runtime exception \"" + oneapi_error_string_ + "\"");
|
||||
}
|
||||
|
@ -138,7 +131,7 @@ void OneapiDevice::generic_alloc(device_memory &mem)
|
|||
* type has been used for oneAPI device in order to better fit in Cycles architecture. */
|
||||
void *device_pointer = nullptr;
|
||||
if (mem.memory_size() + stats.mem_used < max_memory_on_device_)
|
||||
device_pointer = oneapi_dll_.oneapi_usm_alloc_device(device_queue_, memory_size);
|
||||
device_pointer = usm_alloc_device(device_queue_, memory_size);
|
||||
if (device_pointer == nullptr) {
|
||||
set_error("oneAPI kernel - device memory allocation error for " +
|
||||
string_human_readable_size(mem.memory_size()) +
|
||||
|
@ -163,8 +156,7 @@ void OneapiDevice::generic_copy_to(device_memory &mem)
|
|||
/* Copy operation from host shouldn't be requested if there is no memory allocated on host. */
|
||||
assert(mem.host_pointer);
|
||||
assert(device_queue_);
|
||||
oneapi_dll_.oneapi_usm_memcpy(
|
||||
device_queue_, (void *)mem.device_pointer, (void *)mem.host_pointer, memory_size);
|
||||
usm_memcpy(device_queue_, (void *)mem.device_pointer, (void *)mem.host_pointer, memory_size);
|
||||
}
|
||||
|
||||
/* TODO: Make sycl::queue part of OneapiQueue and avoid using pointers to sycl::queue. */
|
||||
|
@ -178,11 +170,6 @@ string OneapiDevice::oneapi_error_message()
|
|||
return string(oneapi_error_string_);
|
||||
}
|
||||
|
||||
OneAPIDLLInterface OneapiDevice::oneapi_dll_object()
|
||||
{
|
||||
return oneapi_dll_;
|
||||
}
|
||||
|
||||
void *OneapiDevice::kernel_globals_device_pointer()
|
||||
{
|
||||
return kg_memory_device_;
|
||||
|
@ -198,7 +185,7 @@ void OneapiDevice::generic_free(device_memory &mem)
|
|||
mem.device_size = 0;
|
||||
|
||||
assert(device_queue_);
|
||||
oneapi_dll_.oneapi_usm_free(device_queue_, (void *)mem.device_pointer);
|
||||
usm_free(device_queue_, (void *)mem.device_pointer);
|
||||
mem.device_pointer = 0;
|
||||
}
|
||||
|
||||
|
@ -266,8 +253,7 @@ void OneapiDevice::mem_copy_from(device_memory &mem, size_t y, size_t w, size_t
|
|||
if (mem.device_pointer) {
|
||||
char *shifted_host = reinterpret_cast<char *>(mem.host_pointer) + offset;
|
||||
char *shifted_device = reinterpret_cast<char *>(mem.device_pointer) + offset;
|
||||
bool is_finished_ok = oneapi_dll_.oneapi_usm_memcpy(
|
||||
device_queue_, shifted_host, shifted_device, size);
|
||||
bool is_finished_ok = usm_memcpy(device_queue_, shifted_host, shifted_device, size);
|
||||
if (is_finished_ok == false) {
|
||||
set_error("oneAPI memory operation error: got runtime exception \"" +
|
||||
oneapi_error_string_ + "\"");
|
||||
|
@ -292,7 +278,7 @@ void OneapiDevice::mem_zero(device_memory &mem)
|
|||
}
|
||||
|
||||
assert(device_queue_);
|
||||
bool is_finished_ok = oneapi_dll_.oneapi_usm_memset(
|
||||
bool is_finished_ok = usm_memset(
|
||||
device_queue_, (void *)mem.device_pointer, 0, mem.memory_size());
|
||||
if (is_finished_ok == false) {
|
||||
set_error("oneAPI memory operation error: got runtime exception \"" + oneapi_error_string_ +
|
||||
|
@ -349,10 +335,9 @@ void OneapiDevice::const_copy_to(const char *name, void *host, size_t size)
|
|||
memcpy(data->data(), host, size);
|
||||
data->copy_to_device();
|
||||
|
||||
oneapi_dll_.oneapi_set_global_memory(
|
||||
device_queue_, kg_memory_, name, (void *)data->device_pointer);
|
||||
set_global_memory(device_queue_, kg_memory_, name, (void *)data->device_pointer);
|
||||
|
||||
oneapi_dll_.oneapi_usm_memcpy(device_queue_, kg_memory_device_, kg_memory_, kg_memory_size_);
|
||||
usm_memcpy(device_queue_, kg_memory_device_, kg_memory_, kg_memory_size_);
|
||||
}
|
||||
|
||||
void OneapiDevice::global_alloc(device_memory &mem)
|
||||
|
@ -367,10 +352,9 @@ void OneapiDevice::global_alloc(device_memory &mem)
|
|||
generic_alloc(mem);
|
||||
generic_copy_to(mem);
|
||||
|
||||
oneapi_dll_.oneapi_set_global_memory(
|
||||
device_queue_, kg_memory_, mem.name, (void *)mem.device_pointer);
|
||||
set_global_memory(device_queue_, kg_memory_, mem.name, (void *)mem.device_pointer);
|
||||
|
||||
oneapi_dll_.oneapi_usm_memcpy(device_queue_, kg_memory_device_, kg_memory_, kg_memory_size_);
|
||||
usm_memcpy(device_queue_, kg_memory_device_, kg_memory_, kg_memory_size_);
|
||||
}
|
||||
|
||||
void OneapiDevice::global_free(device_memory &mem)
|
||||
|
@ -410,18 +394,6 @@ unique_ptr<DeviceQueue> OneapiDevice::gpu_queue_create()
|
|||
return make_unique<OneapiDeviceQueue>(this);
|
||||
}
|
||||
|
||||
int OneapiDevice::get_num_multiprocessors()
|
||||
{
|
||||
assert(device_queue_);
|
||||
return oneapi_dll_.oneapi_get_num_multiprocessors(device_queue_);
|
||||
}
|
||||
|
||||
int OneapiDevice::get_max_num_threads_per_multiprocessor()
|
||||
{
|
||||
assert(device_queue_);
|
||||
return oneapi_dll_.oneapi_get_max_num_threads_per_multiprocessor(device_queue_);
|
||||
}
|
||||
|
||||
bool OneapiDevice::should_use_graphics_interop()
|
||||
{
|
||||
/* NOTE(@nsirgien): oneAPI doesn't yet support direct writing into graphics API objects, so
|
||||
|
@ -432,13 +404,460 @@ bool OneapiDevice::should_use_graphics_interop()
|
|||
void *OneapiDevice::usm_aligned_alloc_host(size_t memory_size, size_t alignment)
|
||||
{
|
||||
assert(device_queue_);
|
||||
return oneapi_dll_.oneapi_usm_aligned_alloc_host(device_queue_, memory_size, alignment);
|
||||
return usm_aligned_alloc_host(device_queue_, memory_size, alignment);
|
||||
}
|
||||
|
||||
void OneapiDevice::usm_free(void *usm_ptr)
|
||||
{
|
||||
assert(device_queue_);
|
||||
return oneapi_dll_.oneapi_usm_free(device_queue_, usm_ptr);
|
||||
return usm_free(device_queue_, usm_ptr);
|
||||
}
|
||||
|
||||
void OneapiDevice::check_usm(SyclQueue *queue_, const void *usm_ptr, bool allow_host = false)
|
||||
{
|
||||
# ifdef _DEBUG
|
||||
sycl::queue *queue = reinterpret_cast<sycl::queue *>(queue_);
|
||||
sycl::info::device_type device_type =
|
||||
queue->get_device().get_info<sycl::info::device::device_type>();
|
||||
sycl::usm::alloc usm_type = get_pointer_type(usm_ptr, queue->get_context());
|
||||
(void)usm_type;
|
||||
assert(usm_type == sycl::usm::alloc::device ||
|
||||
((device_type == sycl::info::device_type::host ||
|
||||
device_type == sycl::info::device_type::cpu || allow_host) &&
|
||||
usm_type == sycl::usm::alloc::host));
|
||||
# endif
|
||||
}
|
||||
|
||||
bool OneapiDevice::create_queue(SyclQueue *&external_queue, int device_index)
|
||||
{
|
||||
bool finished_correct = true;
|
||||
try {
|
||||
std::vector<sycl::device> devices = OneapiDevice::available_devices();
|
||||
if (device_index < 0 || device_index >= devices.size()) {
|
||||
return false;
|
||||
}
|
||||
sycl::queue *created_queue = new sycl::queue(devices[device_index],
|
||||
sycl::property::queue::in_order());
|
||||
external_queue = reinterpret_cast<SyclQueue *>(created_queue);
|
||||
}
|
||||
catch (sycl::exception const &e) {
|
||||
finished_correct = false;
|
||||
oneapi_error_string_ = e.what();
|
||||
}
|
||||
return finished_correct;
|
||||
}
|
||||
|
||||
void OneapiDevice::free_queue(SyclQueue *queue_)
|
||||
{
|
||||
assert(queue_);
|
||||
sycl::queue *queue = reinterpret_cast<sycl::queue *>(queue_);
|
||||
delete queue;
|
||||
}
|
||||
|
||||
void *OneapiDevice::usm_aligned_alloc_host(SyclQueue *queue_, size_t memory_size, size_t alignment)
|
||||
{
|
||||
assert(queue_);
|
||||
sycl::queue *queue = reinterpret_cast<sycl::queue *>(queue_);
|
||||
return sycl::aligned_alloc_host(alignment, memory_size, *queue);
|
||||
}
|
||||
|
||||
void *OneapiDevice::usm_alloc_device(SyclQueue *queue_, size_t memory_size)
|
||||
{
|
||||
assert(queue_);
|
||||
sycl::queue *queue = reinterpret_cast<sycl::queue *>(queue_);
|
||||
return sycl::malloc_device(memory_size, *queue);
|
||||
}
|
||||
|
||||
void OneapiDevice::usm_free(SyclQueue *queue_, void *usm_ptr)
|
||||
{
|
||||
assert(queue_);
|
||||
sycl::queue *queue = reinterpret_cast<sycl::queue *>(queue_);
|
||||
OneapiDevice::check_usm(queue_, usm_ptr, true);
|
||||
sycl::free(usm_ptr, *queue);
|
||||
}
|
||||
|
||||
bool OneapiDevice::usm_memcpy(SyclQueue *queue_, void *dest, void *src, size_t num_bytes)
|
||||
{
|
||||
assert(queue_);
|
||||
sycl::queue *queue = reinterpret_cast<sycl::queue *>(queue_);
|
||||
OneapiDevice::check_usm(queue_, dest, true);
|
||||
OneapiDevice::check_usm(queue_, src, true);
|
||||
sycl::event mem_event = queue->memcpy(dest, src, num_bytes);
|
||||
# ifdef WITH_CYCLES_DEBUG
|
||||
try {
|
||||
/* NOTE(@nsirgien) Waiting on memory operation may give more precise error
|
||||
* messages. Due to impact on occupancy, it makes sense to enable it only during Cycles debug.
|
||||
*/
|
||||
mem_event.wait_and_throw();
|
||||
return true;
|
||||
}
|
||||
catch (sycl::exception const &e) {
|
||||
oneapi_error_string_ = e.what();
|
||||
return false;
|
||||
}
|
||||
# else
|
||||
sycl::usm::alloc dest_type = get_pointer_type(dest, queue->get_context());
|
||||
sycl::usm::alloc src_type = get_pointer_type(src, queue->get_context());
|
||||
bool from_device_to_host = dest_type == sycl::usm::alloc::host &&
|
||||
src_type == sycl::usm::alloc::device;
|
||||
bool host_or_device_memop_with_offset = dest_type == sycl::usm::alloc::unknown ||
|
||||
src_type == sycl::usm::alloc::unknown;
|
||||
/* NOTE(@sirgienko) Host-side blocking wait on this operation is mandatory, otherwise the host
|
||||
* may not wait until the end of the transfer before using the memory.
|
||||
*/
|
||||
if (from_device_to_host || host_or_device_memop_with_offset)
|
||||
mem_event.wait();
|
||||
return true;
|
||||
# endif
|
||||
}
|
||||
|
||||
bool OneapiDevice::usm_memset(SyclQueue *queue_,
|
||||
void *usm_ptr,
|
||||
unsigned char value,
|
||||
size_t num_bytes)
|
||||
{
|
||||
assert(queue_);
|
||||
sycl::queue *queue = reinterpret_cast<sycl::queue *>(queue_);
|
||||
OneapiDevice::check_usm(queue_, usm_ptr, true);
|
||||
sycl::event mem_event = queue->memset(usm_ptr, value, num_bytes);
|
||||
# ifdef WITH_CYCLES_DEBUG
|
||||
try {
|
||||
/* NOTE(@nsirgien) Waiting on memory operation may give more precise error
|
||||
* messages. Due to impact on occupancy, it makes sense to enable it only during Cycles debug.
|
||||
*/
|
||||
mem_event.wait_and_throw();
|
||||
return true;
|
||||
}
|
||||
catch (sycl::exception const &e) {
|
||||
oneapi_error_string_ = e.what();
|
||||
return false;
|
||||
}
|
||||
# else
|
||||
(void)mem_event;
|
||||
return true;
|
||||
# endif
|
||||
}
|
||||
|
||||
bool OneapiDevice::queue_synchronize(SyclQueue *queue_)
|
||||
{
|
||||
assert(queue_);
|
||||
sycl::queue *queue = reinterpret_cast<sycl::queue *>(queue_);
|
||||
try {
|
||||
queue->wait_and_throw();
|
||||
return true;
|
||||
}
|
||||
catch (sycl::exception const &e) {
|
||||
oneapi_error_string_ = e.what();
|
||||
return false;
|
||||
}
|
||||
}
|
||||
|
||||
bool OneapiDevice::kernel_globals_size(SyclQueue *queue_, size_t &kernel_global_size)
|
||||
{
|
||||
kernel_global_size = sizeof(KernelGlobalsGPU);
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
void OneapiDevice::set_global_memory(SyclQueue *queue_,
|
||||
void *kernel_globals,
|
||||
const char *memory_name,
|
||||
void *memory_device_pointer)
|
||||
{
|
||||
assert(queue_);
|
||||
assert(kernel_globals);
|
||||
assert(memory_name);
|
||||
assert(memory_device_pointer);
|
||||
KernelGlobalsGPU *globals = (KernelGlobalsGPU *)kernel_globals;
|
||||
OneapiDevice::check_usm(queue_, memory_device_pointer);
|
||||
OneapiDevice::check_usm(queue_, kernel_globals, true);
|
||||
|
||||
std::string matched_name(memory_name);
|
||||
|
||||
/* This macro will change global ptr of KernelGlobals via name matching. */
|
||||
# define KERNEL_DATA_ARRAY(type, name) \
|
||||
else if (#name == matched_name) \
|
||||
{ \
|
||||
globals->__##name = (type *)memory_device_pointer; \
|
||||
return; \
|
||||
}
|
||||
if (false) {
|
||||
}
|
||||
else if ("integrator_state" == matched_name) {
|
||||
globals->integrator_state = (IntegratorStateGPU *)memory_device_pointer;
|
||||
return;
|
||||
}
|
||||
KERNEL_DATA_ARRAY(KernelData, data)
|
||||
# include "kernel/data_arrays.h"
|
||||
else
|
||||
{
|
||||
std::cerr << "Can't found global/constant memory with name \"" << matched_name << "\"!"
|
||||
<< std::endl;
|
||||
assert(false);
|
||||
}
|
||||
# undef KERNEL_DATA_ARRAY
|
||||
}
|
||||
|
||||
bool OneapiDevice::enqueue_kernel(KernelContext *kernel_context,
|
||||
int kernel,
|
||||
size_t global_size,
|
||||
void **args)
|
||||
{
|
||||
return oneapi_enqueue_kernel(kernel_context, kernel, global_size, args);
|
||||
}
|
||||
|
||||
/* Compute-runtime (ie. NEO) version is what gets returned by sycl/L0 on Windows
|
||||
* since Windows driver 101.3268. */
|
||||
/* The same min compute-runtime version is currently required across Windows and Linux.
|
||||
* For Windows driver 101.3430, compute-runtime version is 23904. */
|
||||
static const int lowest_supported_driver_version_win = 1013430;
|
||||
static const int lowest_supported_driver_version_neo = 23904;
|
||||
|
||||
int OneapiDevice::parse_driver_build_version(const sycl::device &device)
|
||||
{
|
||||
const std::string &driver_version = device.get_info<sycl::info::device::driver_version>();
|
||||
int driver_build_version = 0;
|
||||
|
||||
size_t second_dot_position = driver_version.find('.', driver_version.find('.') + 1);
|
||||
if (second_dot_position == std::string::npos) {
|
||||
std::cerr << "Unable to parse unknown Intel GPU driver version \"" << driver_version
|
||||
<< "\" does not match xx.xx.xxxxx (Linux), x.x.xxxx (L0),"
|
||||
<< " xx.xx.xxx.xxxx (Windows) for device \""
|
||||
<< device.get_info<sycl::info::device::name>() << "\"." << std::endl;
|
||||
}
|
||||
else {
|
||||
try {
|
||||
size_t third_dot_position = driver_version.find('.', second_dot_position + 1);
|
||||
if (third_dot_position != std::string::npos) {
|
||||
const std::string &third_number_substr = driver_version.substr(
|
||||
second_dot_position + 1, third_dot_position - second_dot_position - 1);
|
||||
const std::string &forth_number_substr = driver_version.substr(third_dot_position + 1);
|
||||
if (third_number_substr.length() == 3 && forth_number_substr.length() == 4)
|
||||
driver_build_version = std::stoi(third_number_substr) * 10000 +
|
||||
std::stoi(forth_number_substr);
|
||||
}
|
||||
else {
|
||||
const std::string &third_number_substr = driver_version.substr(second_dot_position + 1);
|
||||
driver_build_version = std::stoi(third_number_substr);
|
||||
}
|
||||
}
|
||||
catch (std::invalid_argument &) {
|
||||
std::cerr << "Unable to parse unknown Intel GPU driver version \"" << driver_version
|
||||
<< "\" does not match xx.xx.xxxxx (Linux), x.x.xxxx (L0),"
|
||||
<< " xx.xx.xxx.xxxx (Windows) for device \""
|
||||
<< device.get_info<sycl::info::device::name>() << "\"." << std::endl;
|
||||
}
|
||||
}
|
||||
|
||||
return driver_build_version;
|
||||
}
|
||||
|
||||
std::vector<sycl::device> OneapiDevice::available_devices()
|
||||
{
|
||||
bool allow_all_devices = false;
|
||||
if (getenv("CYCLES_ONEAPI_ALL_DEVICES") != nullptr)
|
||||
allow_all_devices = true;
|
||||
|
||||
/* Host device is useful only for debugging at the moment
|
||||
* so we hide this device with default build settings. */
|
||||
# ifdef WITH_ONEAPI_SYCL_HOST_ENABLED
|
||||
bool allow_host = true;
|
||||
# else
|
||||
bool allow_host = false;
|
||||
# endif
|
||||
|
||||
const std::vector<sycl::platform> &oneapi_platforms = sycl::platform::get_platforms();
|
||||
|
||||
std::vector<sycl::device> available_devices;
|
||||
for (const sycl::platform &platform : oneapi_platforms) {
|
||||
/* ignore OpenCL platforms to avoid using the same devices through both Level-Zero and OpenCL.
|
||||
*/
|
||||
if (platform.get_backend() == sycl::backend::opencl) {
|
||||
continue;
|
||||
}
|
||||
|
||||
const std::vector<sycl::device> &oneapi_devices =
|
||||
(allow_all_devices || allow_host) ? platform.get_devices(sycl::info::device_type::all) :
|
||||
platform.get_devices(sycl::info::device_type::gpu);
|
||||
|
||||
for (const sycl::device &device : oneapi_devices) {
|
||||
if (allow_all_devices) {
|
||||
/* still filter out host device if build doesn't support it. */
|
||||
if (allow_host || !device.is_host()) {
|
||||
available_devices.push_back(device);
|
||||
}
|
||||
}
|
||||
else {
|
||||
bool filter_out = false;
|
||||
|
||||
/* For now we support all Intel(R) Arc(TM) devices and likely any future GPU,
|
||||
* assuming they have either more than 96 Execution Units or not 7 threads per EU.
|
||||
* Official support can be broaden to older and smaller GPUs once ready. */
|
||||
if (device.is_gpu() && platform.get_backend() == sycl::backend::ext_oneapi_level_zero) {
|
||||
/* Filtered-out defaults in-case these values aren't available through too old L0
|
||||
* runtime. */
|
||||
int number_of_eus = 96;
|
||||
int threads_per_eu = 7;
|
||||
if (device.has(sycl::aspect::ext_intel_gpu_eu_count)) {
|
||||
number_of_eus = device.get_info<sycl::info::device::ext_intel_gpu_eu_count>();
|
||||
}
|
||||
if (device.has(sycl::aspect::ext_intel_gpu_hw_threads_per_eu)) {
|
||||
threads_per_eu =
|
||||
device.get_info<sycl::info::device::ext_intel_gpu_hw_threads_per_eu>();
|
||||
}
|
||||
/* This filters out all Level-Zero supported GPUs from older generation than Arc. */
|
||||
if (number_of_eus <= 96 && threads_per_eu == 7) {
|
||||
filter_out = true;
|
||||
}
|
||||
/* if not already filtered out, check driver version. */
|
||||
if (!filter_out) {
|
||||
int driver_build_version = parse_driver_build_version(device);
|
||||
if ((driver_build_version > 100000 &&
|
||||
driver_build_version < lowest_supported_driver_version_win) ||
|
||||
driver_build_version < lowest_supported_driver_version_neo) {
|
||||
filter_out = true;
|
||||
}
|
||||
}
|
||||
}
|
||||
else if (!allow_host && device.is_host()) {
|
||||
filter_out = true;
|
||||
}
|
||||
else if (!allow_all_devices) {
|
||||
filter_out = true;
|
||||
}
|
||||
|
||||
if (!filter_out) {
|
||||
available_devices.push_back(device);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
return available_devices;
|
||||
}
|
||||
|
||||
char *OneapiDevice::device_capabilities()
|
||||
{
|
||||
std::stringstream capabilities;
|
||||
|
||||
const std::vector<sycl::device> &oneapi_devices = available_devices();
|
||||
for (const sycl::device &device : oneapi_devices) {
|
||||
const std::string &name = device.get_info<sycl::info::device::name>();
|
||||
|
||||
capabilities << std::string("\t") << name << "\n";
|
||||
# define WRITE_ATTR(attribute_name, attribute_variable) \
|
||||
capabilities << "\t\tsycl::info::device::" #attribute_name "\t\t\t" << attribute_variable \
|
||||
<< "\n";
|
||||
# define GET_NUM_ATTR(attribute) \
|
||||
{ \
|
||||
size_t attribute = (size_t)device.get_info<sycl::info::device ::attribute>(); \
|
||||
capabilities << "\t\tsycl::info::device::" #attribute "\t\t\t" << attribute << "\n"; \
|
||||
}
|
||||
|
||||
GET_NUM_ATTR(vendor_id)
|
||||
GET_NUM_ATTR(max_compute_units)
|
||||
GET_NUM_ATTR(max_work_item_dimensions)
|
||||
|
||||
sycl::id<3> max_work_item_sizes =
|
||||
device.get_info<sycl::info::device::max_work_item_sizes<3>>();
|
||||
WRITE_ATTR("max_work_item_sizes_dim0", ((size_t)max_work_item_sizes.get(0)))
|
||||
WRITE_ATTR("max_work_item_sizes_dim1", ((size_t)max_work_item_sizes.get(1)))
|
||||
WRITE_ATTR("max_work_item_sizes_dim2", ((size_t)max_work_item_sizes.get(2)))
|
||||
|
||||
GET_NUM_ATTR(max_work_group_size)
|
||||
GET_NUM_ATTR(max_num_sub_groups)
|
||||
GET_NUM_ATTR(sub_group_independent_forward_progress)
|
||||
|
||||
GET_NUM_ATTR(preferred_vector_width_char)
|
||||
GET_NUM_ATTR(preferred_vector_width_short)
|
||||
GET_NUM_ATTR(preferred_vector_width_int)
|
||||
GET_NUM_ATTR(preferred_vector_width_long)
|
||||
GET_NUM_ATTR(preferred_vector_width_float)
|
||||
GET_NUM_ATTR(preferred_vector_width_double)
|
||||
GET_NUM_ATTR(preferred_vector_width_half)
|
||||
|
||||
GET_NUM_ATTR(native_vector_width_char)
|
||||
GET_NUM_ATTR(native_vector_width_short)
|
||||
GET_NUM_ATTR(native_vector_width_int)
|
||||
GET_NUM_ATTR(native_vector_width_long)
|
||||
GET_NUM_ATTR(native_vector_width_float)
|
||||
GET_NUM_ATTR(native_vector_width_double)
|
||||
GET_NUM_ATTR(native_vector_width_half)
|
||||
|
||||
size_t max_clock_frequency =
|
||||
(size_t)(device.is_host() ? (size_t)0 :
|
||||
device.get_info<sycl::info::device::max_clock_frequency>());
|
||||
WRITE_ATTR("max_clock_frequency", max_clock_frequency)
|
||||
|
||||
GET_NUM_ATTR(address_bits)
|
||||
GET_NUM_ATTR(max_mem_alloc_size)
|
||||
|
||||
/* NOTE(@nsirgien): Implementation doesn't use image support as bindless images aren't
|
||||
* supported so we always return false, even if device supports HW texture usage acceleration.
|
||||
*/
|
||||
bool image_support = false;
|
||||
WRITE_ATTR("image_support", (size_t)image_support)
|
||||
|
||||
GET_NUM_ATTR(max_parameter_size)
|
||||
GET_NUM_ATTR(mem_base_addr_align)
|
||||
GET_NUM_ATTR(global_mem_size)
|
||||
GET_NUM_ATTR(local_mem_size)
|
||||
GET_NUM_ATTR(error_correction_support)
|
||||
GET_NUM_ATTR(profiling_timer_resolution)
|
||||
GET_NUM_ATTR(is_available)
|
||||
|
||||
# undef GET_NUM_ATTR
|
||||
# undef WRITE_ATTR
|
||||
capabilities << "\n";
|
||||
}
|
||||
|
||||
return ::strdup(capabilities.str().c_str());
|
||||
}
|
||||
|
||||
void OneapiDevice::iterate_devices(OneAPIDeviceIteratorCallback cb, void *user_ptr)
|
||||
{
|
||||
int num = 0;
|
||||
std::vector<sycl::device> devices = OneapiDevice::available_devices();
|
||||
for (sycl::device &device : devices) {
|
||||
const std::string &platform_name =
|
||||
device.get_platform().get_info<sycl::info::platform::name>();
|
||||
std::string name = device.get_info<sycl::info::device::name>();
|
||||
std::string id = "ONEAPI_" + platform_name + "_" + name;
|
||||
if (device.has(sycl::aspect::ext_intel_pci_address)) {
|
||||
id.append("_" + device.get_info<sycl::info::device::ext_intel_pci_address>());
|
||||
}
|
||||
(cb)(id.c_str(), name.c_str(), num, user_ptr);
|
||||
num++;
|
||||
}
|
||||
}
|
||||
|
||||
size_t OneapiDevice::get_memcapacity()
|
||||
{
|
||||
return reinterpret_cast<sycl::queue *>(device_queue_)
|
||||
->get_device()
|
||||
.get_info<sycl::info::device::global_mem_size>();
|
||||
}
|
||||
|
||||
int OneapiDevice::get_num_multiprocessors()
|
||||
{
|
||||
const sycl::device &device = reinterpret_cast<sycl::queue *>(device_queue_)->get_device();
|
||||
if (device.has(sycl::aspect::ext_intel_gpu_eu_count)) {
|
||||
return device.get_info<sycl::info::device::ext_intel_gpu_eu_count>();
|
||||
}
|
||||
else
|
||||
return 0;
|
||||
}
|
||||
|
||||
int OneapiDevice::get_max_num_threads_per_multiprocessor()
|
||||
{
|
||||
const sycl::device &device = reinterpret_cast<sycl::queue *>(device_queue_)->get_device();
|
||||
if (device.has(sycl::aspect::ext_intel_gpu_eu_simd_width) &&
|
||||
device.has(sycl::aspect::ext_intel_gpu_hw_threads_per_eu)) {
|
||||
return device.get_info<sycl::info::device::ext_intel_gpu_eu_simd_width>() *
|
||||
device.get_info<sycl::info::device::ext_intel_gpu_hw_threads_per_eu>();
|
||||
}
|
||||
else
|
||||
return 0;
|
||||
}
|
||||
|
||||
CCL_NAMESPACE_END
|
||||
|
|
|
@ -3,9 +3,12 @@
|
|||
|
||||
#ifdef WITH_ONEAPI
|
||||
|
||||
# include <CL/sycl.hpp>
|
||||
|
||||
# include "device/device.h"
|
||||
# include "device/oneapi/device.h"
|
||||
# include "device/oneapi/queue.h"
|
||||
# include "kernel/device/oneapi/kernel.h"
|
||||
|
||||
# include "util/map.h"
|
||||
|
||||
|
@ -13,6 +16,11 @@ CCL_NAMESPACE_BEGIN
|
|||
|
||||
class DeviceQueue;
|
||||
|
||||
typedef void (*OneAPIDeviceIteratorCallback)(const char *id,
|
||||
const char *name,
|
||||
int num,
|
||||
void *user_ptr);
|
||||
|
||||
class OneapiDevice : public Device {
|
||||
private:
|
||||
SyclQueue *device_queue_;
|
||||
|
@ -25,16 +33,12 @@ class OneapiDevice : public Device {
|
|||
void *kg_memory_device_;
|
||||
size_t kg_memory_size_ = (size_t)0;
|
||||
size_t max_memory_on_device_ = (size_t)0;
|
||||
OneAPIDLLInterface oneapi_dll_;
|
||||
std::string oneapi_error_string_;
|
||||
|
||||
public:
|
||||
virtual BVHLayoutMask get_bvh_layout_mask() const override;
|
||||
|
||||
OneapiDevice(const DeviceInfo &info,
|
||||
OneAPIDLLInterface &oneapi_dll_object,
|
||||
Stats &stats,
|
||||
Profiler &profiler);
|
||||
OneapiDevice(const DeviceInfo &info, Stats &stats, Profiler &profiler);
|
||||
|
||||
virtual ~OneapiDevice();
|
||||
|
||||
|
@ -50,12 +54,8 @@ class OneapiDevice : public Device {
|
|||
|
||||
void generic_free(device_memory &mem);
|
||||
|
||||
SyclQueue *sycl_queue();
|
||||
|
||||
string oneapi_error_message();
|
||||
|
||||
OneAPIDLLInterface oneapi_dll_object();
|
||||
|
||||
void *kernel_globals_device_pointer();
|
||||
|
||||
void mem_alloc(device_memory &mem) override;
|
||||
|
@ -90,13 +90,37 @@ class OneapiDevice : public Device {
|
|||
|
||||
virtual unique_ptr<DeviceQueue> gpu_queue_create() override;
|
||||
|
||||
int get_num_multiprocessors();
|
||||
int get_max_num_threads_per_multiprocessor();
|
||||
|
||||
/* NOTE(@nsirgien): Create this methods to avoid some compilation problems on Windows with host
|
||||
* side compilation (MSVC). */
|
||||
void *usm_aligned_alloc_host(size_t memory_size, size_t alignment);
|
||||
void usm_free(void *usm_ptr);
|
||||
|
||||
static std::vector<sycl::device> available_devices();
|
||||
static char *device_capabilities();
|
||||
static int parse_driver_build_version(const sycl::device &device);
|
||||
static void iterate_devices(OneAPIDeviceIteratorCallback cb, void *user_ptr);
|
||||
|
||||
size_t get_memcapacity();
|
||||
int get_num_multiprocessors();
|
||||
int get_max_num_threads_per_multiprocessor();
|
||||
bool queue_synchronize(SyclQueue *queue);
|
||||
bool kernel_globals_size(SyclQueue *queue, size_t &kernel_global_size);
|
||||
void set_global_memory(SyclQueue *queue,
|
||||
void *kernel_globals,
|
||||
const char *memory_name,
|
||||
void *memory_device_pointer);
|
||||
bool enqueue_kernel(KernelContext *kernel_context, int kernel, size_t global_size, void **args);
|
||||
SyclQueue *sycl_queue();
|
||||
|
||||
protected:
|
||||
void check_usm(SyclQueue *queue, const void *usm_ptr, bool allow_host);
|
||||
bool create_queue(SyclQueue *&external_queue, int device_index);
|
||||
void free_queue(SyclQueue *queue);
|
||||
void *usm_aligned_alloc_host(SyclQueue *queue, size_t memory_size, size_t alignment);
|
||||
void *usm_alloc_device(SyclQueue *queue, size_t memory_size);
|
||||
void usm_free(SyclQueue *queue, void *usm_ptr);
|
||||
bool usm_memcpy(SyclQueue *queue, void *dest, void *src, size_t num_bytes);
|
||||
bool usm_memset(SyclQueue *queue, void *usm_ptr, unsigned char value, size_t num_bytes);
|
||||
};
|
||||
|
||||
CCL_NAMESPACE_END
|
||||
|
|
|
@ -1,17 +0,0 @@
|
|||
/* SPDX-License-Identifier: Apache-2.0
|
||||
* Copyright 2011-2022 Blender Foundation */
|
||||
|
||||
#pragma once
|
||||
|
||||
/* Include kernel header to get access to SYCL-specific types, like SyclQueue and
|
||||
* OneAPIDeviceIteratorCallback. */
|
||||
#include "kernel/device/oneapi/kernel.h"
|
||||
|
||||
#ifdef WITH_ONEAPI
|
||||
struct OneAPIDLLInterface {
|
||||
# define DLL_INTERFACE_CALL(function, return_type, ...) \
|
||||
return_type (*function)(__VA_ARGS__) = nullptr;
|
||||
# include "kernel/device/oneapi/dll_interface_template.h"
|
||||
# undef DLL_INTERFACE_CALL
|
||||
};
|
||||
#endif
|
|
@ -22,10 +22,7 @@ struct KernelExecutionInfo {
|
|||
/* OneapiDeviceQueue */
|
||||
|
||||
OneapiDeviceQueue::OneapiDeviceQueue(OneapiDevice *device)
|
||||
: DeviceQueue(device),
|
||||
oneapi_device_(device),
|
||||
oneapi_dll_(device->oneapi_dll_object()),
|
||||
kernel_context_(nullptr)
|
||||
: DeviceQueue(device), oneapi_device_(device), kernel_context_(nullptr)
|
||||
{
|
||||
}
|
||||
|
||||
|
@ -81,14 +78,14 @@ bool OneapiDeviceQueue::enqueue(DeviceKernel kernel,
|
|||
assert(signed_kernel_work_size >= 0);
|
||||
size_t kernel_work_size = (size_t)signed_kernel_work_size;
|
||||
|
||||
size_t kernel_local_size = oneapi_dll_.oneapi_kernel_preferred_local_size(
|
||||
size_t kernel_local_size = oneapi_kernel_preferred_local_size(
|
||||
kernel_context_->queue, (::DeviceKernel)kernel, kernel_work_size);
|
||||
size_t uniformed_kernel_work_size = round_up(kernel_work_size, kernel_local_size);
|
||||
|
||||
assert(kernel_context_);
|
||||
|
||||
/* Call the oneAPI kernel DLL to launch the requested kernel. */
|
||||
bool is_finished_ok = oneapi_dll_.oneapi_enqueue_kernel(
|
||||
bool is_finished_ok = oneapi_device_->enqueue_kernel(
|
||||
kernel_context_, kernel, uniformed_kernel_work_size, args);
|
||||
|
||||
if (is_finished_ok == false) {
|
||||
|
@ -108,7 +105,7 @@ bool OneapiDeviceQueue::synchronize()
|
|||
return false;
|
||||
}
|
||||
|
||||
bool is_finished_ok = oneapi_dll_.oneapi_queue_synchronize(oneapi_device_->sycl_queue());
|
||||
bool is_finished_ok = oneapi_device_->queue_synchronize(oneapi_device_->sycl_queue());
|
||||
if (is_finished_ok == false)
|
||||
oneapi_device_->set_error("oneAPI unknown kernel execution error: got runtime exception \"" +
|
||||
oneapi_device_->oneapi_error_message() + "\"");
|
||||
|
|
|
@ -10,7 +10,7 @@
|
|||
# include "device/queue.h"
|
||||
|
||||
# include "device/oneapi/device.h"
|
||||
# include "device/oneapi/dll_interface.h"
|
||||
# include "kernel/device/oneapi/kernel.h"
|
||||
|
||||
CCL_NAMESPACE_BEGIN
|
||||
|
||||
|
@ -41,9 +41,7 @@ class OneapiDeviceQueue : public DeviceQueue {
|
|||
|
||||
protected:
|
||||
OneapiDevice *oneapi_device_;
|
||||
OneAPIDLLInterface oneapi_dll_;
|
||||
KernelContext *kernel_context_;
|
||||
bool with_kernel_statistics_;
|
||||
};
|
||||
|
||||
CCL_NAMESPACE_END
|
||||
|
|
|
@ -716,7 +716,7 @@ if(WITH_CYCLES_DEVICE_ONEAPI)
|
|||
if(WIN32)
|
||||
set(cycles_kernel_oneapi_lib ${CMAKE_CURRENT_BINARY_DIR}/cycles_kernel_oneapi.dll)
|
||||
else()
|
||||
set(cycles_kernel_oneapi_lib ${CMAKE_CURRENT_BINARY_DIR}/cycles_kernel_oneapi.so)
|
||||
set(cycles_kernel_oneapi_lib ${CMAKE_CURRENT_BINARY_DIR}/libcycles_kernel_oneapi.so)
|
||||
endif()
|
||||
|
||||
set(cycles_oneapi_kernel_sources
|
||||
|
@ -758,7 +758,6 @@ if(WITH_CYCLES_DEVICE_ONEAPI)
|
|||
${SYCL_CPP_FLAGS}
|
||||
)
|
||||
|
||||
|
||||
if (WITH_CYCLES_ONEAPI_SYCL_HOST_ENABLED)
|
||||
list(APPEND sycl_compiler_flags -DWITH_ONEAPI_SYCL_HOST_ENABLED)
|
||||
endif()
|
||||
|
@ -815,6 +814,7 @@ if(WITH_CYCLES_DEVICE_ONEAPI)
|
|||
|
||||
if(WIN32)
|
||||
list(APPEND sycl_compiler_flags
|
||||
-fuse-ld=link
|
||||
-fms-extensions
|
||||
-fms-compatibility
|
||||
-D_WINDLL
|
||||
|
@ -825,36 +825,43 @@ if(WITH_CYCLES_DEVICE_ONEAPI)
|
|||
-D_CRT_SECURE_NO_DEPRECATE
|
||||
-DONEAPI_EXPORT)
|
||||
|
||||
if(sycl_compiler_compiler_name MATCHES "dpcpp")
|
||||
# The oneAPI distribution calls the compiler "dpcpp" and comes with a script that sets environment variables.
|
||||
add_custom_command(
|
||||
OUTPUT ${cycles_kernel_oneapi_lib}
|
||||
COMMAND "${sycl_compiler_root}/../../env/vars.bat"
|
||||
COMMAND ${SYCL_COMPILER} $<$<CONFIG:Debug>:-g>$<$<CONFIG:RelWithDebInfo>:-g> ${sycl_compiler_flags}
|
||||
DEPENDS ${cycles_oneapi_kernel_sources})
|
||||
string(REPLACE /Redist/ /Tools/ MSVC_TOOLS_DIR ${MSVC_REDIST_DIR})
|
||||
if(NOT CMAKE_VS_WINDOWS_TARGET_PLATFORM_VERSION) # case for Ninja on Windows
|
||||
get_filename_component(cmake_mt_dir ${CMAKE_MT} DIRECTORY)
|
||||
string(REPLACE /bin/ /Lib/ WINDOWS_KIT_DIR ${cmake_mt_dir})
|
||||
get_filename_component(WINDOWS_KIT_DIR "${WINDOWS_KIT_DIR}/../" ABSOLUTE)
|
||||
else()
|
||||
# The open source SYCL compiler just goes by clang++ and does not have such a script.
|
||||
# Set the variables manually.
|
||||
string(REPLACE /Redist/ /Tools/ MSVC_TOOLS_DIR ${MSVC_REDIST_DIR})
|
||||
if(NOT CMAKE_VS_WINDOWS_TARGET_PLATFORM_VERSION) # case for Ninja on Windows
|
||||
get_filename_component(cmake_mt_dir ${CMAKE_MT} DIRECTORY)
|
||||
string(REPLACE /bin/ /Lib/ WINDOWS_KIT_DIR ${cmake_mt_dir})
|
||||
get_filename_component(WINDOWS_KIT_DIR "${WINDOWS_KIT_DIR}/../" ABSOLUTE)
|
||||
else()
|
||||
set(WINDOWS_KIT_DIR ${WINDOWS_KITS_DIR}/Lib/${CMAKE_VS_WINDOWS_TARGET_PLATFORM_VERSION})
|
||||
endif()
|
||||
list(APPEND sycl_compiler_flags
|
||||
-L "${MSVC_TOOLS_DIR}/lib/x64"
|
||||
-L "${WINDOWS_KIT_DIR}/um/x64"
|
||||
-L "${WINDOWS_KIT_DIR}/ucrt/x64")
|
||||
add_custom_command(
|
||||
OUTPUT ${cycles_kernel_oneapi_lib}
|
||||
COMMAND ${CMAKE_COMMAND} -E env
|
||||
"LIB=${sycl_compiler_root}/../lib" # for compiler to find sycl.lib
|
||||
"PATH=${OCLOC_INSTALL_DIR};${sycl_compiler_root}"
|
||||
${SYCL_COMPILER} $<$<CONFIG:Debug>:-g>$<$<CONFIG:RelWithDebInfo>:-g> ${sycl_compiler_flags}
|
||||
DEPENDS ${cycles_oneapi_kernel_sources})
|
||||
set(WINDOWS_KIT_DIR ${WINDOWS_KITS_DIR}/Lib/${CMAKE_VS_WINDOWS_TARGET_PLATFORM_VERSION})
|
||||
endif()
|
||||
list(APPEND sycl_compiler_flags
|
||||
-L "${MSVC_TOOLS_DIR}/lib/x64"
|
||||
-L "${WINDOWS_KIT_DIR}/um/x64"
|
||||
-L "${WINDOWS_KIT_DIR}/ucrt/x64")
|
||||
|
||||
set(sycl_compiler_flags_Release ${sycl_compiler_flags})
|
||||
set(sycl_compiler_flags_Debug ${sycl_compiler_flags})
|
||||
set(sycl_compiler_flags_RelWithDebInfo ${sycl_compiler_flags})
|
||||
set(sycl_compiler_flags_MinSizeRel ${sycl_compiler_flags})
|
||||
list(APPEND sycl_compiler_flags_RelWithDebInfo -g)
|
||||
get_filename_component(sycl_library_debug_name ${SYCL_LIBRARY_DEBUG} NAME_WE)
|
||||
list(APPEND sycl_compiler_flags_Debug
|
||||
-g
|
||||
-D_DEBUG
|
||||
-nostdlib -Xclang --dependent-lib=msvcrtd
|
||||
-Xclang --dependent-lib=${sycl_library_debug_name})
|
||||
|
||||
add_custom_command(
|
||||
OUTPUT ${cycles_kernel_oneapi_lib}
|
||||
COMMAND ${CMAKE_COMMAND} -E env
|
||||
"LIB=${sycl_compiler_root}/../lib" # for compiler to find sycl.lib
|
||||
"PATH=${OCLOC_INSTALL_DIR}\;${sycl_compiler_root}"
|
||||
${SYCL_COMPILER}
|
||||
"$<$<CONFIG:Release>:${sycl_compiler_flags_Release}>"
|
||||
"$<$<CONFIG:RelWithDebInfo>:${sycl_compiler_flags_RelWithDebInfo}>"
|
||||
"$<$<CONFIG:Debug>:${sycl_compiler_flags_Debug}>"
|
||||
"$<$<CONFIG:MinSizeRel>:${sycl_compiler_flags_Release}>"
|
||||
COMMAND_EXPAND_LISTS
|
||||
DEPENDS ${cycles_oneapi_kernel_sources})
|
||||
else()
|
||||
list(APPEND sycl_compiler_flags -fPIC)
|
||||
|
||||
|
@ -866,55 +873,36 @@ if(WITH_CYCLES_DEVICE_ONEAPI)
|
|||
# libpi_level_zero.so can be placed next to it and get found.
|
||||
list(APPEND sycl_compiler_flags -Wl,-rpath,'$$ORIGIN')
|
||||
|
||||
# The oneAPI distribution calls the compiler "dpcpp" and comes with a script that sets environment variables.
|
||||
if(sycl_compiler_compiler_name MATCHES "dpcpp")
|
||||
add_custom_command(
|
||||
OUTPUT ${cycles_kernel_oneapi_lib}
|
||||
COMMAND bash -c \"source ${sycl_compiler_root}/../../env/vars.sh&&${SYCL_COMPILER} $<$<CONFIG:Debug>:-g>$<$<CONFIG:RelWithDebInfo>:-g> ${sycl_compiler_flags}\"
|
||||
DEPENDS ${cycles_oneapi_kernel_sources})
|
||||
else()
|
||||
# The open source SYCL compiler just goes by clang++ and does not have such a script.
|
||||
# Set the variables manually.
|
||||
if(NOT IGC_INSTALL_DIR)
|
||||
get_filename_component(IGC_INSTALL_DIR "${sycl_compiler_root}/../lib/igc" ABSOLUTE)
|
||||
endif()
|
||||
add_custom_command(
|
||||
OUTPUT ${cycles_kernel_oneapi_lib}
|
||||
COMMAND ${CMAKE_COMMAND} -E env
|
||||
"LD_LIBRARY_PATH=${sycl_compiler_root}/../lib:${OCLOC_INSTALL_DIR}/lib:${IGC_INSTALL_DIR}/lib"
|
||||
"PATH=${OCLOC_INSTALL_DIR}/bin:${sycl_compiler_root}:$ENV{PATH}" # env PATH is for compiler to find ld
|
||||
${SYCL_COMPILER} $<$<CONFIG:Debug>:-g>$<$<CONFIG:RelWithDebInfo>:-g> ${sycl_compiler_flags}
|
||||
DEPENDS ${cycles_oneapi_kernel_sources})
|
||||
if(NOT IGC_INSTALL_DIR)
|
||||
get_filename_component(IGC_INSTALL_DIR "${sycl_compiler_root}/../lib/igc" ABSOLUTE)
|
||||
endif()
|
||||
add_custom_command(
|
||||
OUTPUT ${cycles_kernel_oneapi_lib}
|
||||
COMMAND ${CMAKE_COMMAND} -E env
|
||||
"LD_LIBRARY_PATH=${sycl_compiler_root}/../lib:${OCLOC_INSTALL_DIR}/lib:${IGC_INSTALL_DIR}/lib"
|
||||
"PATH=${OCLOC_INSTALL_DIR}/bin:${sycl_compiler_root}:$ENV{PATH}" # env PATH is for compiler to find ld
|
||||
${SYCL_COMPILER} $<$<CONFIG:Debug>:-g>$<$<CONFIG:RelWithDebInfo>:-g> ${sycl_compiler_flags}
|
||||
DEPENDS ${cycles_oneapi_kernel_sources})
|
||||
endif()
|
||||
|
||||
if(NOT WITH_BLENDER)
|
||||
# For the Cycles standalone put libraries next to the Cycles application.
|
||||
set(cycles_oneapi_target_path ${CYCLES_INSTALL_PATH})
|
||||
else()
|
||||
# For Blender put the libraries next to the Blender executable.
|
||||
#
|
||||
# Note that the installation path in the delayed_install is relative to the versioned folder,
|
||||
# which means we need to go one level up.
|
||||
set(cycles_oneapi_target_path "../")
|
||||
endif()
|
||||
|
||||
# install dynamic libraries required at runtime
|
||||
if(WIN32)
|
||||
set(SYCL_RUNTIME_DEPENDENCIES
|
||||
sycl.dll
|
||||
pi_level_zero.dll
|
||||
)
|
||||
if(NOT WITH_BLENDER)
|
||||
# For the Cycles standalone put libraries next to the Cycles application.
|
||||
delayed_install("${sycl_compiler_root}" "${SYCL_RUNTIME_DEPENDENCIES}" ${CYCLES_INSTALL_PATH})
|
||||
else()
|
||||
# For Blender put the libraries next to the Blender executable.
|
||||
#
|
||||
# Note that the installation path in the delayed_install is relative to the versioned folder,
|
||||
# which means we need to go one level up.
|
||||
delayed_install("${sycl_compiler_root}" "${SYCL_RUNTIME_DEPENDENCIES}" "../")
|
||||
endif()
|
||||
delayed_install("" "${cycles_kernel_oneapi_lib}" ${cycles_oneapi_target_path})
|
||||
elseif(UNIX AND NOT APPLE)
|
||||
file(GLOB SYCL_RUNTIME_DEPENDENCIES
|
||||
${sycl_compiler_root}/../lib/libsycl.so
|
||||
${sycl_compiler_root}/../lib/libsycl.so.[0-9]
|
||||
${sycl_compiler_root}/../lib/libsycl.so.[0-9].[0-9].[0-9]-[0-9]
|
||||
)
|
||||
list(APPEND SYCL_RUNTIME_DEPENDENCIES ${sycl_compiler_root}/../lib/libpi_level_zero.so)
|
||||
delayed_install("" "${SYCL_RUNTIME_DEPENDENCIES}" ${CYCLES_INSTALL_PATH}/lib)
|
||||
delayed_install("" "${cycles_kernel_oneapi_lib}" ${cycles_oneapi_target_path}/lib)
|
||||
endif()
|
||||
|
||||
delayed_install("${CMAKE_CURRENT_BINARY_DIR}" "${cycles_kernel_oneapi_lib}" ${CYCLES_INSTALL_PATH}/lib)
|
||||
add_custom_target(cycles_kernel_oneapi ALL DEPENDS ${cycles_kernel_oneapi_lib})
|
||||
endif()
|
||||
|
||||
|
|
|
@ -1,54 +0,0 @@
|
|||
/* SPDX-License-Identifier: Apache-2.0
|
||||
* Copyright 2022 Intel Corporation */
|
||||
|
||||
/* device_capabilities() returns a C string that must be free'd with oneapi_free(). */
|
||||
DLL_INTERFACE_CALL(oneapi_device_capabilities, char *)
|
||||
DLL_INTERFACE_CALL(oneapi_free, void, void *)
|
||||
DLL_INTERFACE_CALL(oneapi_get_memcapacity, size_t, SyclQueue *queue)
|
||||
|
||||
DLL_INTERFACE_CALL(oneapi_get_num_multiprocessors, int, SyclQueue *queue)
|
||||
DLL_INTERFACE_CALL(oneapi_get_max_num_threads_per_multiprocessor, int, SyclQueue *queue)
|
||||
DLL_INTERFACE_CALL(oneapi_iterate_devices, void, OneAPIDeviceIteratorCallback cb, void *user_ptr)
|
||||
DLL_INTERFACE_CALL(oneapi_set_error_cb, void, OneAPIErrorCallback, void *user_ptr)
|
||||
|
||||
DLL_INTERFACE_CALL(oneapi_create_queue, bool, SyclQueue *&external_queue, int device_index)
|
||||
DLL_INTERFACE_CALL(oneapi_free_queue, void, SyclQueue *queue)
|
||||
DLL_INTERFACE_CALL(
|
||||
oneapi_usm_aligned_alloc_host, void *, SyclQueue *queue, size_t memory_size, size_t alignment)
|
||||
DLL_INTERFACE_CALL(oneapi_usm_alloc_device, void *, SyclQueue *queue, size_t memory_size)
|
||||
DLL_INTERFACE_CALL(oneapi_usm_free, void, SyclQueue *queue, void *usm_ptr)
|
||||
|
||||
DLL_INTERFACE_CALL(
|
||||
oneapi_usm_memcpy, bool, SyclQueue *queue, void *dest, void *src, size_t num_bytes)
|
||||
DLL_INTERFACE_CALL(oneapi_queue_synchronize, bool, SyclQueue *queue)
|
||||
DLL_INTERFACE_CALL(oneapi_usm_memset,
|
||||
bool,
|
||||
SyclQueue *queue,
|
||||
void *usm_ptr,
|
||||
unsigned char value,
|
||||
size_t num_bytes)
|
||||
|
||||
DLL_INTERFACE_CALL(oneapi_run_test_kernel, bool, SyclQueue *queue)
|
||||
|
||||
/* Operation with Kernel globals structure - map of global/constant allocation - filled before
|
||||
* render/kernel execution As we don't know in cycles `sizeof` this - Cycles will manage just as
|
||||
* pointer. */
|
||||
DLL_INTERFACE_CALL(oneapi_kernel_globals_size, bool, SyclQueue *queue, size_t &kernel_global_size)
|
||||
DLL_INTERFACE_CALL(oneapi_set_global_memory,
|
||||
void,
|
||||
SyclQueue *queue,
|
||||
void *kernel_globals,
|
||||
const char *memory_name,
|
||||
void *memory_device_pointer)
|
||||
|
||||
DLL_INTERFACE_CALL(oneapi_kernel_preferred_local_size,
|
||||
size_t,
|
||||
SyclQueue *queue,
|
||||
const DeviceKernel kernel,
|
||||
const size_t kernel_global_size)
|
||||
DLL_INTERFACE_CALL(oneapi_enqueue_kernel,
|
||||
bool,
|
||||
KernelContext *context,
|
||||
int kernel,
|
||||
size_t global_size,
|
||||
void **args)
|
|
@ -3,7 +3,6 @@
|
|||
|
||||
#ifdef WITH_ONEAPI
|
||||
|
||||
/* clang-format off */
|
||||
# include "kernel.h"
|
||||
# include <iostream>
|
||||
# include <map>
|
||||
|
@ -16,163 +15,16 @@
|
|||
# include "kernel/device/oneapi/kernel_templates.h"
|
||||
|
||||
# include "kernel/device/gpu/kernel.h"
|
||||
/* clang-format on */
|
||||
|
||||
static OneAPIErrorCallback s_error_cb = nullptr;
|
||||
static void *s_error_user_ptr = nullptr;
|
||||
|
||||
static std::vector<sycl::device> oneapi_available_devices();
|
||||
|
||||
void oneapi_set_error_cb(OneAPIErrorCallback cb, void *user_ptr)
|
||||
{
|
||||
s_error_cb = cb;
|
||||
s_error_user_ptr = user_ptr;
|
||||
}
|
||||
|
||||
void oneapi_check_usm(SyclQueue *queue_, const void *usm_ptr, bool allow_host = false)
|
||||
{
|
||||
# ifdef _DEBUG
|
||||
sycl::queue *queue = reinterpret_cast<sycl::queue *>(queue_);
|
||||
sycl::info::device_type device_type =
|
||||
queue->get_device().get_info<sycl::info::device::device_type>();
|
||||
sycl::usm::alloc usm_type = get_pointer_type(usm_ptr, queue->get_context());
|
||||
(void)usm_type;
|
||||
assert(usm_type == sycl::usm::alloc::device ||
|
||||
((device_type == sycl::info::device_type::host ||
|
||||
device_type == sycl::info::device_type::is_cpu || allow_host) &&
|
||||
usm_type == sycl::usm::alloc::host));
|
||||
# endif
|
||||
}
|
||||
|
||||
bool oneapi_create_queue(SyclQueue *&external_queue, int device_index)
|
||||
{
|
||||
bool finished_correct = true;
|
||||
try {
|
||||
std::vector<sycl::device> devices = oneapi_available_devices();
|
||||
if (device_index < 0 || device_index >= devices.size()) {
|
||||
return false;
|
||||
}
|
||||
sycl::queue *created_queue = new sycl::queue(devices[device_index],
|
||||
sycl::property::queue::in_order());
|
||||
external_queue = reinterpret_cast<SyclQueue *>(created_queue);
|
||||
}
|
||||
catch (sycl::exception const &e) {
|
||||
finished_correct = false;
|
||||
if (s_error_cb) {
|
||||
s_error_cb(e.what(), s_error_user_ptr);
|
||||
}
|
||||
}
|
||||
return finished_correct;
|
||||
}
|
||||
|
||||
void oneapi_free_queue(SyclQueue *queue_)
|
||||
{
|
||||
assert(queue_);
|
||||
sycl::queue *queue = reinterpret_cast<sycl::queue *>(queue_);
|
||||
delete queue;
|
||||
}
|
||||
|
||||
void *oneapi_usm_aligned_alloc_host(SyclQueue *queue_, size_t memory_size, size_t alignment)
|
||||
{
|
||||
assert(queue_);
|
||||
sycl::queue *queue = reinterpret_cast<sycl::queue *>(queue_);
|
||||
return sycl::aligned_alloc_host(alignment, memory_size, *queue);
|
||||
}
|
||||
|
||||
void *oneapi_usm_alloc_device(SyclQueue *queue_, size_t memory_size)
|
||||
{
|
||||
assert(queue_);
|
||||
sycl::queue *queue = reinterpret_cast<sycl::queue *>(queue_);
|
||||
return sycl::malloc_device(memory_size, *queue);
|
||||
}
|
||||
|
||||
void oneapi_usm_free(SyclQueue *queue_, void *usm_ptr)
|
||||
{
|
||||
assert(queue_);
|
||||
sycl::queue *queue = reinterpret_cast<sycl::queue *>(queue_);
|
||||
oneapi_check_usm(queue_, usm_ptr, true);
|
||||
sycl::free(usm_ptr, *queue);
|
||||
}
|
||||
|
||||
bool oneapi_usm_memcpy(SyclQueue *queue_, void *dest, void *src, size_t num_bytes)
|
||||
{
|
||||
assert(queue_);
|
||||
sycl::queue *queue = reinterpret_cast<sycl::queue *>(queue_);
|
||||
oneapi_check_usm(queue_, dest, true);
|
||||
oneapi_check_usm(queue_, src, true);
|
||||
sycl::event mem_event = queue->memcpy(dest, src, num_bytes);
|
||||
# ifdef WITH_CYCLES_DEBUG
|
||||
try {
|
||||
/* NOTE(@nsirgien) Waiting on memory operation may give more precise error
|
||||
* messages. Due to impact on occupancy, it makes sense to enable it only during Cycles debug.
|
||||
*/
|
||||
mem_event.wait_and_throw();
|
||||
return true;
|
||||
}
|
||||
catch (sycl::exception const &e) {
|
||||
if (s_error_cb) {
|
||||
s_error_cb(e.what(), s_error_user_ptr);
|
||||
}
|
||||
return false;
|
||||
}
|
||||
# else
|
||||
sycl::usm::alloc dest_type = get_pointer_type(dest, queue->get_context());
|
||||
sycl::usm::alloc src_type = get_pointer_type(src, queue->get_context());
|
||||
bool from_device_to_host = dest_type == sycl::usm::alloc::host &&
|
||||
src_type == sycl::usm::alloc::device;
|
||||
bool host_or_device_memop_with_offset = dest_type == sycl::usm::alloc::unknown ||
|
||||
src_type == sycl::usm::alloc::unknown;
|
||||
/* NOTE(@sirgienko) Host-side blocking wait on this operation is mandatory, otherwise the host
|
||||
* may not wait until the end of the transfer before using the memory.
|
||||
*/
|
||||
if (from_device_to_host || host_or_device_memop_with_offset)
|
||||
mem_event.wait();
|
||||
return true;
|
||||
# endif
|
||||
}
|
||||
|
||||
bool oneapi_usm_memset(SyclQueue *queue_, void *usm_ptr, unsigned char value, size_t num_bytes)
|
||||
{
|
||||
assert(queue_);
|
||||
sycl::queue *queue = reinterpret_cast<sycl::queue *>(queue_);
|
||||
oneapi_check_usm(queue_, usm_ptr, true);
|
||||
sycl::event mem_event = queue->memset(usm_ptr, value, num_bytes);
|
||||
# ifdef WITH_CYCLES_DEBUG
|
||||
try {
|
||||
/* NOTE(@nsirgien) Waiting on memory operation may give more precise error
|
||||
* messages. Due to impact on occupancy, it makes sense to enable it only during Cycles debug.
|
||||
*/
|
||||
mem_event.wait_and_throw();
|
||||
return true;
|
||||
}
|
||||
catch (sycl::exception const &e) {
|
||||
if (s_error_cb) {
|
||||
s_error_cb(e.what(), s_error_user_ptr);
|
||||
}
|
||||
return false;
|
||||
}
|
||||
# else
|
||||
(void)mem_event;
|
||||
return true;
|
||||
# endif
|
||||
}
|
||||
|
||||
bool oneapi_queue_synchronize(SyclQueue *queue_)
|
||||
{
|
||||
assert(queue_);
|
||||
sycl::queue *queue = reinterpret_cast<sycl::queue *>(queue_);
|
||||
try {
|
||||
queue->wait_and_throw();
|
||||
return true;
|
||||
}
|
||||
catch (sycl::exception const &e) {
|
||||
if (s_error_cb) {
|
||||
s_error_cb(e.what(), s_error_user_ptr);
|
||||
}
|
||||
return false;
|
||||
}
|
||||
}
|
||||
|
||||
/* NOTE(@nsirgien): Execution of this simple kernel will check basic functionality and
|
||||
* also trigger runtime compilation of all existing oneAPI kernels */
|
||||
bool oneapi_run_test_kernel(SyclQueue *queue_)
|
||||
|
@ -216,60 +68,13 @@ bool oneapi_run_test_kernel(SyclQueue *queue_)
|
|||
return true;
|
||||
}
|
||||
|
||||
bool oneapi_kernel_globals_size(SyclQueue *queue_, size_t &kernel_global_size)
|
||||
{
|
||||
kernel_global_size = sizeof(KernelGlobalsGPU);
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
void oneapi_set_global_memory(SyclQueue *queue_,
|
||||
void *kernel_globals,
|
||||
const char *memory_name,
|
||||
void *memory_device_pointer)
|
||||
{
|
||||
assert(queue_);
|
||||
assert(kernel_globals);
|
||||
assert(memory_name);
|
||||
assert(memory_device_pointer);
|
||||
KernelGlobalsGPU *globals = (KernelGlobalsGPU *)kernel_globals;
|
||||
oneapi_check_usm(queue_, memory_device_pointer);
|
||||
oneapi_check_usm(queue_, kernel_globals, true);
|
||||
|
||||
std::string matched_name(memory_name);
|
||||
|
||||
/* This macro will change global ptr of KernelGlobals via name matching. */
|
||||
# define KERNEL_DATA_ARRAY(type, name) \
|
||||
else if (#name == matched_name) \
|
||||
{ \
|
||||
globals->__##name = (type *)memory_device_pointer; \
|
||||
return; \
|
||||
}
|
||||
if (false) {
|
||||
}
|
||||
else if ("integrator_state" == matched_name) {
|
||||
globals->integrator_state = (IntegratorStateGPU *)memory_device_pointer;
|
||||
return;
|
||||
}
|
||||
KERNEL_DATA_ARRAY(KernelData, data)
|
||||
# include "kernel/data_arrays.h"
|
||||
else
|
||||
{
|
||||
std::cerr << "Can't found global/constant memory with name \"" << matched_name << "\"!"
|
||||
<< std::endl;
|
||||
assert(false);
|
||||
}
|
||||
# undef KERNEL_DATA_ARRAY
|
||||
}
|
||||
|
||||
/* TODO: Move device information to OneapiDevice initialized on creation and use it. */
|
||||
/* TODO: Move below function to oneapi/queue.cpp. */
|
||||
size_t oneapi_kernel_preferred_local_size(SyclQueue *queue_,
|
||||
size_t oneapi_kernel_preferred_local_size(SyclQueue *queue,
|
||||
const DeviceKernel kernel,
|
||||
const size_t kernel_global_size)
|
||||
{
|
||||
assert(queue_);
|
||||
sycl::queue *queue = reinterpret_cast<sycl::queue *>(queue_);
|
||||
assert(queue);
|
||||
(void)kernel_global_size;
|
||||
const static size_t preferred_work_group_size_intersect_shading = 32;
|
||||
const static size_t preferred_work_group_size_technical = 1024;
|
||||
|
@ -311,8 +116,10 @@ size_t oneapi_kernel_preferred_local_size(SyclQueue *queue_,
|
|||
preferred_work_group_size = 512;
|
||||
}
|
||||
|
||||
const size_t limit_work_group_size =
|
||||
queue->get_device().get_info<sycl::info::device::max_work_group_size>();
|
||||
const size_t limit_work_group_size = reinterpret_cast<sycl::queue *>(queue)
|
||||
->get_device()
|
||||
.get_info<sycl::info::device::max_work_group_size>();
|
||||
|
||||
return std::min(limit_work_group_size, preferred_work_group_size);
|
||||
}
|
||||
|
||||
|
@ -664,266 +471,4 @@ bool oneapi_enqueue_kernel(KernelContext *kernel_context,
|
|||
# endif
|
||||
return success;
|
||||
}
|
||||
|
||||
/* Compute-runtime (ie. NEO) version is what gets returned by sycl/L0 on Windows
|
||||
* since Windows driver 101.3268. */
|
||||
/* The same min compute-runtime version is currently required across Windows and Linux.
|
||||
* For Windows driver 101.3430, compute-runtime version is 23904. */
|
||||
static const int lowest_supported_driver_version_win = 1013430;
|
||||
static const int lowest_supported_driver_version_neo = 23904;
|
||||
|
||||
static int parse_driver_build_version(const sycl::device &device)
|
||||
{
|
||||
const std::string &driver_version = device.get_info<sycl::info::device::driver_version>();
|
||||
int driver_build_version = 0;
|
||||
|
||||
size_t second_dot_position = driver_version.find('.', driver_version.find('.') + 1);
|
||||
if (second_dot_position == std::string::npos) {
|
||||
std::cerr << "Unable to parse unknown Intel GPU driver version \"" << driver_version
|
||||
<< "\" does not match xx.xx.xxxxx (Linux), x.x.xxxx (L0),"
|
||||
<< " xx.xx.xxx.xxxx (Windows) for device \""
|
||||
<< device.get_info<sycl::info::device::name>() << "\"." << std::endl;
|
||||
}
|
||||
else {
|
||||
try {
|
||||
size_t third_dot_position = driver_version.find('.', second_dot_position + 1);
|
||||
if (third_dot_position != std::string::npos) {
|
||||
const std::string &third_number_substr = driver_version.substr(
|
||||
second_dot_position + 1, third_dot_position - second_dot_position - 1);
|
||||
const std::string &forth_number_substr = driver_version.substr(third_dot_position + 1);
|
||||
if (third_number_substr.length() == 3 && forth_number_substr.length() == 4)
|
||||
driver_build_version = std::stoi(third_number_substr) * 10000 +
|
||||
std::stoi(forth_number_substr);
|
||||
}
|
||||
else {
|
||||
const std::string &third_number_substr = driver_version.substr(second_dot_position + 1);
|
||||
driver_build_version = std::stoi(third_number_substr);
|
||||
}
|
||||
}
|
||||
catch (std::invalid_argument &e) {
|
||||
std::cerr << "Unable to parse unknown Intel GPU driver version \"" << driver_version
|
||||
<< "\" does not match xx.xx.xxxxx (Linux), x.x.xxxx (L0),"
|
||||
<< " xx.xx.xxx.xxxx (Windows) for device \""
|
||||
<< device.get_info<sycl::info::device::name>() << "\"." << std::endl;
|
||||
}
|
||||
}
|
||||
|
||||
return driver_build_version;
|
||||
}
|
||||
|
||||
static std::vector<sycl::device> oneapi_available_devices()
|
||||
{
|
||||
bool allow_all_devices = false;
|
||||
if (getenv("CYCLES_ONEAPI_ALL_DEVICES") != nullptr)
|
||||
allow_all_devices = true;
|
||||
|
||||
/* Host device is useful only for debugging at the moment
|
||||
* so we hide this device with default build settings. */
|
||||
# ifdef WITH_ONEAPI_SYCL_HOST_ENABLED
|
||||
bool allow_host = true;
|
||||
# else
|
||||
bool allow_host = false;
|
||||
# endif
|
||||
|
||||
const std::vector<sycl::platform> &oneapi_platforms = sycl::platform::get_platforms();
|
||||
|
||||
std::vector<sycl::device> available_devices;
|
||||
for (const sycl::platform &platform : oneapi_platforms) {
|
||||
/* ignore OpenCL platforms to avoid using the same devices through both Level-Zero and OpenCL.
|
||||
*/
|
||||
if (platform.get_backend() == sycl::backend::opencl) {
|
||||
continue;
|
||||
}
|
||||
|
||||
const std::vector<sycl::device> &oneapi_devices =
|
||||
(allow_all_devices || allow_host) ? platform.get_devices(sycl::info::device_type::all) :
|
||||
platform.get_devices(sycl::info::device_type::gpu);
|
||||
|
||||
for (const sycl::device &device : oneapi_devices) {
|
||||
if (allow_all_devices) {
|
||||
/* still filter out host device if build doesn't support it. */
|
||||
if (allow_host || !device.is_host()) {
|
||||
available_devices.push_back(device);
|
||||
}
|
||||
}
|
||||
else {
|
||||
bool filter_out = false;
|
||||
|
||||
/* For now we support all Intel(R) Arc(TM) devices and likely any future GPU,
|
||||
* assuming they have either more than 96 Execution Units or not 7 threads per EU.
|
||||
* Official support can be broaden to older and smaller GPUs once ready. */
|
||||
if (device.is_gpu() && platform.get_backend() == sycl::backend::ext_oneapi_level_zero) {
|
||||
/* Filtered-out defaults in-case these values aren't available through too old L0
|
||||
* runtime. */
|
||||
int number_of_eus = 96;
|
||||
int threads_per_eu = 7;
|
||||
if (device.has(sycl::aspect::ext_intel_gpu_eu_count)) {
|
||||
number_of_eus = device.get_info<sycl::info::device::ext_intel_gpu_eu_count>();
|
||||
}
|
||||
if (device.has(sycl::aspect::ext_intel_gpu_hw_threads_per_eu)) {
|
||||
threads_per_eu =
|
||||
device.get_info<sycl::info::device::ext_intel_gpu_hw_threads_per_eu>();
|
||||
}
|
||||
/* This filters out all Level-Zero supported GPUs from older generation than Arc. */
|
||||
if (number_of_eus <= 96 && threads_per_eu == 7) {
|
||||
filter_out = true;
|
||||
}
|
||||
/* if not already filtered out, check driver version. */
|
||||
if (!filter_out) {
|
||||
int driver_build_version = parse_driver_build_version(device);
|
||||
if ((driver_build_version > 100000 &&
|
||||
driver_build_version < lowest_supported_driver_version_win) ||
|
||||
driver_build_version < lowest_supported_driver_version_neo) {
|
||||
filter_out = true;
|
||||
}
|
||||
}
|
||||
}
|
||||
else if (!allow_host && device.is_host()) {
|
||||
filter_out = true;
|
||||
}
|
||||
else if (!allow_all_devices) {
|
||||
filter_out = true;
|
||||
}
|
||||
|
||||
if (!filter_out) {
|
||||
available_devices.push_back(device);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
return available_devices;
|
||||
}
|
||||
|
||||
char *oneapi_device_capabilities()
|
||||
{
|
||||
std::stringstream capabilities;
|
||||
|
||||
const std::vector<sycl::device> &oneapi_devices = oneapi_available_devices();
|
||||
for (const sycl::device &device : oneapi_devices) {
|
||||
const std::string &name = device.get_info<sycl::info::device::name>();
|
||||
|
||||
capabilities << std::string("\t") << name << "\n";
|
||||
# define WRITE_ATTR(attribute_name, attribute_variable) \
|
||||
capabilities << "\t\tsycl::info::device::" #attribute_name "\t\t\t" << attribute_variable \
|
||||
<< "\n";
|
||||
# define GET_NUM_ATTR(attribute) \
|
||||
{ \
|
||||
size_t attribute = (size_t)device.get_info<sycl::info::device ::attribute>(); \
|
||||
capabilities << "\t\tsycl::info::device::" #attribute "\t\t\t" << attribute << "\n"; \
|
||||
}
|
||||
|
||||
GET_NUM_ATTR(vendor_id)
|
||||
GET_NUM_ATTR(max_compute_units)
|
||||
GET_NUM_ATTR(max_work_item_dimensions)
|
||||
|
||||
sycl::id<3> max_work_item_sizes =
|
||||
device.get_info<sycl::info::device::max_work_item_sizes<3>>();
|
||||
WRITE_ATTR("max_work_item_sizes_dim0", ((size_t)max_work_item_sizes.get(0)))
|
||||
WRITE_ATTR("max_work_item_sizes_dim1", ((size_t)max_work_item_sizes.get(1)))
|
||||
WRITE_ATTR("max_work_item_sizes_dim2", ((size_t)max_work_item_sizes.get(2)))
|
||||
|
||||
GET_NUM_ATTR(max_work_group_size)
|
||||
GET_NUM_ATTR(max_num_sub_groups)
|
||||
GET_NUM_ATTR(sub_group_independent_forward_progress)
|
||||
|
||||
GET_NUM_ATTR(preferred_vector_width_char)
|
||||
GET_NUM_ATTR(preferred_vector_width_short)
|
||||
GET_NUM_ATTR(preferred_vector_width_int)
|
||||
GET_NUM_ATTR(preferred_vector_width_long)
|
||||
GET_NUM_ATTR(preferred_vector_width_float)
|
||||
GET_NUM_ATTR(preferred_vector_width_double)
|
||||
GET_NUM_ATTR(preferred_vector_width_half)
|
||||
|
||||
GET_NUM_ATTR(native_vector_width_char)
|
||||
GET_NUM_ATTR(native_vector_width_short)
|
||||
GET_NUM_ATTR(native_vector_width_int)
|
||||
GET_NUM_ATTR(native_vector_width_long)
|
||||
GET_NUM_ATTR(native_vector_width_float)
|
||||
GET_NUM_ATTR(native_vector_width_double)
|
||||
GET_NUM_ATTR(native_vector_width_half)
|
||||
|
||||
size_t max_clock_frequency =
|
||||
(size_t)(device.is_host() ? (size_t)0 :
|
||||
device.get_info<sycl::info::device::max_clock_frequency>());
|
||||
WRITE_ATTR("max_clock_frequency", max_clock_frequency)
|
||||
|
||||
GET_NUM_ATTR(address_bits)
|
||||
GET_NUM_ATTR(max_mem_alloc_size)
|
||||
|
||||
/* NOTE(@nsirgien): Implementation doesn't use image support as bindless images aren't
|
||||
* supported so we always return false, even if device supports HW texture usage acceleration.
|
||||
*/
|
||||
bool image_support = false;
|
||||
WRITE_ATTR("image_support", (size_t)image_support)
|
||||
|
||||
GET_NUM_ATTR(max_parameter_size)
|
||||
GET_NUM_ATTR(mem_base_addr_align)
|
||||
GET_NUM_ATTR(global_mem_size)
|
||||
GET_NUM_ATTR(local_mem_size)
|
||||
GET_NUM_ATTR(error_correction_support)
|
||||
GET_NUM_ATTR(profiling_timer_resolution)
|
||||
GET_NUM_ATTR(is_available)
|
||||
|
||||
# undef GET_NUM_ATTR
|
||||
# undef WRITE_ATTR
|
||||
capabilities << "\n";
|
||||
}
|
||||
|
||||
return ::strdup(capabilities.str().c_str());
|
||||
}
|
||||
|
||||
void oneapi_free(void *p)
|
||||
{
|
||||
if (p) {
|
||||
::free(p);
|
||||
}
|
||||
}
|
||||
|
||||
void oneapi_iterate_devices(OneAPIDeviceIteratorCallback cb, void *user_ptr)
|
||||
{
|
||||
int num = 0;
|
||||
std::vector<sycl::device> devices = oneapi_available_devices();
|
||||
for (sycl::device &device : devices) {
|
||||
const std::string &platform_name =
|
||||
device.get_platform().get_info<sycl::info::platform::name>();
|
||||
std::string name = device.get_info<sycl::info::device::name>();
|
||||
std::string id = "ONEAPI_" + platform_name + "_" + name;
|
||||
if (device.has(sycl::aspect::ext_intel_pci_address)) {
|
||||
id.append("_" + device.get_info<sycl::info::device::ext_intel_pci_address>());
|
||||
}
|
||||
(cb)(id.c_str(), name.c_str(), num, user_ptr);
|
||||
num++;
|
||||
}
|
||||
}
|
||||
|
||||
size_t oneapi_get_memcapacity(SyclQueue *queue)
|
||||
{
|
||||
return reinterpret_cast<sycl::queue *>(queue)
|
||||
->get_device()
|
||||
.get_info<sycl::info::device::global_mem_size>();
|
||||
}
|
||||
|
||||
int oneapi_get_num_multiprocessors(SyclQueue *queue)
|
||||
{
|
||||
const sycl::device &device = reinterpret_cast<sycl::queue *>(queue)->get_device();
|
||||
if (device.has(sycl::aspect::ext_intel_gpu_eu_count)) {
|
||||
return device.get_info<sycl::info::device::ext_intel_gpu_eu_count>();
|
||||
}
|
||||
else
|
||||
return 0;
|
||||
}
|
||||
|
||||
int oneapi_get_max_num_threads_per_multiprocessor(SyclQueue *queue)
|
||||
{
|
||||
const sycl::device &device = reinterpret_cast<sycl::queue *>(queue)->get_device();
|
||||
if (device.has(sycl::aspect::ext_intel_gpu_eu_simd_width) &&
|
||||
device.has(sycl::aspect::ext_intel_gpu_hw_threads_per_eu)) {
|
||||
return device.get_info<sycl::info::device::ext_intel_gpu_eu_simd_width>() *
|
||||
device.get_info<sycl::info::device::ext_intel_gpu_hw_threads_per_eu>();
|
||||
}
|
||||
else
|
||||
return 0;
|
||||
}
|
||||
|
||||
#endif /* WITH_ONEAPI */
|
||||
|
|
|
@ -25,11 +25,6 @@ enum DeviceKernel : int;
|
|||
|
||||
class SyclQueue;
|
||||
|
||||
typedef void (*OneAPIDeviceIteratorCallback)(const char *id,
|
||||
const char *name,
|
||||
int num,
|
||||
void *user_ptr);
|
||||
|
||||
typedef void (*OneAPIErrorCallback)(const char *error, void *user_ptr);
|
||||
|
||||
struct KernelContext {
|
||||
|
@ -45,13 +40,15 @@ struct KernelContext {
|
|||
extern "C" {
|
||||
# endif
|
||||
|
||||
# define DLL_INTERFACE_CALL(function, return_type, ...) \
|
||||
CYCLES_KERNEL_ONEAPI_EXPORT return_type function(__VA_ARGS__);
|
||||
# include "kernel/device/oneapi/dll_interface_template.h"
|
||||
# undef DLL_INTERFACE_CALL
|
||||
|
||||
CYCLES_KERNEL_ONEAPI_EXPORT bool oneapi_run_test_kernel(SyclQueue *queue_);
|
||||
CYCLES_KERNEL_ONEAPI_EXPORT void oneapi_set_error_cb(OneAPIErrorCallback cb, void *user_ptr);
|
||||
CYCLES_KERNEL_ONEAPI_EXPORT size_t oneapi_kernel_preferred_local_size(
|
||||
SyclQueue *queue, const DeviceKernel kernel, const size_t kernel_global_size);
|
||||
CYCLES_KERNEL_ONEAPI_EXPORT bool oneapi_enqueue_kernel(KernelContext *context,
|
||||
int kernel,
|
||||
size_t global_size,
|
||||
void **args);
|
||||
# ifdef __cplusplus
|
||||
}
|
||||
# endif
|
||||
|
||||
#endif /* WITH_ONEAPI */
|
||||
|
|
|
@ -68,6 +68,12 @@ static void keyboard_handle_key_repeat_cancel(struct GWL_Seat *seat);
|
|||
|
||||
static void output_handle_done(void *data, struct wl_output *wl_output);
|
||||
|
||||
/* -------------------------------------------------------------------- */
|
||||
/** \name Local Defines
|
||||
*
|
||||
* Control local functionality, compositors specific workarounds.
|
||||
* \{ */
|
||||
|
||||
/**
|
||||
* GNOME (mutter 42.2 had a bug with confine not respecting scale - Hi-DPI), See: T98793.
|
||||
* Even though this has been fixed, at time of writing it's not yet in a release.
|
||||
|
@ -86,6 +92,22 @@ static void output_handle_done(void *data, struct wl_output *wl_output);
|
|||
static bool use_gnome_confine_hack = false;
|
||||
#endif
|
||||
|
||||
/**
|
||||
* GNOME (mutter 42.5) doesn't follow the WAYLAND spec regarding keyboard handling,
|
||||
* unlike (other compositors: KDE-plasma, River & Sway which work without problems).
|
||||
*
|
||||
* This means GNOME can't know which modifiers are held when activating windows,
|
||||
* so we guess the left modifiers are held.
|
||||
*
|
||||
* This define could be removed without changing any functionality,
|
||||
* it just means GNOME users will see verbose warning messages that alert them about
|
||||
* a known problem that needs to be fixed up-stream.
|
||||
* See: https://gitlab.gnome.org/GNOME/mutter/-/issues/2457
|
||||
*/
|
||||
#define USE_GNOME_KEYBOARD_SUPPRESS_WARNING
|
||||
|
||||
/** \} */
|
||||
|
||||
/* -------------------------------------------------------------------- */
|
||||
/** \name Inline Event Codes
|
||||
*
|
||||
|
@ -385,6 +407,13 @@ struct GWL_Seat {
|
|||
/** Keys held matching `xkb_state`. */
|
||||
struct WGL_KeyboardDepressedState key_depressed;
|
||||
|
||||
#ifdef USE_GNOME_KEYBOARD_SUPPRESS_WARNING
|
||||
struct {
|
||||
bool any_mod_held = false;
|
||||
bool any_keys_held_on_enter = false;
|
||||
} key_depressed_suppress_warning;
|
||||
#endif
|
||||
|
||||
/**
|
||||
* Cache result of `xkb_keymap_mod_get_index`
|
||||
* so every time a modifier is accessed a string lookup isn't required.
|
||||
|
@ -2303,6 +2332,7 @@ static void keyboard_handle_enter(void *data,
|
|||
uint32_t *key;
|
||||
WL_ARRAY_FOR_EACH (key, keys) {
|
||||
const xkb_keycode_t key_code = *key + EVDEV_OFFSET;
|
||||
CLOG_INFO(LOG, 2, "enter (key_held=%d)", int(key_code));
|
||||
const xkb_keysym_t sym = xkb_state_key_get_one_sym(seat->xkb_state, key_code);
|
||||
const GHOST_TKey gkey = xkb_map_gkey_or_scan_code(sym, *key);
|
||||
if (gkey != GHOST_kKeyUnknown) {
|
||||
|
@ -2311,6 +2341,10 @@ static void keyboard_handle_enter(void *data,
|
|||
}
|
||||
|
||||
keyboard_depressed_state_push_events_from_change(seat, key_depressed_prev);
|
||||
|
||||
#ifdef USE_GNOME_KEYBOARD_SUPPRESS_WARNING
|
||||
seat->key_depressed_suppress_warning.any_keys_held_on_enter = keys->size != 0;
|
||||
#endif
|
||||
}
|
||||
|
||||
/**
|
||||
|
@ -2336,6 +2370,11 @@ static void keyboard_handle_leave(void *data,
|
|||
if (seat->key_repeat.timer) {
|
||||
keyboard_handle_key_repeat_cancel(seat);
|
||||
}
|
||||
|
||||
#ifdef USE_GNOME_KEYBOARD_SUPPRESS_WARNING
|
||||
seat->key_depressed_suppress_warning.any_mod_held = false;
|
||||
seat->key_depressed_suppress_warning.any_keys_held_on_enter = false;
|
||||
#endif
|
||||
}
|
||||
|
||||
/**
|
||||
|
@ -2407,10 +2446,10 @@ static void keyboard_handle_key(void *data,
|
|||
const xkb_keysym_t sym = xkb_state_key_get_one_sym_without_modifiers(
|
||||
seat->xkb_state_empty, seat->xkb_state_empty_with_numlock, key_code);
|
||||
if (sym == XKB_KEY_NoSymbol) {
|
||||
CLOG_INFO(LOG, 2, "key (no symbol, skipped)");
|
||||
CLOG_INFO(LOG, 2, "key (code=%d, state=%u, no symbol, skipped)", int(key_code), state);
|
||||
return;
|
||||
}
|
||||
CLOG_INFO(LOG, 2, "key");
|
||||
CLOG_INFO(LOG, 2, "key (code=%d, state=%u)", int(key_code), state);
|
||||
|
||||
GHOST_TEventType etype = GHOST_kEventUnknown;
|
||||
switch (state) {
|
||||
|
@ -2554,6 +2593,10 @@ static void keyboard_handle_modifiers(void *data,
|
|||
if (seat->key_repeat.timer) {
|
||||
keyboard_handle_key_repeat_reset(seat, true);
|
||||
}
|
||||
|
||||
#ifdef USE_GNOME_KEYBOARD_SUPPRESS_WARNING
|
||||
seat->key_depressed_suppress_warning.any_mod_held = mods_depressed != 0;
|
||||
#endif
|
||||
}
|
||||
|
||||
static void keyboard_repeat_handle_info(void *data,
|
||||
|
@ -3163,6 +3206,15 @@ GHOST_TSuccess GHOST_SystemWayland::getModifierKeys(GHOST_ModifierKeys &keys) co
|
|||
|
||||
const xkb_mod_mask_t state = xkb_state_serialize_mods(seat->xkb_state, XKB_STATE_MODS_DEPRESSED);
|
||||
|
||||
bool show_warning = true;
|
||||
#ifdef USE_GNOME_KEYBOARD_SUPPRESS_WARNING
|
||||
if ((seat->key_depressed_suppress_warning.any_mod_held == true) &&
|
||||
(seat->key_depressed_suppress_warning.any_keys_held_on_enter == false)) {
|
||||
/* The compositor gave us invalid information, don't show a warning. */
|
||||
show_warning = false;
|
||||
}
|
||||
#endif
|
||||
|
||||
/* Use local #WGL_KeyboardDepressedState to check which key is pressed.
|
||||
* Use XKB as the source of truth, if there is any discrepancy. */
|
||||
for (int i = 0; i < MOD_INDEX_NUM; i++) {
|
||||
|
@ -3178,18 +3230,22 @@ GHOST_TSuccess GHOST_SystemWayland::getModifierKeys(GHOST_ModifierKeys &keys) co
|
|||
* Warn so if this happens it can be investigated. */
|
||||
if (val) {
|
||||
if (UNLIKELY(!(val_l || val_r))) {
|
||||
CLOG_WARN(&LOG_WL_KEYBOARD_DEPRESSED_STATE,
|
||||
"modifier (%s) state is inconsistent (held keys do not match XKB)",
|
||||
mod_info.display_name);
|
||||
if (show_warning) {
|
||||
CLOG_WARN(&LOG_WL_KEYBOARD_DEPRESSED_STATE,
|
||||
"modifier (%s) state is inconsistent (GHOST held keys do not match XKB)",
|
||||
mod_info.display_name);
|
||||
}
|
||||
/* Picking the left is arbitrary. */
|
||||
val_l = true;
|
||||
}
|
||||
}
|
||||
else {
|
||||
if (UNLIKELY(val_l || val_r)) {
|
||||
CLOG_WARN(&LOG_WL_KEYBOARD_DEPRESSED_STATE,
|
||||
"modifier (%s) state is inconsistent (released keys do not match XKB)",
|
||||
mod_info.display_name);
|
||||
if (show_warning) {
|
||||
CLOG_WARN(&LOG_WL_KEYBOARD_DEPRESSED_STATE,
|
||||
"modifier (%s) state is inconsistent (GHOST released keys do not match XKB)",
|
||||
mod_info.display_name);
|
||||
}
|
||||
val_l = false;
|
||||
val_r = false;
|
||||
}
|
||||
|
|
|
@ -98,7 +98,7 @@ static int blf_search(const char *name)
|
|||
{
|
||||
for (int i = 0; i < BLF_MAX_FONT; i++) {
|
||||
const FontBLF *font = global_font[i];
|
||||
if (font && (STREQ(font->name, name))) {
|
||||
if (font && STREQ(font->name, name)) {
|
||||
return i;
|
||||
}
|
||||
}
|
||||
|
@ -226,7 +226,7 @@ void BLF_unload(const char *name)
|
|||
for (int i = 0; i < BLF_MAX_FONT; i++) {
|
||||
FontBLF *font = global_font[i];
|
||||
|
||||
if (font && (STREQ(font->name, name))) {
|
||||
if (font && STREQ(font->name, name)) {
|
||||
BLI_assert(font->reference_count > 0);
|
||||
font->reference_count--;
|
||||
|
||||
|
|
|
@ -54,7 +54,7 @@
|
|||
*/
|
||||
static FT_Fixed to_16dot16(double val)
|
||||
{
|
||||
return (FT_Fixed)(lround(val * 65536.0));
|
||||
return (FT_Fixed)lround(val * 65536.0);
|
||||
}
|
||||
|
||||
/** \} */
|
||||
|
|
|
@ -3198,7 +3198,7 @@ static void animsys_create_action_track_strip(const AnimData *adt,
|
|||
* (which making new strips doesn't do due to the troublesome nature of that). */
|
||||
calc_action_range(r_action_strip->act, &r_action_strip->actstart, &r_action_strip->actend, 1);
|
||||
r_action_strip->start = r_action_strip->actstart;
|
||||
r_action_strip->end = (IS_EQF(r_action_strip->actstart, r_action_strip->actend)) ?
|
||||
r_action_strip->end = IS_EQF(r_action_strip->actstart, r_action_strip->actend) ?
|
||||
(r_action_strip->actstart + 1.0f) :
|
||||
(r_action_strip->actend);
|
||||
|
||||
|
|
|
@ -301,12 +301,22 @@ static void armature_vert_task_with_dvert(const ArmatureUserdata *data,
|
|||
}
|
||||
|
||||
/* check if there's any point in calculating for this vert */
|
||||
if (armature_weight == 0.0f) {
|
||||
return;
|
||||
}
|
||||
if (vert_coords_prev) {
|
||||
if (prevco_weight == 1.0f) {
|
||||
return;
|
||||
}
|
||||
|
||||
/* get the coord we work on */
|
||||
co = vert_coords_prev ? vert_coords_prev[i] : vert_coords[i];
|
||||
/* get the coord we work on */
|
||||
co = vert_coords_prev[i];
|
||||
}
|
||||
else {
|
||||
if (armature_weight == 0.0f) {
|
||||
return;
|
||||
}
|
||||
|
||||
/* get the coord we work on */
|
||||
co = vert_coords[i];
|
||||
}
|
||||
|
||||
/* Apply the object's matrix */
|
||||
mul_m4_v3(data->premat, co);
|
||||
|
|
|
@ -25,7 +25,7 @@
|
|||
bUserMenu *BKE_blender_user_menu_find(ListBase *lb, char space_type, const char *context)
|
||||
{
|
||||
LISTBASE_FOREACH (bUserMenu *, um, lb) {
|
||||
if ((space_type == um->space_type) && (STREQ(context, um->context))) {
|
||||
if ((space_type == um->space_type) && STREQ(context, um->context)) {
|
||||
return um;
|
||||
}
|
||||
}
|
||||
|
|
|
@ -713,10 +713,8 @@ static bool camera_frame_fit_calc_from_data(CameraParams *params,
|
|||
plane_from_point_normal_v3(plane_tx[i], co, data->plane_tx[i]);
|
||||
}
|
||||
|
||||
if ((!isect_plane_plane_v3(
|
||||
plane_tx[Y_MIN], plane_tx[Y_MAX], plane_isect_1, plane_isect_1_no)) ||
|
||||
(!isect_plane_plane_v3(
|
||||
plane_tx[Z_MIN], plane_tx[Z_MAX], plane_isect_2, plane_isect_2_no))) {
|
||||
if (!isect_plane_plane_v3(plane_tx[Y_MIN], plane_tx[Y_MAX], plane_isect_1, plane_isect_1_no) ||
|
||||
!isect_plane_plane_v3(plane_tx[Z_MIN], plane_tx[Z_MAX], plane_isect_2, plane_isect_2_no)) {
|
||||
return false;
|
||||
}
|
||||
|
||||
|
|
|
@ -1129,7 +1129,7 @@ static void cloth_update_springs(ClothModifierData *clmd)
|
|||
* because implicit solver would need reset! */
|
||||
|
||||
/* Activate / Deactivate existing springs */
|
||||
if ((!(cloth->verts[spring->ij].flags & CLOTH_VERT_FLAG_PINNED)) &&
|
||||
if (!(cloth->verts[spring->ij].flags & CLOTH_VERT_FLAG_PINNED) &&
|
||||
(cloth->verts[spring->ij].goal > ALMOST_ZERO)) {
|
||||
spring->flags &= ~CLOTH_SPRING_FLAG_DEACTIVATE;
|
||||
}
|
||||
|
|
|
@ -946,7 +946,7 @@ bool BKE_collection_has_object(Collection *collection, const Object *ob)
|
|||
return false;
|
||||
}
|
||||
|
||||
return (BLI_findptr(&collection->gobject, ob, offsetof(CollectionObject, ob)));
|
||||
return BLI_findptr(&collection->gobject, ob, offsetof(CollectionObject, ob));
|
||||
}
|
||||
|
||||
bool BKE_collection_has_object_recursive(Collection *collection, Object *ob)
|
||||
|
@ -956,7 +956,7 @@ bool BKE_collection_has_object_recursive(Collection *collection, Object *ob)
|
|||
}
|
||||
|
||||
const ListBase objects = BKE_collection_object_cache_get(collection);
|
||||
return (BLI_findptr(&objects, ob, offsetof(Base, object)));
|
||||
return BLI_findptr(&objects, ob, offsetof(Base, object));
|
||||
}
|
||||
|
||||
bool BKE_collection_has_object_recursive_instanced(Collection *collection, Object *ob)
|
||||
|
@ -966,7 +966,7 @@ bool BKE_collection_has_object_recursive_instanced(Collection *collection, Objec
|
|||
}
|
||||
|
||||
const ListBase objects = BKE_collection_object_cache_instanced_get(collection);
|
||||
return (BLI_findptr(&objects, ob, offsetof(Base, object)));
|
||||
return BLI_findptr(&objects, ob, offsetof(Base, object));
|
||||
}
|
||||
|
||||
static Collection *collection_next_find(Main *bmain, Scene *scene, Collection *collection)
|
||||
|
|
|
@ -1629,7 +1629,7 @@ void BKE_scopes_update(Scopes *scopes,
|
|||
}
|
||||
|
||||
/* hmmmm */
|
||||
if (!(ELEM(ibuf->channels, 3, 4))) {
|
||||
if (!ELEM(ibuf->channels, 3, 4)) {
|
||||
return;
|
||||
}
|
||||
|
||||
|
|
|
@ -634,7 +634,7 @@ ListBase CTX_data_dir_get(const bContext *C)
|
|||
|
||||
bool CTX_data_equals(const char *member, const char *str)
|
||||
{
|
||||
return (STREQ(member, str));
|
||||
return STREQ(member, str);
|
||||
}
|
||||
|
||||
bool CTX_data_dir(const char *member)
|
||||
|
|
|
@ -464,7 +464,7 @@ static std::string to_manifest(const CryptomatteLayer *layer)
|
|||
else {
|
||||
manifest << ",";
|
||||
}
|
||||
manifest << quoted(item.key) << ":\"" << (item.value.hex_encoded()) << "\"";
|
||||
manifest << quoted(item.key) << ":\"" << item.value.hex_encoded() << "\"";
|
||||
}
|
||||
manifest << "}";
|
||||
return manifest.str();
|
||||
|
|
|
@ -1935,7 +1935,7 @@ static void calc_bevel_sin_cos(
|
|||
t02 = M_PI_2;
|
||||
}
|
||||
else {
|
||||
t02 = (saacos(t02)) / 2.0f;
|
||||
t02 = saacos(t02) / 2.0f;
|
||||
}
|
||||
|
||||
t02 = sinf(t02);
|
||||
|
@ -4085,12 +4085,12 @@ void BKE_nurb_bezt_handle_test(BezTriple *bezt,
|
|||
}
|
||||
|
||||
if (bezt->h1 == HD_VECT) {
|
||||
if ((!(flag & SEL_F1)) != (!(flag & SEL_F2))) {
|
||||
if (!(flag & SEL_F1) != !(flag & SEL_F2)) {
|
||||
bezt->h1 = HD_FREE;
|
||||
}
|
||||
}
|
||||
if (bezt->h2 == HD_VECT) {
|
||||
if ((!(flag & SEL_F3)) != (!(flag & SEL_F2))) {
|
||||
if (!(flag & SEL_F3) != !(flag & SEL_F2)) {
|
||||
bezt->h2 = HD_FREE;
|
||||
}
|
||||
}
|
||||
|
@ -5101,7 +5101,7 @@ bool BKE_curve_minmax(Curve *cu, bool use_radius, float min[3], float max[3])
|
|||
{
|
||||
ListBase *nurb_lb = BKE_curve_nurbs_get(cu);
|
||||
ListBase temp_nurb_lb = {nullptr, nullptr};
|
||||
const bool is_font = (BLI_listbase_is_empty(nurb_lb)) && (cu->len != 0);
|
||||
const bool is_font = BLI_listbase_is_empty(nurb_lb) && (cu->len != 0);
|
||||
/* For font curves we generate temp list of splines.
|
||||
*
|
||||
* This is likely to be fine, this function is not supposed to be called
|
||||
|
|
|
@ -160,7 +160,7 @@ static bool calc_curve_deform(
|
|||
/* Zero the axis which is not used,
|
||||
* the big block of text above now applies to these 3 lines.
|
||||
* The `upflag` argument may be a dummy, set so no rotation is done. */
|
||||
quat_apply_track(quat, axis, (ELEM(axis, 0, 2)) ? 1 : 0);
|
||||
quat_apply_track(quat, axis, ELEM(axis, 0, 2) ? 1 : 0);
|
||||
vec_apply_track(cent, axis);
|
||||
cent[index] = 0.0f;
|
||||
|
||||
|
|
|
@ -121,7 +121,7 @@ BLI_INLINE void value_dissolve(float *r_value,
|
|||
const float scale,
|
||||
const bool is_log)
|
||||
{
|
||||
*r_value = (is_log) ? (*r_value) * (powf(MIN_WETNESS, 1.0f / (1.2f * time / scale))) :
|
||||
*r_value = (is_log) ? (*r_value) * powf(MIN_WETNESS, 1.0f / (1.2f * time / scale)) :
|
||||
(*r_value) - 1.0f / time * scale;
|
||||
}
|
||||
|
||||
|
|
|
@ -193,7 +193,7 @@ static void cage_mapped_verts_callback(void *userData,
|
|||
{
|
||||
CageUserData *data = static_cast<CageUserData *>(userData);
|
||||
|
||||
if ((index >= 0 && index < data->totvert) && (!BLI_BITMAP_TEST(data->visit_bitmap, index))) {
|
||||
if ((index >= 0 && index < data->totvert) && !BLI_BITMAP_TEST(data->visit_bitmap, index)) {
|
||||
BLI_BITMAP_ENABLE(data->visit_bitmap, index);
|
||||
copy_v3_v3(data->cos_cage[index], co);
|
||||
}
|
||||
|
|
|
@ -79,7 +79,7 @@ PartDeflect *BKE_partdeflect_new(int type)
|
|||
pd->pdef_sbift = 0.2f;
|
||||
pd->pdef_sboft = 0.02f;
|
||||
pd->pdef_cfrict = 5.0f;
|
||||
pd->seed = ((uint)(ceil(PIL_check_seconds_timer())) + 1) % 128;
|
||||
pd->seed = ((uint)ceil(PIL_check_seconds_timer()) + 1) % 128;
|
||||
pd->f_strength = 1.0f;
|
||||
pd->f_damp = 1.0f;
|
||||
|
||||
|
|
|
@ -333,7 +333,7 @@ static float dvar_eval_rotDiff(ChannelDriver *driver, DriverVar *dvar)
|
|||
|
||||
invert_qt_normalized(q1);
|
||||
mul_qt_qtqt(quat, q1, q2);
|
||||
angle = 2.0f * (saacos(quat[0]));
|
||||
angle = 2.0f * saacos(quat[0]);
|
||||
angle = fabsf(angle);
|
||||
|
||||
return (angle > (float)M_PI) ? (float)((2.0f * (float)M_PI) - angle) : (float)(angle);
|
||||
|
|
|
@ -2204,9 +2204,9 @@ static void adaptive_domain_adjust(
|
|||
/* add to total shift */
|
||||
add_v3_v3(fds->shift_f, frame_shift_f);
|
||||
/* convert to integer */
|
||||
total_shift[0] = (int)(floorf(fds->shift_f[0]));
|
||||
total_shift[1] = (int)(floorf(fds->shift_f[1]));
|
||||
total_shift[2] = (int)(floorf(fds->shift_f[2]));
|
||||
total_shift[0] = (int)floorf(fds->shift_f[0]);
|
||||
total_shift[1] = (int)floorf(fds->shift_f[1]);
|
||||
total_shift[2] = (int)floorf(fds->shift_f[2]);
|
||||
int temp_shift[3];
|
||||
copy_v3_v3_int(temp_shift, fds->shift);
|
||||
sub_v3_v3v3_int(new_shift, total_shift, fds->shift);
|
||||
|
|
|
@ -83,8 +83,8 @@ static void greasepencil_copy_data(Main *UNUSED(bmain),
|
|||
/* Apply local layer transform to all frames. Calc the active frame is not enough
|
||||
* because onion skin can use more frames. This is more slow but required here. */
|
||||
if (gpl_dst->actframe != NULL) {
|
||||
bool transformed = ((!is_zero_v3(gpl_dst->location)) || (!is_zero_v3(gpl_dst->rotation)) ||
|
||||
(!is_one_v3(gpl_dst->scale)));
|
||||
bool transformed = (!is_zero_v3(gpl_dst->location) || !is_zero_v3(gpl_dst->rotation) ||
|
||||
!is_one_v3(gpl_dst->scale));
|
||||
if (transformed) {
|
||||
loc_eul_size_to_mat4(
|
||||
gpl_dst->layer_mat, gpl_dst->location, gpl_dst->rotation, gpl_dst->scale);
|
||||
|
@ -2013,7 +2013,7 @@ bool BKE_gpencil_merge_materials_table_get(Object *ob,
|
|||
/* Read secondary material to compare with primary material. */
|
||||
ma_secondary = BKE_gpencil_material(ob, idx_secondary + 1);
|
||||
if ((ma_secondary == NULL) ||
|
||||
(BLI_ghash_haskey(r_mat_table, POINTER_FROM_INT(idx_secondary)))) {
|
||||
BLI_ghash_haskey(r_mat_table, POINTER_FROM_INT(idx_secondary))) {
|
||||
continue;
|
||||
}
|
||||
gp_style_primary = ma_primary->gp_style;
|
||||
|
@ -2063,17 +2063,17 @@ bool BKE_gpencil_merge_materials_table_get(Object *ob,
|
|||
rgb_to_hsv_compat_v(col, f_hsv_b);
|
||||
|
||||
/* Check stroke and fill color. */
|
||||
if ((!compare_ff(s_hsv_a[0], s_hsv_b[0], hue_threshold)) ||
|
||||
(!compare_ff(s_hsv_a[1], s_hsv_b[1], sat_threshold)) ||
|
||||
(!compare_ff(s_hsv_a[2], s_hsv_b[2], val_threshold)) ||
|
||||
(!compare_ff(f_hsv_a[0], f_hsv_b[0], hue_threshold)) ||
|
||||
(!compare_ff(f_hsv_a[1], f_hsv_b[1], sat_threshold)) ||
|
||||
(!compare_ff(f_hsv_a[2], f_hsv_b[2], val_threshold)) ||
|
||||
(!compare_ff(gp_style_primary->stroke_rgba[3],
|
||||
gp_style_secondary->stroke_rgba[3],
|
||||
val_threshold)) ||
|
||||
(!compare_ff(
|
||||
gp_style_primary->fill_rgba[3], gp_style_secondary->fill_rgba[3], val_threshold))) {
|
||||
if (!compare_ff(s_hsv_a[0], s_hsv_b[0], hue_threshold) ||
|
||||
!compare_ff(s_hsv_a[1], s_hsv_b[1], sat_threshold) ||
|
||||
!compare_ff(s_hsv_a[2], s_hsv_b[2], val_threshold) ||
|
||||
!compare_ff(f_hsv_a[0], f_hsv_b[0], hue_threshold) ||
|
||||
!compare_ff(f_hsv_a[1], f_hsv_b[1], sat_threshold) ||
|
||||
!compare_ff(f_hsv_a[2], f_hsv_b[2], val_threshold) ||
|
||||
!compare_ff(gp_style_primary->stroke_rgba[3],
|
||||
gp_style_secondary->stroke_rgba[3],
|
||||
val_threshold) ||
|
||||
!compare_ff(
|
||||
gp_style_primary->fill_rgba[3], gp_style_secondary->fill_rgba[3], val_threshold)) {
|
||||
continue;
|
||||
}
|
||||
|
||||
|
@ -2337,7 +2337,7 @@ bool BKE_gpencil_from_image(
|
|||
static bool gpencil_is_layer_mask(ViewLayer *view_layer, bGPdata *gpd, bGPDlayer *gpl_mask)
|
||||
{
|
||||
LISTBASE_FOREACH (bGPDlayer *, gpl, &gpd->layers) {
|
||||
if ((gpl->viewlayername[0] != '\0') && (!STREQ(view_layer->name, gpl->viewlayername))) {
|
||||
if ((gpl->viewlayername[0] != '\0') && !STREQ(view_layer->name, gpl->viewlayername)) {
|
||||
continue;
|
||||
}
|
||||
|
||||
|
@ -2412,7 +2412,7 @@ void BKE_gpencil_visible_stroke_advanced_iter(ViewLayer *view_layer,
|
|||
int cfra)
|
||||
{
|
||||
bGPdata *gpd = (bGPdata *)ob->data;
|
||||
const bool is_multiedit = (GPENCIL_MULTIEDIT_SESSIONS_ON(gpd) && (!GPENCIL_PLAY_ON(gpd)));
|
||||
const bool is_multiedit = (GPENCIL_MULTIEDIT_SESSIONS_ON(gpd) && !GPENCIL_PLAY_ON(gpd));
|
||||
const bool is_onion = do_onion && ((gpd->flag & GP_DATA_STROKE_WEIGHTMODE) == 0);
|
||||
const bool is_drawing = (gpd->runtime.sbuffer_used > 0);
|
||||
|
||||
|
@ -2444,7 +2444,7 @@ void BKE_gpencil_visible_stroke_advanced_iter(ViewLayer *view_layer,
|
|||
* generate renders, putting only selected GP layers for each View Layer.
|
||||
* This is used only in final render and never in Viewport. */
|
||||
if ((view_layer != NULL) && (gpl->viewlayername[0] != '\0') &&
|
||||
(!STREQ(view_layer->name, gpl->viewlayername))) {
|
||||
!STREQ(view_layer->name, gpl->viewlayername)) {
|
||||
/* Do not skip masks when rendering the view-layer so that it can still be used to clip
|
||||
* other layers. Instead set their opacity to zero. */
|
||||
if (gpencil_is_layer_mask(view_layer, gpd, gpl)) {
|
||||
|
@ -2786,8 +2786,8 @@ void BKE_gpencil_update_layer_transforms(const Depsgraph *depsgraph, Object *ob)
|
|||
}
|
||||
|
||||
/* Calc local layer transform. Early out if we have non-animated zero transforms. */
|
||||
bool transformed = ((!is_zero_v3(gpl->location)) || (!is_zero_v3(gpl->rotation)) ||
|
||||
(!is_one_v3(gpl->scale)));
|
||||
bool transformed = (!is_zero_v3(gpl->location) || !is_zero_v3(gpl->rotation) ||
|
||||
!is_one_v3(gpl->scale));
|
||||
float tmp_mat[4][4];
|
||||
loc_eul_size_to_mat4(tmp_mat, gpl->location, gpl->rotation, gpl->scale);
|
||||
transformed |= !equals_m4m4(gpl->layer_mat, tmp_mat);
|
||||
|
@ -2834,7 +2834,7 @@ int BKE_gpencil_material_find_index_by_name_prefix(Object *ob, const char *name_
|
|||
for (int i = 0; i < ob->totcol; i++) {
|
||||
Material *ma = BKE_object_material_get(ob, i + 1);
|
||||
if ((ma != NULL) && (ma->gp_style != NULL) &&
|
||||
(STREQLEN(ma->id.name + 2, name_prefix, name_prefix_len))) {
|
||||
STREQLEN(ma->id.name + 2, name_prefix, name_prefix_len)) {
|
||||
return i;
|
||||
}
|
||||
}
|
||||
|
|
|
@ -722,8 +722,8 @@ void BKE_gpencil_prepare_eval_data(Depsgraph *depsgraph, Scene *scene, Object *o
|
|||
}
|
||||
|
||||
/* Only do layer transformations for non-zero or animated transforms. */
|
||||
bool transformed = ((!is_zero_v3(gpl->location)) || (!is_zero_v3(gpl->rotation)) ||
|
||||
(!is_one_v3(gpl->scale)));
|
||||
bool transformed = (!is_zero_v3(gpl->location) || !is_zero_v3(gpl->rotation) ||
|
||||
!is_one_v3(gpl->scale));
|
||||
float tmp_mat[4][4];
|
||||
loc_eul_size_to_mat4(tmp_mat, gpl->location, gpl->rotation, gpl->scale);
|
||||
transformed |= !equals_m4m4(gpl->layer_mat, tmp_mat);
|
||||
|
@ -752,7 +752,7 @@ void BKE_gpencil_prepare_eval_data(Depsgraph *depsgraph, Scene *scene, Object *o
|
|||
const bool is_curve_edit = (bool)GPENCIL_CURVE_EDIT_SESSIONS_ON(gpd_orig);
|
||||
const bool do_modifiers = (bool)((!is_multiedit) && (!is_curve_edit) &&
|
||||
(ob_orig->greasepencil_modifiers.first != NULL) &&
|
||||
(!GPENCIL_SIMPLIFY_MODIF(scene)));
|
||||
!GPENCIL_SIMPLIFY_MODIF(scene));
|
||||
if ((!do_modifiers) && (!do_parent) && (!do_transform)) {
|
||||
BLI_assert(ob->data != NULL);
|
||||
return;
|
||||
|
@ -782,7 +782,7 @@ void BKE_gpencil_modifiers_calc(Depsgraph *depsgraph, Scene *scene, Object *ob)
|
|||
const bool is_multiedit = (bool)(GPENCIL_MULTIEDIT_SESSIONS_ON(gpd) && !is_render);
|
||||
const bool do_modifiers = (bool)((!is_multiedit) && (!is_curve_edit) &&
|
||||
(ob->greasepencil_modifiers.first != NULL) &&
|
||||
(!GPENCIL_SIMPLIFY_MODIF(scene)));
|
||||
!GPENCIL_SIMPLIFY_MODIF(scene));
|
||||
if (!do_modifiers) {
|
||||
return;
|
||||
}
|
||||
|
|
|
@ -520,7 +520,7 @@ static bool do_add_image_extension(char *string,
|
|||
}
|
||||
#endif
|
||||
else { // R_IMF_IMTYPE_AVIRAW, R_IMF_IMTYPE_AVIJPEG, R_IMF_IMTYPE_JPEG90 etc
|
||||
if (!(BLI_path_extension_check_n(string, extension_test = ".jpg", ".jpeg", nullptr))) {
|
||||
if (!BLI_path_extension_check_n(string, extension_test = ".jpg", ".jpeg", nullptr)) {
|
||||
extension = extension_test;
|
||||
}
|
||||
}
|
||||
|
|
|
@ -824,7 +824,7 @@ bool BKE_image_render_write_exr(ReportList *reports,
|
|||
|
||||
/* We only store RGBA passes as half float, for
|
||||
* others precision loss can be problematic. */
|
||||
const bool pass_RGBA = (STR_ELEM(rp->chan_id, "RGB", "RGBA", "R", "G", "B", "A"));
|
||||
const bool pass_RGBA = STR_ELEM(rp->chan_id, "RGB", "RGBA", "R", "G", "B", "A");
|
||||
const bool pass_half_float = half_float && pass_RGBA;
|
||||
|
||||
/* Color-space conversion only happens on RGBA passes. */
|
||||
|
|
|
@ -1542,7 +1542,7 @@ static void icu_to_fcurves(ID *id,
|
|||
}
|
||||
|
||||
/* correct values, by checking if the flag of interest is set */
|
||||
if (((int)(dst->vec[1][1])) & (abp->bit)) {
|
||||
if ((int)(dst->vec[1][1]) & (abp->bit)) {
|
||||
dst->vec[0][1] = dst->vec[1][1] = dst->vec[2][1] = 1.0f;
|
||||
}
|
||||
else {
|
||||
|
|
|
@ -1282,8 +1282,8 @@ BLI_INLINE uint layer_bucket_index_from_xy(MaskRasterLayer *layer, const float x
|
|||
{
|
||||
BLI_assert(BLI_rctf_isect_pt_v(&layer->bounds, xy));
|
||||
|
||||
return ((uint)((xy[0] - layer->bounds.xmin) * layer->buckets_xy_scalar[0])) +
|
||||
(((uint)((xy[1] - layer->bounds.ymin) * layer->buckets_xy_scalar[1])) * layer->buckets_x);
|
||||
return (uint)((xy[0] - layer->bounds.xmin) * layer->buckets_xy_scalar[0]) +
|
||||
((uint)((xy[1] - layer->bounds.ymin) * layer->buckets_xy_scalar[1]) * layer->buckets_x);
|
||||
}
|
||||
|
||||
static float layer_bucket_depth_from_xy(MaskRasterLayer *layer, const float xy[2])
|
||||
|
|
|
@ -297,7 +297,7 @@ bool BKE_mball_is_basis(const Object *ob)
|
|||
|
||||
/* Just a quick test. */
|
||||
const int len = strlen(ob->id.name);
|
||||
return (!isdigit(ob->id.name[len - 1]));
|
||||
return !isdigit(ob->id.name[len - 1]);
|
||||
}
|
||||
|
||||
bool BKE_mball_is_same_group(const Object *ob1, const Object *ob2)
|
||||
|
|
|
@ -431,8 +431,7 @@ NlaStrip *BKE_nlastrip_new(bAction *act)
|
|||
BKE_action_get_frame_range(strip->act, &strip->actstart, &strip->actend);
|
||||
|
||||
strip->start = strip->actstart;
|
||||
strip->end = (IS_EQF(strip->actstart, strip->actend)) ? (strip->actstart + 1.0f) :
|
||||
(strip->actend);
|
||||
strip->end = IS_EQF(strip->actstart, strip->actend) ? (strip->actstart + 1.0f) : strip->actend;
|
||||
|
||||
/* strip should be referenced as-is */
|
||||
strip->scale = 1.0f;
|
||||
|
|
|
@ -3352,7 +3352,7 @@ void BKE_object_get_parent_matrix(Object *ob, Object *par, float r_parentmat[4][
|
|||
case PAROBJECT: {
|
||||
bool ok = false;
|
||||
if (par->type == OB_CURVES_LEGACY) {
|
||||
if ((((Curve *)par->data)->flag & CU_PATH) && (ob_parcurve(ob, par, tmat))) {
|
||||
if ((((Curve *)par->data)->flag & CU_PATH) && ob_parcurve(ob, par, tmat)) {
|
||||
ok = true;
|
||||
}
|
||||
}
|
||||
|
|
|
@ -1850,7 +1850,7 @@ static bool find_rna_property_rgba(PointerRNA *id_ptr, const char *name, float r
|
|||
value = RNA_property_float_get(&ptr, prop);
|
||||
}
|
||||
else if (type == PROP_INT) {
|
||||
value = static_cast<float>(RNA_property_int_get(&ptr, prop));
|
||||
value = float(RNA_property_int_get(&ptr, prop));
|
||||
}
|
||||
else if (type == PROP_BOOLEAN) {
|
||||
value = RNA_property_boolean_get(&ptr, prop) ? 1.0f : 0.0f;
|
||||
|
@ -1873,7 +1873,7 @@ static bool find_rna_property_rgba(PointerRNA *id_ptr, const char *name, float r
|
|||
int tmp[4] = {0, 0, 0, 1};
|
||||
RNA_property_int_get_array(&ptr, prop, tmp);
|
||||
for (int i = 0; i < 4; i++) {
|
||||
r_data[i] = static_cast<float>(tmp[i]);
|
||||
r_data[i] = float(tmp[i]);
|
||||
}
|
||||
return true;
|
||||
}
|
||||
|
|
|
@ -55,8 +55,8 @@ static float gaussRand(RNG *rng)
|
|||
float length2;
|
||||
|
||||
do {
|
||||
x = (float)(nextfr(rng, -1, 1));
|
||||
y = (float)(nextfr(rng, -1, 1));
|
||||
x = (float)nextfr(rng, -1, 1);
|
||||
y = (float)nextfr(rng, -1, 1);
|
||||
length2 = x * x + y * y;
|
||||
} while (length2 >= 1 || length2 == 0);
|
||||
|
||||
|
@ -930,40 +930,37 @@ bool BKE_ocean_init(struct Ocean *o,
|
|||
case MOD_OCEAN_SPECTRUM_JONSWAP:
|
||||
mul_complex_f(o->_h0[i * o->_N + j],
|
||||
r1r2,
|
||||
(float)(sqrt(BLI_ocean_spectrum_jonswap(o, o->_kx[i], o->_kz[j]) / 2.0f)));
|
||||
mul_complex_f(
|
||||
o->_h0_minus[i * o->_N + j],
|
||||
r1r2,
|
||||
(float)(sqrt(BLI_ocean_spectrum_jonswap(o, -o->_kx[i], -o->_kz[j]) / 2.0f)));
|
||||
(float)sqrt(BLI_ocean_spectrum_jonswap(o, o->_kx[i], o->_kz[j]) / 2.0f));
|
||||
mul_complex_f(o->_h0_minus[i * o->_N + j],
|
||||
r1r2,
|
||||
(float)sqrt(BLI_ocean_spectrum_jonswap(o, -o->_kx[i], -o->_kz[j]) / 2.0f));
|
||||
break;
|
||||
case MOD_OCEAN_SPECTRUM_TEXEL_MARSEN_ARSLOE:
|
||||
mul_complex_f(
|
||||
o->_h0[i * o->_N + j],
|
||||
r1r2,
|
||||
(float)(sqrt(BLI_ocean_spectrum_texelmarsenarsloe(o, o->_kx[i], o->_kz[j]) / 2.0f)));
|
||||
(float)sqrt(BLI_ocean_spectrum_texelmarsenarsloe(o, o->_kx[i], o->_kz[j]) / 2.0f));
|
||||
mul_complex_f(
|
||||
o->_h0_minus[i * o->_N + j],
|
||||
r1r2,
|
||||
(float)(sqrt(BLI_ocean_spectrum_texelmarsenarsloe(o, -o->_kx[i], -o->_kz[j]) /
|
||||
2.0f)));
|
||||
(float)sqrt(BLI_ocean_spectrum_texelmarsenarsloe(o, -o->_kx[i], -o->_kz[j]) / 2.0f));
|
||||
break;
|
||||
case MOD_OCEAN_SPECTRUM_PIERSON_MOSKOWITZ:
|
||||
mul_complex_f(
|
||||
o->_h0[i * o->_N + j],
|
||||
r1r2,
|
||||
(float)(sqrt(BLI_ocean_spectrum_piersonmoskowitz(o, o->_kx[i], o->_kz[j]) / 2.0f)));
|
||||
(float)sqrt(BLI_ocean_spectrum_piersonmoskowitz(o, o->_kx[i], o->_kz[j]) / 2.0f));
|
||||
mul_complex_f(
|
||||
o->_h0_minus[i * o->_N + j],
|
||||
r1r2,
|
||||
(float)(sqrt(BLI_ocean_spectrum_piersonmoskowitz(o, -o->_kx[i], -o->_kz[j]) /
|
||||
2.0f)));
|
||||
(float)sqrt(BLI_ocean_spectrum_piersonmoskowitz(o, -o->_kx[i], -o->_kz[j]) / 2.0f));
|
||||
break;
|
||||
default:
|
||||
mul_complex_f(
|
||||
o->_h0[i * o->_N + j], r1r2, (float)(sqrt(Ph(o, o->_kx[i], o->_kz[j]) / 2.0f)));
|
||||
o->_h0[i * o->_N + j], r1r2, (float)sqrt(Ph(o, o->_kx[i], o->_kz[j]) / 2.0f));
|
||||
mul_complex_f(o->_h0_minus[i * o->_N + j],
|
||||
r1r2,
|
||||
(float)(sqrt(Ph(o, -o->_kx[i], -o->_kz[j]) / 2.0f)));
|
||||
(float)sqrt(Ph(o, -o->_kx[i], -o->_kz[j]) / 2.0f));
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
|
|
@ -2462,7 +2462,7 @@ bool do_guides(Depsgraph *depsgraph,
|
|||
if (guidetime != 0.0f) {
|
||||
/* curve direction */
|
||||
cross_v3_v3v3(temp, eff->guide_dir, guidedir);
|
||||
angle = dot_v3v3(eff->guide_dir, guidedir) / (len_v3(eff->guide_dir));
|
||||
angle = dot_v3v3(eff->guide_dir, guidedir) / len_v3(eff->guide_dir);
|
||||
angle = saacos(angle);
|
||||
axis_angle_to_quat(rot2, temp, angle);
|
||||
mul_qt_v3(rot2, vec_to_point);
|
||||
|
@ -2987,8 +2987,7 @@ static void psys_thread_create_path(ParticleTask *task,
|
|||
* pa->num, pa->fuv,
|
||||
* NULL);
|
||||
*/
|
||||
cpa_num = (ELEM(pa->num_dmcache, DMCACHE_ISCHILD, DMCACHE_NOTFOUND)) ? pa->num :
|
||||
pa->num_dmcache;
|
||||
cpa_num = ELEM(pa->num_dmcache, DMCACHE_ISCHILD, DMCACHE_NOTFOUND) ? pa->num : pa->num_dmcache;
|
||||
|
||||
/* XXX hack to avoid messed up particle num and subsequent crash (T40733) */
|
||||
if (cpa_num > ctx->sim.psmd->mesh_final->totface) {
|
||||
|
@ -3849,7 +3848,7 @@ static void psys_face_mat(Object *ob, Mesh *mesh, ParticleData *pa, float mat[4]
|
|||
MFace *mface;
|
||||
const float(*orcodata)[3];
|
||||
|
||||
int i = (ELEM(pa->num_dmcache, DMCACHE_ISCHILD, DMCACHE_NOTFOUND)) ? pa->num : pa->num_dmcache;
|
||||
int i = ELEM(pa->num_dmcache, DMCACHE_ISCHILD, DMCACHE_NOTFOUND) ? pa->num : pa->num_dmcache;
|
||||
if (i == -1 || i >= mesh->totface) {
|
||||
unit_m4(mat);
|
||||
return;
|
||||
|
|
|
@ -1096,7 +1096,7 @@ static int psys_thread_context_init_distribute(ParticleThreadContext *ctx,
|
|||
maxweight /= totarea;
|
||||
}
|
||||
else {
|
||||
float min = 1.0f / (float)(MIN2(totelem, totpart));
|
||||
float min = 1.0f / (float)MIN2(totelem, totpart);
|
||||
for (i = 0; i < totelem; i++) {
|
||||
element_weight[i] = min;
|
||||
}
|
||||
|
|
|
@ -603,7 +603,7 @@ static void initialize_all_particles(ParticleSimulationData *sim)
|
|||
* UNEXIST flag.
|
||||
*/
|
||||
const bool emit_from_volume_grid = (part->distr == PART_DISTR_GRID) &&
|
||||
(!ELEM(part->from, PART_FROM_VERT, PART_FROM_CHILD));
|
||||
!ELEM(part->from, PART_FROM_VERT, PART_FROM_CHILD);
|
||||
PARTICLE_P;
|
||||
LOOP_PARTICLES
|
||||
{
|
||||
|
@ -4151,17 +4151,17 @@ static bool particles_has_tracer(short parttype)
|
|||
|
||||
static bool particles_has_spray(short parttype)
|
||||
{
|
||||
return (ELEM(parttype, PART_FLUID_SPRAY, PART_FLUID_SPRAYFOAM, PART_FLUID_SPRAYFOAMBUBBLE));
|
||||
return ELEM(parttype, PART_FLUID_SPRAY, PART_FLUID_SPRAYFOAM, PART_FLUID_SPRAYFOAMBUBBLE);
|
||||
}
|
||||
|
||||
static bool particles_has_bubble(short parttype)
|
||||
{
|
||||
return (ELEM(parttype, PART_FLUID_BUBBLE, PART_FLUID_FOAMBUBBLE, PART_FLUID_SPRAYFOAMBUBBLE));
|
||||
return ELEM(parttype, PART_FLUID_BUBBLE, PART_FLUID_FOAMBUBBLE, PART_FLUID_SPRAYFOAMBUBBLE);
|
||||
}
|
||||
|
||||
static bool particles_has_foam(short parttype)
|
||||
{
|
||||
return (ELEM(parttype, PART_FLUID_FOAM, PART_FLUID_SPRAYFOAM, PART_FLUID_SPRAYFOAMBUBBLE));
|
||||
return ELEM(parttype, PART_FLUID_FOAM, PART_FLUID_SPRAYFOAM, PART_FLUID_SPRAYFOAMBUBBLE);
|
||||
}
|
||||
|
||||
static void particles_fluid_step(ParticleSimulationData *sim,
|
||||
|
|
|
@ -199,7 +199,7 @@ static ShrinkwrapBoundaryData *shrinkwrap_build_boundary_data(Mesh *mesh)
|
|||
|
||||
/* Count faces per edge (up to 2). */
|
||||
char *edge_mode = static_cast<char *>(
|
||||
MEM_calloc_arrayN((size_t)mesh->totedge, sizeof(char), __func__));
|
||||
MEM_calloc_arrayN(size_t(mesh->totedge), sizeof(char), __func__));
|
||||
|
||||
for (int i = 0; i < mesh->totloop; i++) {
|
||||
uint eidx = mloop[i].e;
|
||||
|
@ -258,7 +258,7 @@ static ShrinkwrapBoundaryData *shrinkwrap_build_boundary_data(Mesh *mesh)
|
|||
|
||||
/* Find boundary vertices and build a mapping table for compact storage of data. */
|
||||
int *vert_boundary_id = static_cast<int *>(
|
||||
MEM_calloc_arrayN((size_t)mesh->totvert, sizeof(int), __func__));
|
||||
MEM_calloc_arrayN(size_t(mesh->totvert), sizeof(int), __func__));
|
||||
|
||||
for (int i = 0; i < mesh->totedge; i++) {
|
||||
if (edge_mode[i]) {
|
||||
|
@ -272,7 +272,7 @@ static ShrinkwrapBoundaryData *shrinkwrap_build_boundary_data(Mesh *mesh)
|
|||
uint num_boundary_verts = 0;
|
||||
|
||||
for (int i = 0; i < mesh->totvert; i++) {
|
||||
vert_boundary_id[i] = (vert_boundary_id[i] != 0) ? (int)num_boundary_verts++ : -1;
|
||||
vert_boundary_id[i] = (vert_boundary_id[i] != 0) ? int(num_boundary_verts++) : -1;
|
||||
}
|
||||
|
||||
data->vert_boundary_id = vert_boundary_id;
|
||||
|
|
|
@ -217,8 +217,8 @@ static void studiolight_load_solid_light(StudioLight *sl)
|
|||
#undef READ_IVAL
|
||||
#undef READ_FVAL
|
||||
|
||||
#define WRITE_FVAL(str, id, val) (BLI_dynstr_appendf(str, id " %f\n", val))
|
||||
#define WRITE_IVAL(str, id, val) (BLI_dynstr_appendf(str, id " %d\n", val))
|
||||
#define WRITE_FVAL(str, id, val) BLI_dynstr_appendf(str, id " %f\n", val)
|
||||
#define WRITE_IVAL(str, id, val) BLI_dynstr_appendf(str, id " %d\n", val)
|
||||
|
||||
#define WRITE_VEC3(str, id, val) \
|
||||
do { \
|
||||
|
@ -273,7 +273,7 @@ static void direction_to_equirect(float r[2], const float dir[3])
|
|||
|
||||
static void equirect_to_direction(float r[3], float u, float v)
|
||||
{
|
||||
float phi = (-(M_PI * 2)) * u + M_PI;
|
||||
float phi = -(M_PI * 2) * u + M_PI;
|
||||
float theta = -M_PI * v + M_PI;
|
||||
float sin_theta = sinf(theta);
|
||||
r[0] = sin_theta * cosf(phi);
|
||||
|
|
|
@ -128,9 +128,9 @@ static void subdiv_ccg_alloc_elements(SubdivCCG *subdiv_ccg, Subdiv *subdiv)
|
|||
subdiv_ccg->num_grids = num_grids;
|
||||
subdiv_ccg->grids = static_cast<CCGElem **>(
|
||||
MEM_calloc_arrayN(num_grids, sizeof(CCGElem *), "subdiv ccg grids"));
|
||||
subdiv_ccg->grids_storage = static_cast<unsigned char *>(MEM_calloc_arrayN(
|
||||
num_grids, ((size_t)grid_area) * element_size, "subdiv ccg grids storage"));
|
||||
const size_t grid_size_in_bytes = (size_t)grid_area * element_size;
|
||||
subdiv_ccg->grids_storage = static_cast<unsigned char *>(
|
||||
MEM_calloc_arrayN(num_grids, size_t(grid_area) * element_size, "subdiv ccg grids storage"));
|
||||
const size_t grid_size_in_bytes = size_t(grid_area) * element_size;
|
||||
for (int grid_index = 0; grid_index < num_grids; grid_index++) {
|
||||
const size_t grid_offset = grid_size_in_bytes * grid_index;
|
||||
subdiv_ccg->grids[grid_index] = (CCGElem *)&subdiv_ccg->grids_storage[grid_offset];
|
||||
|
@ -241,7 +241,7 @@ static void subdiv_ccg_eval_regular_grid(CCGEvalGridsData *data, const int face_
|
|||
const float grid_u = x * grid_size_1_inv;
|
||||
float u, v;
|
||||
BKE_subdiv_rotate_grid_to_quad(corner, grid_u, grid_v, &u, &v);
|
||||
const size_t grid_element_index = (size_t)y * grid_size + x;
|
||||
const size_t grid_element_index = size_t(y) * grid_size + x;
|
||||
const size_t grid_element_offset = grid_element_index * element_size;
|
||||
subdiv_ccg_eval_grid_element(data, ptex_face_index, u, v, &grid[grid_element_offset]);
|
||||
}
|
||||
|
@ -271,7 +271,7 @@ static void subdiv_ccg_eval_special_grid(CCGEvalGridsData *data, const int face_
|
|||
const float u = 1.0f - (y * grid_size_1_inv);
|
||||
for (int x = 0; x < grid_size; x++) {
|
||||
const float v = 1.0f - (x * grid_size_1_inv);
|
||||
const size_t grid_element_index = (size_t)y * grid_size + x;
|
||||
const size_t grid_element_index = size_t(y) * grid_size + x;
|
||||
const size_t grid_element_offset = grid_element_index * element_size;
|
||||
subdiv_ccg_eval_grid_element(data, ptex_face_index, u, v, &grid[grid_element_offset]);
|
||||
}
|
||||
|
|
|
@ -2660,9 +2660,9 @@ ImBuf *BKE_tracking_sample_pattern(int frame_width,
|
|||
if (from_anchor) {
|
||||
for (int a = 0; a < 5; a++) {
|
||||
src_pixel_x[a] += (double)((track->offset[0] * frame_width) -
|
||||
((int)(track->offset[0] * frame_width)));
|
||||
(int)(track->offset[0] * frame_width));
|
||||
src_pixel_y[a] += (double)((track->offset[1] * frame_height) -
|
||||
((int)(track->offset[1] * frame_height)));
|
||||
(int)(track->offset[1] * frame_height));
|
||||
|
||||
/* when offset is negative, rounding happens in opposite direction */
|
||||
if (track->offset[0] < 0.0f) {
|
||||
|
|
|
@ -1044,7 +1044,7 @@ static bool vfont_to_curve(Object *ob,
|
|||
|
||||
CLAMP_MIN(maxlen, lineinfo[lnr].x_min);
|
||||
|
||||
if ((tb_scale.h != 0.0f) && ((-(yof - tb_scale.y)) > (tb_scale.h - linedist) - yof_scale)) {
|
||||
if ((tb_scale.h != 0.0f) && (-(yof - tb_scale.y) > (tb_scale.h - linedist) - yof_scale)) {
|
||||
if (cu->totbox > (curbox + 1)) {
|
||||
maxlen = 0;
|
||||
curbox++;
|
||||
|
@ -1175,7 +1175,7 @@ static bool vfont_to_curve(Object *ob,
|
|||
}
|
||||
}
|
||||
for (i = 0; i <= slen; i++) {
|
||||
for (j = i; (!ELEM(mem[j], '\0', '\n')) && (chartransdata[j].dobreak == 0) && (j < slen);
|
||||
for (j = i; !ELEM(mem[j], '\0', '\n') && (chartransdata[j].dobreak == 0) && (j < slen);
|
||||
j++) {
|
||||
/* do nothing */
|
||||
}
|
||||
|
|
|
@ -34,10 +34,6 @@
|
|||
#include "DNA_packedFile_types.h"
|
||||
#include "DNA_vfont_types.h"
|
||||
|
||||
/* local variables */
|
||||
static FT_Library library;
|
||||
static FT_Error err;
|
||||
|
||||
static VChar *freetypechar_to_vchar(FT_Face face, FT_ULong charcode, VFontData *vfd)
|
||||
{
|
||||
const float scale = vfd->scale;
|
||||
|
@ -60,7 +56,7 @@ static VChar *freetypechar_to_vchar(FT_Face face, FT_ULong charcode, VFontData *
|
|||
*
|
||||
* Get the FT Glyph index and load the Glyph */
|
||||
glyph_index = FT_Get_Char_Index(face, charcode);
|
||||
err = FT_Load_Glyph(face, glyph_index, FT_LOAD_NO_SCALE | FT_LOAD_NO_BITMAP);
|
||||
FT_Error err = FT_Load_Glyph(face, glyph_index, FT_LOAD_NO_SCALE | FT_LOAD_NO_BITMAP);
|
||||
|
||||
/* If loading succeeded, convert the FT glyph to the internal format */
|
||||
if (!err) {
|
||||
|
@ -240,7 +236,7 @@ static VChar *freetypechar_to_vchar(FT_Face face, FT_ULong charcode, VFontData *
|
|||
return NULL;
|
||||
}
|
||||
|
||||
static VChar *objchr_to_ftvfontdata(VFont *vfont, FT_ULong charcode)
|
||||
static VChar *objchr_to_ftvfontdata(FT_Library library, VFont *vfont, FT_ULong charcode)
|
||||
{
|
||||
VChar *che;
|
||||
|
||||
|
@ -249,13 +245,13 @@ static VChar *objchr_to_ftvfontdata(VFont *vfont, FT_ULong charcode)
|
|||
|
||||
/* Load the font to memory */
|
||||
if (vfont->temp_pf) {
|
||||
err = FT_New_Memory_Face(library, vfont->temp_pf->data, vfont->temp_pf->size, 0, &face);
|
||||
FT_Error err = FT_New_Memory_Face(
|
||||
library, vfont->temp_pf->data, vfont->temp_pf->size, 0, &face);
|
||||
if (err) {
|
||||
return NULL;
|
||||
}
|
||||
}
|
||||
else {
|
||||
err = true;
|
||||
return NULL;
|
||||
}
|
||||
|
||||
|
@ -266,7 +262,7 @@ static VChar *objchr_to_ftvfontdata(VFont *vfont, FT_ULong charcode)
|
|||
return che;
|
||||
}
|
||||
|
||||
static FT_Face vfont_face_load_from_packed_file(PackedFile *pf)
|
||||
static FT_Face vfont_face_load_from_packed_file(FT_Library library, PackedFile *pf)
|
||||
{
|
||||
FT_Face face = NULL;
|
||||
FT_New_Memory_Face(library, pf->data, pf->size, 0, &face);
|
||||
|
@ -281,14 +277,14 @@ static FT_Face vfont_face_load_from_packed_file(PackedFile *pf)
|
|||
}
|
||||
|
||||
/* Select a character map. */
|
||||
FT_Error err_charmap = FT_Select_Charmap(face, FT_ENCODING_UNICODE);
|
||||
if (err_charmap) {
|
||||
err_charmap = FT_Select_Charmap(face, FT_ENCODING_APPLE_ROMAN);
|
||||
FT_Error err = FT_Select_Charmap(face, FT_ENCODING_UNICODE);
|
||||
if (err) {
|
||||
err = FT_Select_Charmap(face, FT_ENCODING_APPLE_ROMAN);
|
||||
}
|
||||
if (err_charmap && face->num_charmaps > 0) {
|
||||
err_charmap = FT_Select_Charmap(face, face->charmaps[0]->encoding);
|
||||
if (err && face->num_charmaps > 0) {
|
||||
err = FT_Select_Charmap(face, face->charmaps[0]->encoding);
|
||||
}
|
||||
if (err_charmap) {
|
||||
if (err) {
|
||||
FT_Done_Face(face);
|
||||
return NULL;
|
||||
}
|
||||
|
@ -307,11 +303,12 @@ static FT_Face vfont_face_load_from_packed_file(PackedFile *pf)
|
|||
|
||||
VFontData *BKE_vfontdata_from_freetypefont(PackedFile *pf)
|
||||
{
|
||||
FT_Library library = NULL;
|
||||
if (FT_Init_FreeType(&library) != FT_Err_Ok) {
|
||||
return NULL;
|
||||
}
|
||||
|
||||
FT_Face face = vfont_face_load_from_packed_file(pf);
|
||||
FT_Face face = vfont_face_load_from_packed_file(library, pf);
|
||||
if (!face) {
|
||||
FT_Done_FreeType(library);
|
||||
return NULL;
|
||||
|
@ -399,14 +396,15 @@ VChar *BKE_vfontdata_char_from_freetypefont(VFont *vfont, ulong character)
|
|||
}
|
||||
|
||||
/* Init Freetype */
|
||||
err = FT_Init_FreeType(&library);
|
||||
FT_Library library = NULL;
|
||||
FT_Error err = FT_Init_FreeType(&library);
|
||||
if (err) {
|
||||
/* XXX error("Failed to load the Freetype font library"); */
|
||||
return NULL;
|
||||
}
|
||||
|
||||
/* Load the character */
|
||||
che = objchr_to_ftvfontdata(vfont, character);
|
||||
che = objchr_to_ftvfontdata(library, vfont, character);
|
||||
|
||||
/* Free Freetype */
|
||||
FT_Done_FreeType(library);
|
||||
|
|
|
@ -80,7 +80,7 @@ static bool keycmp(const void *a, const void *b)
|
|||
if (ka->case_str == 1 || kb->case_str == 1) {
|
||||
return (BLI_strcasecmp(ka->arg, kb->arg) != 0);
|
||||
}
|
||||
return (!STREQ(ka->arg, kb->arg));
|
||||
return !STREQ(ka->arg, kb->arg);
|
||||
}
|
||||
return BLI_ghashutil_intcmp((const void *)ka->pass, (const void *)kb->pass);
|
||||
}
|
||||
|
|
|
@ -93,7 +93,7 @@ static int bli_compare(struct direntry *entry1, struct direntry *entry2)
|
|||
return 1;
|
||||
}
|
||||
|
||||
return (BLI_strcasecmp_natural(entry1->relname, entry2->relname));
|
||||
return BLI_strcasecmp_natural(entry1->relname, entry2->relname);
|
||||
}
|
||||
|
||||
struct BuildDirCtx {
|
||||
|
|
|
@ -491,7 +491,7 @@ void BLI_box_pack_2d(BoxPack *boxarray, const uint len, float *r_tot_x, float *r
|
|||
* flag verts on one or both of the boxes
|
||||
* as being used by checking the width or
|
||||
* height of both boxes */
|
||||
if (vert->tlb && vert->trb && (ELEM(box, vert->tlb, vert->trb))) {
|
||||
if (vert->tlb && vert->trb && ELEM(box, vert->tlb, vert->trb)) {
|
||||
if (UNLIKELY(fabsf(vert->tlb->h - vert->trb->h) < EPSILON_MERGE)) {
|
||||
#ifdef USE_MERGE
|
||||
# define A (vert->trb->v[TL])
|
||||
|
@ -522,7 +522,7 @@ void BLI_box_pack_2d(BoxPack *boxarray, const uint len, float *r_tot_x, float *r
|
|||
vert->tlb->v[TR]->free &= ~(TRF | BRF);
|
||||
}
|
||||
}
|
||||
else if (vert->blb && vert->brb && (ELEM(box, vert->blb, vert->brb))) {
|
||||
else if (vert->blb && vert->brb && ELEM(box, vert->blb, vert->brb)) {
|
||||
if (UNLIKELY(fabsf(vert->blb->h - vert->brb->h) < EPSILON_MERGE)) {
|
||||
#ifdef USE_MERGE
|
||||
# define A (vert->blb->v[BR])
|
||||
|
@ -554,7 +554,7 @@ void BLI_box_pack_2d(BoxPack *boxarray, const uint len, float *r_tot_x, float *r
|
|||
}
|
||||
}
|
||||
/* Horizontal */
|
||||
if (vert->tlb && vert->blb && (ELEM(box, vert->tlb, vert->blb))) {
|
||||
if (vert->tlb && vert->blb && ELEM(box, vert->tlb, vert->blb)) {
|
||||
if (UNLIKELY(fabsf(vert->tlb->w - vert->blb->w) < EPSILON_MERGE)) {
|
||||
#ifdef USE_MERGE
|
||||
# define A (vert->blb->v[TL])
|
||||
|
@ -585,7 +585,7 @@ void BLI_box_pack_2d(BoxPack *boxarray, const uint len, float *r_tot_x, float *r
|
|||
vert->tlb->v[BL]->free &= ~(BLF | BRF);
|
||||
}
|
||||
}
|
||||
else if (vert->trb && vert->brb && (ELEM(box, vert->trb, vert->brb))) {
|
||||
else if (vert->trb && vert->brb && ELEM(box, vert->trb, vert->brb)) {
|
||||
if (UNLIKELY(fabsf(vert->trb->w - vert->brb->w) < EPSILON_MERGE)) {
|
||||
|
||||
#ifdef USE_MERGE
|
||||
|
|
|
@ -65,7 +65,7 @@ float floor_power_of_10(float f)
|
|||
{
|
||||
BLI_assert(!(f < 0.0f));
|
||||
if (f != 0.0f) {
|
||||
return 1.0f / (powf(10.0f, ceilf(log10f(1.0f / f))));
|
||||
return 1.0f / powf(10.0f, ceilf(log10f(1.0f / f)));
|
||||
}
|
||||
return 0.0f;
|
||||
}
|
||||
|
@ -74,7 +74,7 @@ float ceil_power_of_10(float f)
|
|||
{
|
||||
BLI_assert(!(f < 0.0f));
|
||||
if (f != 0.0f) {
|
||||
return 1.0f / (powf(10.0f, floorf(log10f(1.0f / f))));
|
||||
return 1.0f / powf(10.0f, floorf(log10f(1.0f / f)));
|
||||
}
|
||||
return 0.0f;
|
||||
}
|
||||
|
|
|
@ -368,9 +368,9 @@ uint rgb_to_cpack(float r, float g, float b)
|
|||
|
||||
void cpack_to_rgb(uint col, float *r_r, float *r_g, float *r_b)
|
||||
{
|
||||
*r_r = ((float)(col & 0xFF)) * (1.0f / 255.0f);
|
||||
*r_g = ((float)((col >> 8) & 0xFF)) * (1.0f / 255.0f);
|
||||
*r_b = ((float)((col >> 16) & 0xFF)) * (1.0f / 255.0f);
|
||||
*r_r = (float)(col & 0xFF) * (1.0f / 255.0f);
|
||||
*r_g = (float)((col >> 8) & 0xFF) * (1.0f / 255.0f);
|
||||
*r_b = (float)((col >> 16) & 0xFF) * (1.0f / 255.0f);
|
||||
}
|
||||
|
||||
void rgb_uchar_to_float(float r_col[3], const uchar col_ub[3])
|
||||
|
|
|
@ -2748,7 +2748,7 @@ bool isect_sweeping_sphere_tri_v3(const float p1[3],
|
|||
edotv = dot_v3v3(e1, vel);
|
||||
edotbv = dot_v3v3(e1, bv);
|
||||
|
||||
a = elen2 * (-dot_v3v3(vel, vel)) + edotv * edotv;
|
||||
a = elen2 * -dot_v3v3(vel, vel) + edotv * edotv;
|
||||
b = 2.0f * (elen2 * dot_v3v3(vel, bv) - edotv * edotbv);
|
||||
c = elen2 * (radius2 - dot_v3v3(bv, bv)) + edotbv * edotbv;
|
||||
|
||||
|
@ -2770,7 +2770,7 @@ bool isect_sweeping_sphere_tri_v3(const float p1[3],
|
|||
edotv = dot_v3v3(e2, vel);
|
||||
edotbv = dot_v3v3(e2, bv);
|
||||
|
||||
a = elen2 * (-dot_v3v3(vel, vel)) + edotv * edotv;
|
||||
a = elen2 * -dot_v3v3(vel, vel) + edotv * edotv;
|
||||
b = 2.0f * (elen2 * dot_v3v3(vel, bv) - edotv * edotbv);
|
||||
c = elen2 * (radius2 - dot_v3v3(bv, bv)) + edotbv * edotbv;
|
||||
|
||||
|
@ -2797,7 +2797,7 @@ bool isect_sweeping_sphere_tri_v3(const float p1[3],
|
|||
edotv = dot_v3v3(e3, vel);
|
||||
edotbv = dot_v3v3(e3, bv);
|
||||
|
||||
a = elen2 * (-dot_v3v3(vel, vel)) + edotv * edotv;
|
||||
a = elen2 * -dot_v3v3(vel, vel) + edotv * edotv;
|
||||
b = 2.0f * (elen2 * dot_v3v3(vel, bv) - edotv * edotbv);
|
||||
c = elen2 * (radius2 - dot_v3v3(bv, bv)) + edotbv * edotbv;
|
||||
|
||||
|
|
|
@ -624,10 +624,10 @@ void BLI_ewa_filter(const int width,
|
|||
|
||||
U0 = uv[0] * (float)width;
|
||||
V0 = uv[1] * (float)height;
|
||||
u1 = (int)(floorf(U0 - ue));
|
||||
u2 = (int)(ceilf(U0 + ue));
|
||||
v1 = (int)(floorf(V0 - ve));
|
||||
v2 = (int)(ceilf(V0 + ve));
|
||||
u1 = (int)floorf(U0 - ue);
|
||||
u2 = (int)ceilf(U0 + ue);
|
||||
v1 = (int)floorf(V0 - ve);
|
||||
v2 = (int)ceilf(V0 + ve);
|
||||
|
||||
/* sane clamping to avoid unnecessarily huge loops */
|
||||
/* NOTE: if eccentricity gets clamped (see above),
|
||||
|
|
|
@ -2360,8 +2360,8 @@ bool mat3_from_axis_conversion(
|
|||
value = ((src_forward << (0 * 3)) | (src_up << (1 * 3)) | (dst_forward << (2 * 3)) |
|
||||
(dst_up << (3 * 3)));
|
||||
|
||||
for (uint i = 0; i < (ARRAY_SIZE(_axis_convert_matrix)); i++) {
|
||||
for (uint j = 0; j < (ARRAY_SIZE(*_axis_convert_lut)); j++) {
|
||||
for (uint i = 0; i < ARRAY_SIZE(_axis_convert_matrix); i++) {
|
||||
for (uint j = 0; j < ARRAY_SIZE(*_axis_convert_lut); j++) {
|
||||
if (_axis_convert_lut[i][j] == value) {
|
||||
copy_m3_m3(r_mat, _axis_convert_matrix[i]);
|
||||
return true;
|
||||
|
|
|
@ -499,7 +499,7 @@ float angle_signed_on_axis_v3v3_v3(const float v1[3], const float v2[3], const f
|
|||
/* calculate the sign (reuse 'tproj') */
|
||||
cross_v3_v3v3(tproj, v2_proj, v1_proj);
|
||||
if (dot_v3v3(tproj, axis) < 0.0f) {
|
||||
angle = ((float)(M_PI * 2.0)) - angle;
|
||||
angle = (float)(M_PI * 2.0) - angle;
|
||||
}
|
||||
|
||||
return angle;
|
||||
|
|
|
@ -939,9 +939,9 @@ void BLI_noise_voronoi(float x, float y, float z, float *da, float *pa, float me
|
|||
break;
|
||||
}
|
||||
|
||||
int xi = (int)(floor(x));
|
||||
int yi = (int)(floor(y));
|
||||
int zi = (int)(floor(z));
|
||||
int xi = (int)floor(x);
|
||||
int yi = (int)floor(y);
|
||||
int zi = (int)floor(z);
|
||||
da[0] = da[1] = da[2] = da[3] = 1e10f;
|
||||
for (int xx = xi - 1; xx <= xi + 1; xx++) {
|
||||
for (int yy = yi - 1; yy <= yi + 1; yy++) {
|
||||
|
@ -1112,9 +1112,9 @@ static float BLI_cellNoiseU(float x, float y, float z)
|
|||
y = (y + 0.000001f) * 1.00001f;
|
||||
z = (z + 0.000001f) * 1.00001f;
|
||||
|
||||
int xi = (int)(floor(x));
|
||||
int yi = (int)(floor(y));
|
||||
int zi = (int)(floor(z));
|
||||
int xi = (int)floor(x);
|
||||
int yi = (int)floor(y);
|
||||
int zi = (int)floor(z);
|
||||
uint n = xi + yi * 1301 + zi * 314159;
|
||||
n ^= (n << 13);
|
||||
return ((float)(n * (n * n * 15731 + 789221) + 1376312589) / 4294967296.0f);
|
||||
|
@ -1132,9 +1132,9 @@ void BLI_noise_cell_v3(float x, float y, float z, float r_ca[3])
|
|||
y = (y + 0.000001f) * 1.00001f;
|
||||
z = (z + 0.000001f) * 1.00001f;
|
||||
|
||||
int xi = (int)(floor(x));
|
||||
int yi = (int)(floor(y));
|
||||
int zi = (int)(floor(z));
|
||||
int xi = (int)floor(x);
|
||||
int yi = (int)floor(y);
|
||||
int zi = (int)floor(z);
|
||||
const float *p = HASHPNT(xi, yi, zi);
|
||||
r_ca[0] = p[0];
|
||||
r_ca[1] = p[1];
|
||||
|
|
|
@ -1347,7 +1347,7 @@ bool BLI_path_extension_ensure(char *path, size_t maxlen, const char *ext)
|
|||
ssize_t a;
|
||||
|
||||
/* first check the extension is already there */
|
||||
if ((ext_len <= path_len) && (STREQ(path + (path_len - ext_len), ext))) {
|
||||
if ((ext_len <= path_len) && STREQ(path + (path_len - ext_len), ext)) {
|
||||
return true;
|
||||
}
|
||||
|
||||
|
|
|
@ -373,12 +373,12 @@ static bool kdtree2d_isect_tri_recursive(const struct KDTree2D *tree,
|
|||
|
||||
# define KDTREE2D_ISECT_TRI_RECURSE_NEG \
|
||||
(((node->neg != KDNODE_UNSET) && (co[node->axis] >= bounds[node->axis].min)) && \
|
||||
(kdtree2d_isect_tri_recursive( \
|
||||
tree, tri_index, tri_coords, tri_center, bounds, &tree->nodes[node->neg])))
|
||||
kdtree2d_isect_tri_recursive( \
|
||||
tree, tri_index, tri_coords, tri_center, bounds, &tree->nodes[node->neg]))
|
||||
# define KDTREE2D_ISECT_TRI_RECURSE_POS \
|
||||
(((node->pos != KDNODE_UNSET) && (co[node->axis] <= bounds[node->axis].max)) && \
|
||||
(kdtree2d_isect_tri_recursive( \
|
||||
tree, tri_index, tri_coords, tri_center, bounds, &tree->nodes[node->pos])))
|
||||
kdtree2d_isect_tri_recursive( \
|
||||
tree, tri_index, tri_coords, tri_center, bounds, &tree->nodes[node->pos]))
|
||||
|
||||
if (tri_center[node->axis] > co[node->axis]) {
|
||||
if (KDTREE2D_ISECT_TRI_RECURSE_POS) {
|
||||
|
|
|
@ -314,7 +314,7 @@ size_t BLI_str_unescape_ex(char *__restrict dst,
|
|||
break;
|
||||
}
|
||||
char c = *src;
|
||||
if (UNLIKELY(c == '\\') && (str_unescape_pair(*(src + 1), &c))) {
|
||||
if (UNLIKELY(c == '\\') && str_unescape_pair(*(src + 1), &c)) {
|
||||
src++;
|
||||
}
|
||||
dst[len++] = c;
|
||||
|
@ -329,7 +329,7 @@ size_t BLI_str_unescape(char *__restrict dst, const char *__restrict src, const
|
|||
size_t len = 0;
|
||||
for (const char *src_end = src + src_maxncpy; (src < src_end) && *src; src++) {
|
||||
char c = *src;
|
||||
if (UNLIKELY(c == '\\') && (str_unescape_pair(*(src + 1), &c))) {
|
||||
if (UNLIKELY(c == '\\') && str_unescape_pair(*(src + 1), &c)) {
|
||||
src++;
|
||||
}
|
||||
dst[len++] = c;
|
||||
|
|
|
@ -176,7 +176,7 @@ size_t BLI_timecode_string_from_time_simple(char *str,
|
|||
const int hr = ((int)time_seconds) / (60 * 60);
|
||||
const int min = (((int)time_seconds) / 60) % 60;
|
||||
const int sec = ((int)time_seconds) % 60;
|
||||
const int hun = ((int)(fmod(time_seconds, 1.0) * 100));
|
||||
const int hun = (int)(fmod(time_seconds, 1.0) * 100);
|
||||
|
||||
if (hr) {
|
||||
rlen = BLI_snprintf_rlen(str, maxncpy, "%.2d:%.2d:%.2d.%.2d", hr, min, sec, hun);
|
||||
|
|
|
@ -130,7 +130,7 @@ TEST(linear_allocator, ManyAllocations)
|
|||
RandomNumberGenerator rng;
|
||||
for (int i = 0; i < 1000; i++) {
|
||||
int size = rng.get_int32(10000);
|
||||
int alignment = 1 << (rng.get_int32(7));
|
||||
int alignment = 1 << rng.get_int32(7);
|
||||
void *buffer = allocator.allocate(size, alignment);
|
||||
EXPECT_NE(buffer, nullptr);
|
||||
}
|
||||
|
|
|
@ -1112,7 +1112,7 @@ static int *read_file_thumbnail(FileData *fd)
|
|||
const bool do_endian_swap = (fd->flags & FD_FLAGS_SWITCH_ENDIAN) != 0;
|
||||
int *data = (int *)(bhead + 1);
|
||||
|
||||
if (bhead->len < (sizeof(int[2]))) {
|
||||
if (bhead->len < sizeof(int[2])) {
|
||||
break;
|
||||
}
|
||||
|
||||
|
|
|
@ -595,7 +595,7 @@ static void do_versions_fix_annotations(bGPdata *gpd)
|
|||
|
||||
LISTBASE_FOREACH (bGPDframe *, gpf, &gpl->frames) {
|
||||
LISTBASE_FOREACH (bGPDstroke *, gps, &gpf->strokes) {
|
||||
if ((gps->colorname[0] != '\0') && (STREQ(gps->colorname, palcolor->info))) {
|
||||
if ((gps->colorname[0] != '\0') && STREQ(gps->colorname, palcolor->info)) {
|
||||
/* copy color settings */
|
||||
copy_v4_v4(gpl->color, palcolor->color);
|
||||
}
|
||||
|
|
|
@ -1611,8 +1611,8 @@ void blo_do_versions_290(FileData *fd, Library *UNUSED(lib), Main *bmain)
|
|||
}
|
||||
}
|
||||
|
||||
if ((!MAIN_VERSION_ATLEAST(bmain, 292, 14)) ||
|
||||
((bmain->versionfile == 293) && (!MAIN_VERSION_ATLEAST(bmain, 293, 1)))) {
|
||||
if (!MAIN_VERSION_ATLEAST(bmain, 292, 14) ||
|
||||
((bmain->versionfile == 293) && !MAIN_VERSION_ATLEAST(bmain, 293, 1))) {
|
||||
FOREACH_NODETREE_BEGIN (bmain, ntree, id) {
|
||||
if (ntree->type != NTREE_GEOMETRY) {
|
||||
continue;
|
||||
|
|
|
@ -1194,7 +1194,7 @@ static void update_voronoi_node_square_distance(bNodeTree *ntree)
|
|||
NodeTexVoronoi *tex = (NodeTexVoronoi *)node->storage;
|
||||
bNodeSocket *sockDistance = nodeFindSocket(node, SOCK_OUT, "Distance");
|
||||
if (tex->distance == SHD_VORONOI_EUCLIDEAN &&
|
||||
(ELEM(tex->feature, SHD_VORONOI_F1, SHD_VORONOI_F2)) && socket_is_used(sockDistance)) {
|
||||
ELEM(tex->feature, SHD_VORONOI_F1, SHD_VORONOI_F2) && socket_is_used(sockDistance)) {
|
||||
bNode *multiplyNode = nodeAddStaticNode(NULL, ntree, SH_NODE_MATH);
|
||||
multiplyNode->custom1 = NODE_MATH_MULTIPLY;
|
||||
multiplyNode->locx = node->locx + node->width + 20.0f;
|
||||
|
|
|
@ -2531,7 +2531,7 @@ void blo_do_versions_pre250(FileData *fd, Library *lib, Main *bmain)
|
|||
Object *ob;
|
||||
for (ob = bmain->objects.first; ob; ob = ob->id.next) {
|
||||
if (ob->pd) {
|
||||
ob->pd->seed = ((uint)(ceil(PIL_check_seconds_timer())) + 1) % 128;
|
||||
ob->pd->seed = ((uint)ceil(PIL_check_seconds_timer()) + 1) % 128;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
|
|
@ -1780,7 +1780,7 @@ void BM_lnorspace_invalidate(BMesh *bm, const bool do_invalidate_all)
|
|||
|
||||
/* Note that we only handle unselected neighbor vertices here, main loop will take care of
|
||||
* selected ones. */
|
||||
if ((!BM_elem_flag_test(l->prev->v, BM_ELEM_SELECT)) &&
|
||||
if (!BM_elem_flag_test(l->prev->v, BM_ELEM_SELECT) &&
|
||||
!BLI_BITMAP_TEST(done_verts, BM_elem_index_get(l->prev->v))) {
|
||||
|
||||
BMLoop *l_prev;
|
||||
|
@ -1791,7 +1791,7 @@ void BM_lnorspace_invalidate(BMesh *bm, const bool do_invalidate_all)
|
|||
BLI_BITMAP_ENABLE(done_verts, BM_elem_index_get(l_prev->v));
|
||||
}
|
||||
|
||||
if ((!BM_elem_flag_test(l->next->v, BM_ELEM_SELECT)) &&
|
||||
if (!BM_elem_flag_test(l->next->v, BM_ELEM_SELECT) &&
|
||||
!BLI_BITMAP_TEST(done_verts, BM_elem_index_get(l->next->v))) {
|
||||
|
||||
BMLoop *l_next;
|
||||
|
|
|
@ -52,7 +52,7 @@
|
|||
#define GROW(len_alloc) ((len_alloc) + ((len_alloc) - ((len_alloc) / 2)))
|
||||
#define GROW_ARRAY(mem, len_alloc) \
|
||||
{ \
|
||||
mem = MEM_reallocN(mem, (sizeof(*mem)) * ((len_alloc) = GROW(len_alloc))); \
|
||||
mem = MEM_reallocN(mem, sizeof(*mem) * ((len_alloc) = GROW(len_alloc))); \
|
||||
} \
|
||||
((void)0)
|
||||
|
||||
|
|
|
@ -62,7 +62,7 @@ bool BM_disk_dissolve(BMesh *bm, BMVert *v)
|
|||
e = v->e;
|
||||
do {
|
||||
e = bmesh_disk_edge_next(e, v);
|
||||
if (!(BM_edge_share_face_check(e, v->e))) {
|
||||
if (!BM_edge_share_face_check(e, v->e)) {
|
||||
keepedge = e;
|
||||
baseedge = v->e;
|
||||
break;
|
||||
|
|
|
@ -721,7 +721,7 @@ void bmo_inset_region_exec(BMesh *bm, BMOperator *op)
|
|||
(use_boundary && BM_edge_is_boundary(e) && BM_elem_flag_test(e->l->f, BM_ELEM_TAG)) ||
|
||||
|
||||
/* tag if edge is an interior edge in between a tagged and untagged face */
|
||||
(bm_edge_is_mixed_face_tag(e->l))) {
|
||||
bm_edge_is_mixed_face_tag(e->l)) {
|
||||
/* tag */
|
||||
BM_elem_flag_enable(e->v1, BM_ELEM_TAG);
|
||||
BM_elem_flag_enable(e->v2, BM_ELEM_TAG);
|
||||
|
|
|
@ -43,11 +43,11 @@ static float quad_calc_error(const float v1[3],
|
|||
|
||||
normal_tri_v3(n1, v1, v2, v3);
|
||||
normal_tri_v3(n2, v1, v3, v4);
|
||||
angle_a = (compare_v3v3(n1, n2, FLT_EPSILON)) ? 0.0f : angle_normalized_v3v3(n1, n2);
|
||||
angle_a = compare_v3v3(n1, n2, FLT_EPSILON) ? 0.0f : angle_normalized_v3v3(n1, n2);
|
||||
|
||||
normal_tri_v3(n1, v2, v3, v4);
|
||||
normal_tri_v3(n2, v4, v1, v2);
|
||||
angle_b = (compare_v3v3(n1, n2, FLT_EPSILON)) ? 0.0f : angle_normalized_v3v3(n1, n2);
|
||||
angle_b = compare_v3v3(n1, n2, FLT_EPSILON) ? 0.0f : angle_normalized_v3v3(n1, n2);
|
||||
|
||||
diff = (angle_a + angle_b) / (float)(M_PI * 2);
|
||||
|
||||
|
@ -166,15 +166,15 @@ static float bm_edge_is_delimit(const BMEdge *e, const struct DelimitData *delim
|
|||
float angle;
|
||||
#endif
|
||||
|
||||
if ((delimit_data->do_seam) && (BM_elem_flag_test(e, BM_ELEM_SEAM))) {
|
||||
if (delimit_data->do_seam && BM_elem_flag_test(e, BM_ELEM_SEAM)) {
|
||||
goto fail;
|
||||
}
|
||||
|
||||
if ((delimit_data->do_sharp) && (BM_elem_flag_test(e, BM_ELEM_SMOOTH) == 0)) {
|
||||
if (delimit_data->do_sharp && (BM_elem_flag_test(e, BM_ELEM_SMOOTH) == 0)) {
|
||||
goto fail;
|
||||
}
|
||||
|
||||
if ((delimit_data->do_mat) && (f_a->mat_nr != f_b->mat_nr)) {
|
||||
if (delimit_data->do_mat && (f_a->mat_nr != f_b->mat_nr)) {
|
||||
goto fail;
|
||||
}
|
||||
|
||||
|
|
|
@ -4697,7 +4697,7 @@ static VMesh *pipe_adj_vmesh(BevelParams *bp, BevVert *bv, BoundVert *vpipe)
|
|||
* vertices to snap to the midline on the pipe, not just to one plane or the other. */
|
||||
bool even = (ns % 2) == 0;
|
||||
bool midline = even && k == half_ns &&
|
||||
((i == 0 && j == half_ns) || (ELEM(i, ipipe1, ipipe2)));
|
||||
((i == 0 && j == half_ns) || ELEM(i, ipipe1, ipipe2));
|
||||
snap_to_pipe_profile(vpipe, midline, mesh_vert(vm, i, j, k)->co);
|
||||
}
|
||||
}
|
||||
|
@ -5376,7 +5376,7 @@ static void bevel_build_rings(BevelParams *bp, BMesh *bm, BevVert *bv, BoundVert
|
|||
for (int i = 0; i < n_bndv; i++) {
|
||||
for (int j = 0; j <= ns2; j++) {
|
||||
for (int k = 0; k <= ns; k++) {
|
||||
if (j == 0 && (ELEM(k, 0, ns))) {
|
||||
if (j == 0 && ELEM(k, 0, ns)) {
|
||||
continue; /* Boundary corners already made. */
|
||||
}
|
||||
if (!is_canon(vm, i, j, k)) {
|
||||
|
|
|
@ -463,7 +463,7 @@ void BM_mesh_bisect_plane(BMesh *bm,
|
|||
}
|
||||
|
||||
vert_is_center_disable(v);
|
||||
BM_VERT_DIR(v) = plane_point_test_v3(plane, v->co, eps, &(BM_VERT_DIST(v)));
|
||||
BM_VERT_DIR(v) = plane_point_test_v3(plane, v->co, eps, &BM_VERT_DIST(v));
|
||||
|
||||
if (BM_VERT_DIR(v) == 0) {
|
||||
if (oflag_center) {
|
||||
|
|
|
@ -618,9 +618,9 @@ static void bm_decim_triangulate_end(BMesh *bm, const int edges_tri_tot)
|
|||
(BM_loop_is_manifold(l) && ((l)->v != (l)->radial_next->v) && \
|
||||
(l_a_index == BM_elem_index_get(l)) && (l_a_index == BM_elem_index_get((l)->radial_next)))
|
||||
|
||||
if ((l_a->f->len == 3 && l_b->f->len == 3) && (!CAN_LOOP_MERGE(l_a->next)) &&
|
||||
(!CAN_LOOP_MERGE(l_a->prev)) && (!CAN_LOOP_MERGE(l_b->next)) &&
|
||||
(!CAN_LOOP_MERGE(l_b->prev))) {
|
||||
if ((l_a->f->len == 3 && l_b->f->len == 3) && !CAN_LOOP_MERGE(l_a->next) &&
|
||||
!CAN_LOOP_MERGE(l_a->prev) && !CAN_LOOP_MERGE(l_b->next) &&
|
||||
!CAN_LOOP_MERGE(l_b->prev)) {
|
||||
BMVert *vquad[4] = {
|
||||
e->v1,
|
||||
BM_vert_in_edge(e, l_a->next->v) ? l_a->prev->v : l_a->next->v,
|
||||
|
|
|
@ -37,7 +37,7 @@ enum {
|
|||
*/
|
||||
static bool bm_edge_step_ok(BMEdge *e)
|
||||
{
|
||||
return BM_elem_flag_test(e, BM_ELEM_TAG) && (ELEM(e->l, NULL, e->l->radial_next));
|
||||
return BM_elem_flag_test(e, BM_ELEM_TAG) && ELEM(e->l, NULL, e->l->radial_next);
|
||||
}
|
||||
|
||||
static int bm_edge_face(BMEdge *e)
|
||||
|
|
|
@ -1263,7 +1263,7 @@ bool BM_mesh_intersect(BMesh *bm,
|
|||
/* only start on an edge-case */
|
||||
/* pass */
|
||||
}
|
||||
else if ((!BM_elem_flag_test(v_a, BM_ELEM_TAG)) && (!BM_elem_flag_test(v_b, BM_ELEM_TAG))) {
|
||||
else if (!BM_elem_flag_test(v_a, BM_ELEM_TAG) && !BM_elem_flag_test(v_b, BM_ELEM_TAG)) {
|
||||
/* simple case, single edge spans face */
|
||||
BMVert **splice_pair;
|
||||
BM_elem_flag_enable(e_pair[1], BM_ELEM_TAG);
|
||||
|
|
|
@ -231,7 +231,7 @@ static void edgetag_add_adjacent(HeapSimple *heap,
|
|||
BM_ITER_ELEM (v, &viter, e_a, BM_VERTS_OF_EDGE) {
|
||||
|
||||
/* Don't walk over previous vertex. */
|
||||
if ((edges_prev[e_a_index]) && (BM_vert_in_edge(edges_prev[e_a_index], v))) {
|
||||
if ((edges_prev[e_a_index]) && BM_vert_in_edge(edges_prev[e_a_index], v)) {
|
||||
continue;
|
||||
}
|
||||
|
||||
|
|
|
@ -354,7 +354,7 @@ struct LinkNode *BM_mesh_calc_path_uv_edge(BMesh *bm,
|
|||
while (!BLI_heapsimple_is_empty(heap)) {
|
||||
l = BLI_heapsimple_pop_min(heap);
|
||||
|
||||
if ((l->e == l_dst->e) && (BM_loop_uv_share_edge_check(l, l_dst, params->cd_loop_uv_offset))) {
|
||||
if ((l->e == l_dst->e) && BM_loop_uv_share_edge_check(l, l_dst, params->cd_loop_uv_offset)) {
|
||||
break;
|
||||
}
|
||||
|
||||
|
@ -364,7 +364,7 @@ struct LinkNode *BM_mesh_calc_path_uv_edge(BMesh *bm,
|
|||
}
|
||||
}
|
||||
|
||||
if ((l->e == l_dst->e) && (BM_loop_uv_share_edge_check(l, l_dst, params->cd_loop_uv_offset))) {
|
||||
if ((l->e == l_dst->e) && BM_loop_uv_share_edge_check(l, l_dst, params->cd_loop_uv_offset)) {
|
||||
do {
|
||||
BLI_linklist_prepend(&path, l);
|
||||
} while ((l = loops_prev[BM_elem_index_get(l)]));
|
||||
|
|
|
@ -510,7 +510,7 @@ static void bm_uuidwalk_pass_add(UUIDWalk *uuidwalk,
|
|||
do {
|
||||
if (!BLI_ghash_haskey(uuidwalk->faces_uuid, l_iter_radial->f) &&
|
||||
!BLI_gset_haskey(faces_step_next, l_iter_radial->f) &&
|
||||
(bm_uuidwalk_face_test(uuidwalk, l_iter_radial->f))) {
|
||||
bm_uuidwalk_face_test(uuidwalk, l_iter_radial->f)) {
|
||||
BLI_gset_insert(faces_step_next, l_iter_radial->f);
|
||||
|
||||
/* add to fstep */
|
||||
|
|
|
@ -53,10 +53,10 @@ Node::Node(bNode *editor_node, bool create_sockets)
|
|||
Node::~Node()
|
||||
{
|
||||
while (!outputs_.is_empty()) {
|
||||
delete (outputs_.pop_last());
|
||||
delete outputs_.pop_last();
|
||||
}
|
||||
while (!inputs_.is_empty()) {
|
||||
delete (inputs_.pop_last());
|
||||
delete inputs_.pop_last();
|
||||
}
|
||||
}
|
||||
|
||||
|
|
|
@ -576,7 +576,7 @@ void NodeOperationBuilder::add_output_buffers(NodeOperation *operation,
|
|||
/* try to find existing write buffer operation */
|
||||
if (target->get_operation().get_flags().is_write_buffer_operation) {
|
||||
BLI_assert(write_operation == nullptr); /* there should only be one write op connected */
|
||||
write_operation = (WriteBufferOperation *)(&target->get_operation());
|
||||
write_operation = (WriteBufferOperation *)&target->get_operation();
|
||||
}
|
||||
else {
|
||||
/* remove all links to other nodes */
|
||||
|
|
|
@ -191,7 +191,7 @@ bool check_id_has_anim_component(ID *id)
|
|||
if (adt == nullptr) {
|
||||
return false;
|
||||
}
|
||||
return (adt->action != nullptr) || (!BLI_listbase_is_empty(&adt->nla_tracks));
|
||||
return (adt->action != nullptr) || !BLI_listbase_is_empty(&adt->nla_tracks);
|
||||
}
|
||||
|
||||
bool check_id_has_driver_component(ID *id)
|
||||
|
|
|
@ -490,7 +490,6 @@ set(GLSL_SRC
|
|||
intern/shaders/common_debug_shape_lib.glsl
|
||||
intern/shaders/common_fullscreen_vert.glsl
|
||||
intern/shaders/common_fxaa_lib.glsl
|
||||
intern/shaders/common_globals_lib.glsl
|
||||
intern/shaders/common_gpencil_lib.glsl
|
||||
intern/shaders/common_hair_lib.glsl
|
||||
intern/shaders/common_hair_refine_comp.glsl
|
||||
|
|
|
@ -85,8 +85,7 @@ static void basic_cache_init(void *vedata)
|
|||
|
||||
DRW_PASS_CREATE(psl->depth_pass[i], state | clip_state | infront_state);
|
||||
stl->g_data->depth_shgrp[i] = grp = DRW_shgroup_create(sh, psl->depth_pass[i]);
|
||||
DRW_shgroup_uniform_vec2(grp, "sizeViewport", DRW_viewport_size_get(), 1);
|
||||
DRW_shgroup_uniform_vec2(grp, "sizeViewportInv", DRW_viewport_invert_size_get(), 1);
|
||||
DRW_shgroup_uniform_block(grp, "globalsBlock", G_draw.block_ubo);
|
||||
|
||||
sh = DRW_state_is_select() ?
|
||||
BASIC_shaders_pointcloud_depth_conservative_sh_get(draw_ctx->sh_cfg) :
|
||||
|
@ -94,22 +93,22 @@ static void basic_cache_init(void *vedata)
|
|||
DRW_PASS_CREATE(psl->depth_pass_pointcloud[i], state | clip_state | infront_state);
|
||||
stl->g_data->depth_pointcloud_shgrp[i] = grp = DRW_shgroup_create(
|
||||
sh, psl->depth_pass_pointcloud[i]);
|
||||
DRW_shgroup_uniform_vec2(grp, "sizeViewport", DRW_viewport_size_get(), 1);
|
||||
DRW_shgroup_uniform_vec2(grp, "sizeViewportInv", DRW_viewport_invert_size_get(), 1);
|
||||
DRW_shgroup_uniform_block(grp, "globalsBlock", G_draw.block_ubo);
|
||||
|
||||
stl->g_data->depth_hair_shgrp[i] = grp = DRW_shgroup_create(
|
||||
BASIC_shaders_depth_sh_get(draw_ctx->sh_cfg), psl->depth_pass[i]);
|
||||
DRW_shgroup_uniform_block(grp, "globalsBlock", G_draw.block_ubo);
|
||||
|
||||
stl->g_data->depth_curves_shgrp[i] = grp = DRW_shgroup_create(
|
||||
BASIC_shaders_curves_depth_sh_get(draw_ctx->sh_cfg), psl->depth_pass[i]);
|
||||
DRW_shgroup_uniform_block(grp, "globalsBlock", G_draw.block_ubo);
|
||||
|
||||
sh = DRW_state_is_select() ? BASIC_shaders_depth_conservative_sh_get(draw_ctx->sh_cfg) :
|
||||
BASIC_shaders_depth_sh_get(draw_ctx->sh_cfg);
|
||||
state |= DRW_STATE_CULL_BACK;
|
||||
DRW_PASS_CREATE(psl->depth_pass_cull[i], state | clip_state | infront_state);
|
||||
stl->g_data->depth_shgrp_cull[i] = grp = DRW_shgroup_create(sh, psl->depth_pass_cull[i]);
|
||||
DRW_shgroup_uniform_vec2(grp, "sizeViewport", DRW_viewport_size_get(), 1);
|
||||
DRW_shgroup_uniform_vec2(grp, "sizeViewportInv", DRW_viewport_invert_size_get(), 1);
|
||||
DRW_shgroup_uniform_block(grp, "globalsBlock", G_draw.block_ubo);
|
||||
}
|
||||
}
|
||||
|
||||
|
|
|
@ -19,7 +19,7 @@ void main()
|
|||
/* Compute NDC bound box. */
|
||||
vec4 bbox = vec4(min(min(pos0.xy, pos1.xy), pos2.xy), max(max(pos0.xy, pos1.xy), pos2.xy));
|
||||
/* Convert to pixel space. */
|
||||
bbox = (bbox * 0.5 + 0.5) * drw_view.viewport_size.xyxy;
|
||||
bbox = (bbox * 0.5 + 0.5) * sizeViewport.xyxy;
|
||||
/* Detect failure cases where triangles would produce no fragments. */
|
||||
bvec2 is_subpixel = lessThan(bbox.zw - bbox.xy, vec2(1.0));
|
||||
/* View aligned triangle. */
|
||||
|
@ -31,13 +31,13 @@ void main()
|
|||
if (all(is_subpixel)) {
|
||||
vec2 ofs = (i == 0) ? vec2(-1.0) : ((i == 1) ? vec2(2.0, -1.0) : vec2(-1.0, 2.0));
|
||||
/* HACK: Fix cases where the triangle is too small make it cover at least one pixel. */
|
||||
gl_Position.xy += drw_view.viewport_size_inverse * gl_Position.w * ofs;
|
||||
gl_Position.xy += sizeViewportInv * gl_Position.w * ofs;
|
||||
}
|
||||
/* Test if the triangle is almost parallel with the view to avoid precision issues. */
|
||||
else if (any(is_subpixel) || is_coplanar) {
|
||||
/* HACK: Fix cases where the triangle is Parallel to the view by deforming it slightly. */
|
||||
vec2 ofs = (i == 0) ? vec2(-1.0) : ((i == 1) ? vec2(1.0, -1.0) : vec2(1.0));
|
||||
gl_Position.xy += drw_view.viewport_size_inverse * gl_Position.w * ofs;
|
||||
gl_Position.xy += sizeViewportInv * gl_Position.w * ofs;
|
||||
}
|
||||
else {
|
||||
/* Triangle expansion should happen here, but we decide to not implement it for
|
||||
|
|
|
@ -60,6 +60,6 @@ GPU_SHADER_CREATE_INFO(basic_curves)
|
|||
|
||||
GPU_SHADER_CREATE_INFO(basic_depth).fragment_source("basic_depth_frag.glsl");
|
||||
|
||||
BASIC_OBTYPE_VARIATIONS(basic_depth, "basic_depth");
|
||||
BASIC_OBTYPE_VARIATIONS(basic_depth, "basic_depth", "draw_globals");
|
||||
|
||||
/** \} */
|
||||
|
|
Some files were not shown because too many files have changed in this diff Show More
Loading…
Reference in New Issue