Compare commits

..
Author SHA1 Message Date
Opitz, Daniel 4d78c72a77 test interop in ci 2025-08-11 14:56:32 +02:00
Opitz, Daniel 442f84f3d6 fixed interop test vk errors 2025-07-31 14:36:59 +02:00
Daniel Opitz 977da20d32 added dx interop testing app 2025-07-31 14:08:04 +02:00
Opitz, Daniel 872ae0504e added vulkan interop testing app 2025-07-17 09:40:48 +02:00
191 changed files with 7532 additions and 9452 deletions
+1 -1
View File
@@ -6,6 +6,6 @@
version: 2
updates:
- package-ecosystem: "pip" # See documentation for possible values
directory: "/training" # Location of package manifests
directory: "/" # Location of package manifests
schedule:
interval: "monthly"
+16 -16
View File
@@ -31,9 +31,9 @@ jobs:
mkdir /tmp/deps
cd /tmp/deps
wget https://github.com/ispc/ispc/releases/download/v1.30.0/ispc-v1.30.0-linux.tar.gz
tar -xvf ispc-v1.30.0-linux.tar.gz
echo "PATH=$PATH:`pwd`/ispc-v1.30.0-linux/bin" >> $GITHUB_ENV
wget https://github.com/ispc/ispc/releases/download/v1.24.0/ispc-v1.24.0-linux.tar.gz
tar -xvf ispc-v1.24.0-linux.tar.gz
echo "PATH=$PATH:`pwd`/ispc-v1.24.0-linux/bin" >> $GITHUB_ENV
- name: Checkout Repository
uses: actions/checkout@v4
@@ -90,9 +90,9 @@ jobs:
mkdir /tmp/deps
cd /tmp/deps
wget https://github.com/ispc/ispc/releases/download/v1.30.0/ispc-v1.30.0-linux.tar.gz
tar -xvf ispc-v1.30.0-linux.tar.gz
echo "PATH=$PATH:`pwd`/ispc-v1.30.0-linux/bin" >> $GITHUB_ENV
wget https://github.com/ispc/ispc/releases/download/v1.24.0/ispc-v1.24.0-linux.tar.gz
tar -xvf ispc-v1.24.0-linux.tar.gz
echo "PATH=$PATH:`pwd`/ispc-v1.24.0-linux/bin" >> $GITHUB_ENV
- name: Checkout Repository
uses: actions/checkout@v4
@@ -149,9 +149,9 @@ jobs:
mkdir /tmp/deps
cd /tmp/deps
wget https://github.com/ispc/ispc/releases/download/v1.30.0/ispc-v1.30.0-linux.aarch64.tar.gz
tar -xvf ispc-v1.30.0-linux.aarch64.tar.gz
echo "PATH=$PATH:`pwd`/ispc-v1.30.0-linux.aarch64/bin" >> $GITHUB_ENV
wget https://github.com/ispc/ispc/releases/download/v1.24.0/ispc-v1.24.0-linux.aarch64.tar.gz
tar -xvf ispc-v1.24.0-linux.aarch64.tar.gz
echo "PATH=$PATH:`pwd`/ispc-v1.24.0-linux.aarch64/bin" >> $GITHUB_ENV
- name: Checkout Repository
uses: actions/checkout@v4
@@ -171,7 +171,7 @@ jobs:
with:
name: build-cpu-rocky-8-arm
path: build/install
test-cpu-rocky-8-arm:
needs: build-cpu-rocky-8-arm
runs-on: ubuntu-24.04-arm
@@ -191,7 +191,7 @@ jobs:
chmod +x ./bin/*
./bin/oidnTest
./bin/oidnBenchmark -v 1
build-cpu-ubuntu-2204-arm:
runs-on: ubuntu-24.04-arm
container:
@@ -208,9 +208,9 @@ jobs:
mkdir /tmp/deps
cd /tmp/deps
wget https://github.com/ispc/ispc/releases/download/v1.30.0/ispc-v1.30.0-linux.aarch64.tar.gz
tar -xvf ispc-v1.30.0-linux.aarch64.tar.gz
echo "PATH=$PATH:`pwd`/ispc-v1.30.0-linux.aarch64/bin" >> $GITHUB_ENV
wget https://github.com/ispc/ispc/releases/download/v1.24.0/ispc-v1.24.0-linux.aarch64.tar.gz
tar -xvf ispc-v1.24.0-linux.aarch64.tar.gz
echo "PATH=$PATH:`pwd`/ispc-v1.24.0-linux.aarch64/bin" >> $GITHUB_ENV
- name: Checkout Repository
uses: actions/checkout@v4
@@ -230,7 +230,7 @@ jobs:
with:
name: build-cpu-ubuntu-2204-arm
path: build/install
test-cpu-ubuntu-2204-arm:
needs: build-cpu-ubuntu-2204-arm
runs-on: ubuntu-24.04-arm
@@ -250,4 +250,4 @@ jobs:
chmod +x ./bin/*
./bin/oidnTest
./bin/oidnBenchmark -v 1
@@ -1 +0,0 @@
GFX_DRIVER_VERSION=windows-101.7080
@@ -0,0 +1 @@
GFX_DRIVER_VERSION=windows-101.6256
+1 -1
View File
@@ -1 +1 @@
GFX_DRIVER_VERSION=windows-latest
GFX_DRIVER_VERSION=windows-101.6449
File diff suppressed because it is too large Load Diff
+3
View File
@@ -1,3 +1,6 @@
[submodule "mkl-dnn"]
path = external/mkl-dnn
url = ../mkl-dnn.git
[submodule "weights"]
path = weights
url = ../oidn-weights.git
-34
View File
@@ -1,40 +1,6 @@
Version History
---------------
### Changes in v2.5.0:
- Significantly improved performance and reduced memory usage on Intel GPUs
with XMX support and CPUs with AMX-FP16 support
- Added API for importing external semaphores from graphics APIs (e.g. Vulkan,
Direct3D 12). Currently this is supported only by CUDA (Windows and Linux)
and HIP (Windows only) devices. SYCL device support will be added in a
future version
- Added the `OIDN_EXTERNAL_MEMORY_TYPE_FLAG_DEDICATED` flag which must be
combined with the handle type flag when importing external memory with
dedicated allocation
- Fixed corrupted output on Apple M5 Pro/Max GPUs using Metal
- Fixed a crash caused by over-releasing the MTLDevice, which could occur
after creating and destroying multiple devices
- Fixed device detection failure or crash on Windows if some old Intel
integrated GPU drivers are installed (fix requires building with oneAPI
DPC++ Compiler 6.1.0)
### Changes in v2.4.1:
- Added AMD RDNA 3.5 GFX1152 GPU support
### Changes in v2.4.0:
- Added Intel AMX-FP16 support, dramatically improving performance on Intel
Granite Rapids CPUs
- Added Intel BMG-G31, Wildcat Lake, Nova Lake, and Crescent Island GPU support
- Added AMD RDNA 3.5 GPU support, extended RDNA 2 support
- Fixed integer overflow and out-of-bounds write issues in image loaders (only
affects the `oidnDenoise` example application)
- Added support for compilation with CUDA 13
- Added support for compilation with ROCm 7
- Removed support for NVIDIA Volta GPUs
### Changes in v2.3.3:
- Added NVIDIA Blackwell GPU support
+1
View File
@@ -134,6 +134,7 @@ install(
${PROJECT_SOURCE_DIR}/CHANGELOG.md
${PROJECT_SOURCE_DIR}/third-party-programs.txt
${PROJECT_SOURCE_DIR}/third-party-programs-DPCPP.txt
${PROJECT_SOURCE_DIR}/third-party-programs-oneDNN.txt
${PROJECT_SOURCE_DIR}/third-party-programs-oneTBB.txt
DESTINATION ${CMAKE_INSTALL_DOCDIR}
COMPONENT lib
+113 -186
View File
@@ -1,6 +1,6 @@
# Intel® Open Image Denoise
This is release v2.5.0 of Intel Open Image Denoise. For changes and new
This is release v2.3.3 of Intel Open Image Denoise. For changes and new
features see the [changelog](CHANGELOG.md). Visit
https://www.openimagedenoise.org for more information.
@@ -58,12 +58,13 @@ different vendors:
Graphics, Intel® Core™ Ultra Processors with Intel® Arc™ Graphics,
11th-14th Gen Intel® Core™ processor graphics, and related Intel
Pentium® and Celeron® processors (Xe-LP, Xe-LPG, Xe-LPG+, Xe-HPG,
Xe-HPC, Xe2-LPG, Xe2-HPG, Xe3-LPG, and Xe3p-XPC microarchitectures)
Xe-HPC, Xe2-LPG, Xe2-HPG, and Xe3-LPG microarchitectures)
- NVIDIA GPUs with Turing, Ampere, Ada Lovelace, Hopper, and Blackwell
architectures
- NVIDIA GPUs with Volta, Turing, Ampere, Ada Lovelace, Hopper, and
Blackwell architectures
- AMD GPUs with RDNA 2, RDNA 3, RDNA 3.5, and RDNA 4 architectures
- AMD GPUs with RDNA2 (Navi 21 only), RDNA3 (Navi 3x), and RDNA4 (Navi
4x) architectures
- Apple silicon GPUs (M1 and newer)
@@ -73,10 +74,9 @@ only for offline rendering, but, depending on the hardware used, also
for interactive or even real-time ray tracing.
Intel Open Image Denoise exploits modern instruction sets like SSE4,
AVX2, AVX-512, Intel® Advanced Matrix Extensions (Intel® AMX), and NEON
on CPUs, Intel® Xe Matrix Extensions (Intel® XMX) on Intel GPUs, and
various other AI acceleration capabilities on NVIDIA, AMD, and Apple
GPUs.
AVX2, AVX-512, and NEON on CPUs, Intel® Xe Matrix Extensions (Intel®
XMX) on Intel GPUs, and tensor cores on NVIDIA GPUs to achieve high
denoising performance.
## System Requirements
@@ -88,12 +88,14 @@ For Intel GPU support, please also install the latest Intel graphics
drivers:
- Windows: [Intel® Graphics
Driver](https://www.intel.com/content/www/us/en/download-center/home.html)
Driver](https://www.intel.com/content/www/us/en/download/726609/intel-arc-iris-xe-graphics-whql-windows.html)
31.0.101.4953 or newer
- Linux: [Intel® software for General Purpose GPU
capabilities](https://dgpu-docs.intel.com/driver/overview/overview.html)
release 20230323 or newer
capabilities](https://dgpu-docs.intel.com/driver/installation.html)
release
[20230323](https://dgpu-docs.intel.com/releases/stable_602_20230323.html)
or newer
Using older driver versions is *not* supported and Intel Open Image
Denoise might run with only limited capabilities, have suboptimal
@@ -111,11 +113,11 @@ drivers](https://www.nvidia.com/en-us/geforce/drivers/):
For AMD GPU support, please also install the latest [AMD graphics
drivers](https://www.amd.com/en/support):
- Windows: AMD Software: Adrenalin Edition 25.30.02.01 or newer
- Windows: AMD Software: Adrenalin Edition 25.3.1 or newer
- Linux: [Radeon Software for
Linux](https://www.amd.com/en/support/download/linux-drivers.html)
version 25.30.1 or newer
Linux](https://www.amd.com/en/support/linux-drivers) version 24.30.4
or newer
For Apple GPU support, macOS Ventura or newer is required.
@@ -144,7 +146,7 @@ cite the project using the following BibTeX entry:
@misc{OpenImageDenoise,
author = {Attila T. {\'A}fra},
title = {{Intel\textsuperscript{\textregistered} Open Image Denoise}},
year = {2026},
year = {2025},
note = {\url{https://www.openimagedenoise.org}}
}
```
@@ -185,18 +187,18 @@ additional prerequisites are needed:
#### CPU device:
- [Intel® SPMD Program Compiler (ISPC)](http://ispc.github.io) 1.30.0
- [Intel® SPMD Program Compiler (ISPC)](http://ispc.github.io) 1.21.0
or newer. Please obtain a release of ISPC from the [ISPC downloads
page](https://ispc.github.io/downloads.html). The build system looks
for ISPC in the `PATH` and in the directory right “next to” the
checked-out Intel Open Image Denoise sources. For example, if Intel
Open Image Denoise is in `~/Projects/oidn`, ISPC will also be
searched in `~/Projects/ispc-v1.30.0-linux`. Alternatively set the
searched in `~/Projects/ispc-v1.21.0-linux`. Alternatively set the
CMake variable `ISPC_EXECUTABLE` to the location of the ISPC
compiler.
- [Intel® Threading Building
Blocks](https://github.com/uxlfoundation/oneTBB) (TBB) 2017 or newer
Blocks](https://github.com/oneapi-src/oneTBB) (TBB) 2017 or newer
#### SYCL device for Intel GPUs:
@@ -205,19 +207,19 @@ additional prerequisites are needed:
Image Denoise):
- [oneAPI DPC++
Compiler 6.1.0](https://github.com/intel/llvm/releases/tag/v6.1.0).
This is the open source version of the compiler. Versions 6.2.x
and 6.3.0 may cause crashes or device detection failures on
Windows so it is highly recommended to avoid these.
Compiler 6.0.1](https://github.com/intel/llvm/releases/tag/v6.0.1).
This is the open source version of the compiler, which needs to
be built from source. We recommend building it with the
`--disable-jit` flag to minimize the size of the SYCL runtime
binaries.
- [Intel® oneAPI DPC++/C++
Compiler](https://www.intel.com/content/www/us/en/developer/tools/oneapi/dpc-compiler.html)
2025.3 or newer
2024.1 or newer
- *Optional*: Intel® Graphics Offline Compiler for OpenCL™ Code
(OCLOC), if building with `OIDN_DEVICE_SYCL_AOT` enabled
- Intel® Graphics Offline Compiler for OpenCL™ Code (OCLOC)
- Windows: Version [2025.3.3
/ 32.0.101.8331](https://registrationcenter-download.intel.com/akdlm/IRC_NAS/cb17f6e4-6e61-47c7-bb27-1008b23f1c7b/intel-ocloc-2025.3.3.4_offline.exe)
- Windows: Version [2025.0.0
/ 32.0.101.6129](https://registrationcenter-download.intel.com/akdlm/IRC_NAS/7000f8d2-dda8-4dd6-8b63-3917e4476fa5/intel-ocloc-2025.0.0.257_offline.exe)
or newer as a [standalone component of Intel® oneAPI
Toolkits](https://www.intel.com/content/www/us/en/developer/articles/tool/oneapi-standalone-components.html),
which must be extracted and its contents added to the `PATH`.
@@ -226,11 +228,11 @@ additional prerequisites are needed:
- Linux: Included with [Intel® software for General Purpose GPU
capabilities](https://dgpu-docs.intel.com) release
[LTS 2523.x](https://dgpu-docs.intel.com/releases/LTS-release-notes.html#release-2025-12-11)
[2441.19](https://dgpu-docs.intel.com/releases/rolling-release-notes.html#release-2024-10-31)
or newer (install at least `intel-opencl-icd` on Ubuntu,
`intel-ocloc` on RHEL or SLES). For more recent versions please
refer to [Intel® Graphics Compute Runtime for oneAPI Level Zero
and OpenCL™ Driver](https://github.com/intel/compute-runtime).
`intel-ocloc` on RHEL or SLES). Also available with [Intel®
Graphics Compute Runtime for oneAPI Level Zero and OpenCL™
Driver](https://github.com/intel/compute-runtime).
- If using Intel® oneAPI DPC++/C++ Compiler:
[CMake](http://www.cmake.org) 3.25.2 or newer
@@ -252,7 +254,9 @@ additional prerequisites are needed:
- [Ninja](https://ninja-build.org) or Make as the CMake generator. The
Visual Studio generator is *not* supported.
- [AMD ROCm (HIP SDK)](https://rocm.docs.amd.com) v6.4.2 or newer.
- [AMD ROCm (HIP SDK)](https://rocm.docs.amd.com) v6.2.4 or newer.
- Perl (e.g. [Strawberry Perl](https://strawberryperl.com) on Windows)
#### Metal device for Apple GPUs:
@@ -272,8 +276,12 @@ follow the included instructions.
If you are building with SYCL support on Linux, make sure that the DPC++
compiler is properly set up. The open source oneAPI DPC++ Compiler can
be downloaded and simply extracted. However, before using the compiler,
the environment must be set up as described in the [Get Started
Guide](https://github.com/intel/llvm/blob/sycl/sycl/doc/GetStartedGuide.md).
the environment must be set up as well with the following command:
source ./dpcpp_compiler/startup.sh
The `startup.sh` script will put `clang` and `clang++` from the oneAPI
DPC++ Compiler into your `PATH`.
Alternatively, if you have installed Intel® oneAPI DPC++/C++ Compiler
instead, you can set up the compiler by sourcing the `vars.sh` script in
@@ -319,8 +327,19 @@ Intel(R) oneAPI DPC++/C++ Compiler in your `PATH`.
If you are building with SYCL support, make sure that the DPC++ compiler
is properly set up. The open source oneAPI DPC++ Compiler can be
downloaded and simply extracted. However, before using the compiler, the
environment must be set up as described in the [Get Started
Guide](https://github.com/intel/llvm/blob/sycl/sycl/doc/GetStartedGuide.md).
environment must be set up. To achieve this, open the “x64 Native Tools
Command Prompt for VS” that ships with Visual Studio and execute the
following commands:
set "DPCPP_DIR=path_to_dpcpp_compiler"
set "PATH=%DPCPP_DIR%\bin;%PATH%"
set "PATH=%DPCPP_DIR%\lib;%PATH%"
set "CPATH=%DPCPP_DIR%\include;%CPATH%"
set "INCLUDE=%DPCPP_DIR%\include;%INCLUDE%"
set "LIB=%DPCPP_DIR%\lib;%LIB%"
The `path_to_dpcpp_compiler` should point to the unpacked oneAPI DPC++
Compiler.
Alternatively, if you have installed Intel® oneAPI DPC++/C++ Compiler
instead, you can either open a regular “Command Prompt” and execute the
@@ -399,9 +418,10 @@ CMake:
by default).
- `OIDN_DEVICE_SYCL_AOT`: Enable ahead-of-time (AOT) compilation for
SYCL kernels (OFF by default). Turning this on adds dependency on
OCLOC at build time and increases binary size but decreases
initialization time at runtime.
SYCL kernels (ON by default). Turning this off removes dependency on
OCLOC at build time and decreases binary size but significantly
increases initialization time at runtime, so it is recommended only
for development.
- `OIDN_DEVICE_CUDA`: Enable CUDA device support for NVIDIA GPUs (OFF
by default).
@@ -428,9 +448,9 @@ CMake:
default).
- `OIDN_APPS_OPENIMAGEIO`: Enable
[OpenImageIO](http://openimageio.org/) 2.1 or later support in the
example and test applications to be able to load/save OpenEXR, PNG,
and other image file formats (OFF by default).
[OpenImageIO](http://openimageio.org/) support in the example and
test applications to be able to load/save OpenEXR, PNG, and other
image file formats (OFF by default).
- `OIDN_INSTALL_DEPENDENCIES`: Enable installing the dependencies
(e.g. TBB, SYCL runtime) as well.
@@ -548,7 +568,7 @@ oidn::BufferRef colorBuf = device.newBuffer(width * height * 3 * sizeof(float))
oidn::BufferRef albedoBuf = ...
// Create a filter for denoising a beauty (color) image using optional auxiliary images too
// This can be an expensive operation, so try not to create a new filter for every image!
// This can be an expensive operation, so try no to create a new filter for every image!
oidn::FilterRef filter = device.newFilter("RT"); // generic ray tracing filter
filter.setImage("color", colorBuf, oidn::Format::Float3, width, height); // beauty
filter.setImage("albedo", albedoBuf, oidn::Format::Float3, width, height); // auxiliary
@@ -644,7 +664,7 @@ overhead of copying as much as possible:
buffers should be directly passed to filters as image parameters
instead of the original pointers using `oidnSetFilterImage`.
- Data should be copied asynchronously using the new
- Data should be copied asynchronously using using the new
`oidnReadBufferAsync` and `oidnWriteBufferAsync` functions, which
may achieve higher performance than plain `memcpy`.
@@ -674,14 +694,14 @@ compute API (e.g. `sycl::malloc_device`, `cudaMalloc`) instead of using
buffers. This way, it is the responsibility of the user to correctly
allocate the memory for the device.
In such cases, it is often necessary to have more control over the
device creation as well, to ensure that filtering is running on the
intended device and command queues or streams from the application can
be shared to improve performance. If the application is using the same
compute or graphics API as the Open Image Denoise device, this can be
achieved by creating devices with `oidnNewSYCLDevice`,
`oidnNewCUDADevice`, etc. For some APIs there are additional
interoperability functions as well, e.g. `oidnExecuteSYCLFilterAsync`.
In such cases, it often necessary to have more control over the device
creation as well, to ensure that filtering is running on the intended
device and command queues or streams from the application can be shared
to improve performance. If the application is using the same compute or
graphics API as the Open Image Denoise device, this can be achieved by
creating devices with `oidnNewSYCLDevice`, `oidnNewCUDADevice`, etc. For
some APIs there are additional interoperability functions as well,
e.g. `oidnExecuteSYCLFilterAsync`.
If the application is using a graphics API which does not support
unified memory allocations, e.g. DX12 or Vulkan, it may be still
@@ -707,21 +727,10 @@ possible avoid copying texture data by using a linear texture layout
backs this data. In this case, you should ensure that the row stride of
the linear texture data is correctly set.
For efficiently synchronizing access to shared memory, external
semaphores should be also imported from the graphics API using the
`oidnNewSharedSemaphoreFromFD` and
`oidnNewSharedSemaphoreFromWin32Handle` functions. The external memory
types supported by Open Image Denoise can be queried using the
`externalSemaphoreTypes` device parameter. Signaling and waiting on
semaphores are executed asynchronously with the
`oidnSignalSemaphoresAsync` and `oidnWaitSemaphoresAsync` functions.
If importing external semaphores is not possible, synchronizing access
to shared memory should be done on the host using `oidnSyncDevice` and
the graphics API but this typically introduces a performance overhead.
Since external semaphore support is not guaranteed on all device types,
driver versions and OSs, it is strongly recommended to implement
host-blocking synchronization as well in the application, as a fallback.
Importing external synchronization primitives (e.g. semaphores) from
graphics APIs is not yet supported either but it is planned for a future
release. Meanwhile, synchronizing access to shared memory should be done
on the host using `oidnSyncDevice` and the used graphics API.
When importing external memory, the application also needs to make sure
that the Open Image Denoise device is running on the same *physical*
@@ -740,8 +749,8 @@ Open Image Denoise 2 introduces a simple *physical device* API, which
enables the application to query the list of supported physical devices
in the system, including their name, type, UUID, LUID, PCI address, etc.
(see `oidnGetNumPhysicalDevices`, `oidnGetPhysicalDeviceString`, etc.).
New logical device (i.e. `OIDNDevice`) creation functions have also been
introduced, which enable creating a logical device on a specific
New logical device (i.e. `OIDNDevice`) creation functions for have been
also introduced, which enable creating a logical device on a specific
physical device: `oidnNewDeviceByID`, `oidnNewDeviceByUUID`, etc.
Creating a logical device on a physical device having a particular UUID,
@@ -795,11 +804,11 @@ multi-vendor GPU support. If the library is built with GPU support as
well, the `OIDN_STATIC_LIB` option is still available but enabling it
results in a hybrid static/shared library.
If the main reason for building as a static library is the ability to
use multiple versions of Open Image Denoise in the same process, please
use the existing `OIDN_API_NAMESPACE` CMake option instead. With this
feature all symbols of the library will be put into a custom namespace,
which can prevent symbol clashes.
If the main reason for building as a static library would be is the
ability to use multiple versions of Open Image Denoise in the same
process, please use the existing `OIDN_API_NAMESPACE` CMake option
instead. With this feature all symbols of the library will be put into a
custom namespace, which can prevent symbol clashes.
## Physical Devices
@@ -953,12 +962,12 @@ For Metal, a single command queue is supported.
Once a device is created, you can call
``` cpp
bool oidnGetDeviceBool(OIDNDevice device, const char* name);
void oidnSetDeviceBool(OIDNDevice device, const char* name, bool value);
int oidnGetDeviceInt (OIDNDevice device, const char* name);
void oidnSetDeviceInt (OIDNDevice device, const char* name, int value);
unsigned int oidnGetDeviceUInt(OIDNDevice device, const char* name);
void oidnSetDeviceUInt(OIDNDevice device, const char* name, unsigned int value);
bool oidnGetDeviceBool(OIDNDevice device, const char* name);
void oidnSetDeviceBool(OIDNDevice device, const char* name, bool value);
int oidnGetDeviceInt (OIDNDevice device, const char* name);
void oidnSetDeviceInt (OIDNDevice device, const char* name, int value);
int oidnGetDeviceUInt(OIDNDevice device, const char* name);
void oidnSetDeviceUInt(OIDNDevice device, const char* name, unsigned int value);
```
to set and get parameter values on the device. Note that some parameters
@@ -974,8 +983,7 @@ for the parameters supported by devices.
| `Int` | `versionPatch` | *constant* | patch version number |
| `Bool` | `systemMemorySupported` | *constant* | device can directly access memory allocated with the system allocator (e.g. `malloc`) |
| `Bool` | `managedMemorySupported` | *constant* | device supports buffers created with managed storage (`OIDN_STORAGE_MANAGED`) |
| `UInt` | `externalMemoryTypes` | *constant* | bitfield of `OIDNExternalMemoryTypeFlag` values representing the external memory types supported by the device |
| `UInt` | `externalSemaphoreTypes` | *constant* | bitfield of `OIDNExternalSemaphoreTypeFlag` values representing the external semaphore types supported by the device |
| `Int` | `externalMemoryTypes` | *constant* | bitfield of `OIDNExternalMemoryTypeFlag` values representing the external memory types supported by the device |
| `Int` | `verbose` | 0 | verbosity level of the console output between 0–4; when set to 0, no output is printed, when set to a higher level more output is printed |
Parameters supported by all devices.
@@ -1199,17 +1207,17 @@ Win32 handles using
``` cpp
OIDNBuffer oidnNewSharedBufferFromFD(OIDNDevice device,
OIDNExternalMemoryTypeFlags fdType,
OIDNExternalMemoryTypeFlag fdType,
int fd, size_t byteSize);
OIDNBuffer oidnNewSharedBufferFromWin32Handle(OIDNDevice device,
OIDNExternalMemoryTypeFlags handleType,
OIDNExternalMemoryTypeFlag handleType,
void* handle, const void* name, size_t byteSize);
```
Before exporting memory from the graphics API, the application should
find a handle type which is supported by both the Open Image Denoise
device (see the `externalMemoryTypes` device parameter) and the graphics
device (see `externalMemoryTypes` device parameter) and the graphics
API. Note that different GPU vendors may support different handle types.
To ensure compatibility with all device types, applications should
support at least `OIDN_EXTERNAL_MEMORY_TYPE_FLAG_OPAQUE_WIN32` on
@@ -1217,32 +1225,23 @@ Windows and both `OIDN_EXTERNAL_MEMORY_TYPE_FLAG_OPAQUE_FD` and
`OIDN_EXTERNAL_MEMORY_TYPE_FLAG_DMA_BUF` on Linux. All possible external
memory types are listed in the table below.
| Name | Description |
| :-------------------------------------------------- | :-------------------------------------------------------------------------------------------------------------------------------------------- |
| `OIDN_EXTERNAL_MEMORY_TYPE_FLAG_NONE` | |
| `OIDN_EXTERNAL_MEMORY_TYPE_FLAG_OPAQUE_FD` | opaque POSIX file descriptor handle (recommended on Linux) |
| `OIDN_EXTERNAL_MEMORY_TYPE_FLAG_DMA_BUF` | file descriptor handle for a Linux dma\_buf (recommended on Linux) |
| `OIDN_EXTERNAL_MEMORY_TYPE_FLAG_OPAQUE_WIN32` | NT handle (recommended on Windows) |
| `OIDN_EXTERNAL_MEMORY_TYPE_FLAG_OPAQUE_WIN32_KMT` | global share (KMT) handle |
| `OIDN_EXTERNAL_MEMORY_TYPE_FLAG_D3D11_TEXTURE` | NT handle returned by `IDXGIResource1::CreateSharedHandle` referring to a Direct3D 11 texture resource |
| `OIDN_EXTERNAL_MEMORY_TYPE_FLAG_D3D11_TEXTURE_KMT` | global share (KMT) handle returned by `IDXGIResource::GetSharedHandle` referring to a Direct3D 11 texture resource |
| `OIDN_EXTERNAL_MEMORY_TYPE_FLAG_D3D11_RESOURCE` | NT handle returned by `IDXGIResource1::CreateSharedHandle` referring to a Direct3D 11 resource |
| `OIDN_EXTERNAL_MEMORY_TYPE_FLAG_D3D11_RESOURCE_KMT` | global share (KMT) handle returned by `IDXGIResource::GetSharedHandle` referring to a Direct3D 11 resource |
| `OIDN_EXTERNAL_MEMORY_TYPE_FLAG_D3D12_HEAP` | NT handle returned by `ID3D12Device::CreateSharedHandle` referring to a Direct3D 12 heap resource |
| `OIDN_EXTERNAL_MEMORY_TYPE_FLAG_D3D12_RESOURCE` | NT handle returned by `ID3D12Device::CreateSharedHandle` referring to a Direct3D 12 committed resource |
| `OIDN_EXTERNAL_MEMORY_TYPE_FLAG_DEDICATED` | modifier flag indicating that the external memory has dedicated allocation (used only in combination with one of the handle type flags above) |
| Name | Description |
| :-------------------------------------------------- | :----------------------------------------------------------------------------------------------------------------- |
| `OIDN_EXTERNAL_MEMORY_TYPE_FLAG_NONE` | |
| `OIDN_EXTERNAL_MEMORY_TYPE_FLAG_OPAQUE_FD` | opaque POSIX file descriptor handle (recommended on Linux) |
| `OIDN_EXTERNAL_MEMORY_TYPE_FLAG_DMA_BUF` | file descriptor handle for a Linux dma\_buf (recommended on Linux) |
| `OIDN_EXTERNAL_MEMORY_TYPE_FLAG_OPAQUE_WIN32` | NT handle (recommended on Windows) |
| `OIDN_EXTERNAL_MEMORY_TYPE_FLAG_OPAQUE_WIN32_KMT` | global share (KMT) handle |
| `OIDN_EXTERNAL_MEMORY_TYPE_FLAG_D3D11_TEXTURE` | NT handle returned by `IDXGIResource1::CreateSharedHandle` referring to a Direct3D 11 texture resource |
| `OIDN_EXTERNAL_MEMORY_TYPE_FLAG_D3D11_TEXTURE_KMT` | global share (KMT) handle returned by `IDXGIResource::GetSharedHandle` referring to a Direct3D 11 texture resource |
| `OIDN_EXTERNAL_MEMORY_TYPE_FLAG_D3D11_RESOURCE` | NT handle returned by `IDXGIResource1::CreateSharedHandle` referring to a Direct3D 11 resource |
| `OIDN_EXTERNAL_MEMORY_TYPE_FLAG_D3D11_RESOURCE_KMT` | global share (KMT) handle returned by `IDXGIResource::GetSharedHandle` referring to a Direct3D 11 resource |
| `OIDN_EXTERNAL_MEMORY_TYPE_FLAG_D3D12_HEAP` | NT handle returned by `ID3D12Device::CreateSharedHandle` referring to a Direct3D 12 heap resource |
| `OIDN_EXTERNAL_MEMORY_TYPE_FLAG_D3D12_RESOURCE` | NT handle returned by `ID3D12Device::CreateSharedHandle` referring to a Direct3D 12 committed resource |
Supported external memory type flags, i.e., valid constants of type
`OIDNExternalMemoryTypeFlag`.
Please note that if the external memory uses dedicated allocation, the
`OIDN_EXTERNAL_MEMORY_TYPE_FLAG_DEDICATED` flag must be combined with
the handle type flag (e.g., `OIDN_EXTERNAL_MEMORY_TYPE_FLAG_OPAQUE_WIN32
| OIDN_EXTERNAL_MEMORY_TYPE_FLAG_DEDICATED` as `handleType`). To
maximize compatibility, we recommend to always use dedicated allocations
if possible because some backends support only dedicated allocations for
certain external memory types.
Metal buffers can be imported directly with
``` cpp
@@ -1327,78 +1326,6 @@ pointers. This can be done using the `OIDNFormat` enumeration type:
Supported data formats, i.e., valid constants of type `OIDNFormat`.
## Semaphores
Access to external buffers imported from graphics APIs needs to be
explicitly synchronized between Open Image Denoise and graphics
operations to avoid race conditions. The most efficient way is by using
external semaphores also imported from the graphics API, which enable
synchronization on the GPU device instead of the host (but this depends
on the implementation).
External semaphores can be imported similarly to buffers:
``` cpp
OIDNSemaphore oidnNewSharedSemaphoreFromFD(OIDNDevice device,
OIDNExternalSemaphoreTypeFlags fdType,
int fd);
OIDNSemaphore oidnNewSharedSemaphoreFromWin32Handle(OIDNDevice device,
OIDNExternalSemaphoreTypeFlags handleType,
void* handle, const void* name);
```
The semaphore handle types supported by Open Image Denoise can be
queried using the `externalSemaphoreTypes` device parameter. The
possible semaphore type flags are listed in the following table.
| Name | Description |
| :----------------------------------------------------------- | :--------------------------------------------------------------------- |
| `OIDN_EXTERNAL_SEMAPHORE_TYPE_FLAG_NONE` | |
| `OIDN_EXTERNAL_SEMAPHORE_TYPE_FLAG_OPAQUE_FD` | opaque POSIX file descriptor handle |
| `OIDN_EXTERNAL_SEMAPHORE_TYPE_FLAG_OPAQUE_WIN32` | opaque NT handle |
| `OIDN_EXTERNAL_SEMAPHORE_TYPE_FLAG_OPAQUE_WIN32_KMT` | opaque global share (KMT) handle |
| `OIDN_EXTERNAL_SEMAPHORE_TYPE_FLAG_D3D11_FENCE` | NT handle referencing a Direct3D 11 fence object |
| `OIDN_EXTERNAL_SEMAPHORE_TYPE_FLAG_D3D12_FENCE` | NT handle referencing a Direct3D 12 fence object |
| `OIDN_EXTERNAL_SEMAPHORE_TYPE_FLAG_KEYED_MUTEX` | NT handle referencing a Direct3D 11 keyed mutex object |
| `OIDN_EXTERNAL_SEMAPHORE_TYPE_FLAG_KEYED_MUTEX_KMT` | global share (KMT) handle referencing a Direct3D 11 keyed mutex object |
| `OIDN_EXTERNAL_SEMAPHORE_TYPE_FLAG_TIMELINE_SEMAPHORE_FD` | POSIX file descriptor referencing a timeline semaphore |
| `OIDN_EXTERNAL_SEMAPHORE_TYPE_FLAG_TIMELINE_SEMAPHORE_WIN32` | NT handle referencing a timeline semaphore |
Supported external semaphore type flags, i.e., valid constants of type
`OIDNExternalSemaphoreTypeFlag`.
The reference counted semaphore objects can be retained and released
with
``` cpp
void oidnRetainSemaphore(OIDNSemaphore semaphore);
void oidnReleaseSemaphore(OIDNSemaphore semaphore);
```
One or more semaphores can be signaled or waited on asynchronously by
calling the following functions:
``` cpp
void oidnSignalSemaphoresAsync(OIDNDevice device,
const OIDNSemaphore* semaphores,
const uint64_t* values,
int numSemaphores);
void oidnWaitSemaphoresAsync(OIDNDevice device,
const OIDNSemaphore* semaphores,
const uint64_t* values,
const uint32_t* timeoutsMs,
int numSemaphores);
```
The `values` and `timeoutMs` parameters are needed only for certain
kinds of semaphores (e.g. fences, keyed mutexes), for other semaphores
these values are ignored and the pointers may be `NULL`. The meaning of
the `values` parameter depends on the respective semaphore types
(e.g. for fences it is the value to set/wait for, for keyed mutexes it
is the key).
## Filters
Filters are the main objects in Open Image Denoise that are responsible
@@ -1415,10 +1342,10 @@ filter types are documented later in this section.
Creating filter objects can be very expensive, therefore it is
*strongly* recommended to reuse the same filter for denoising as many
images as possible, as long as these images have the same size, format,
and features (i.e., only the memory locations and pixel values may be
different). Otherwise (e.g. for images with different resolutions),
reusing the same filter would not have any benefits.
images as possible, as long as the these images have the same same size,
format, and features (i.e., only the memory locations and pixel values
may be different). Otherwise (e.g. for images with different
resolutions), reusing the same filter would not have any benefits.
Once created, filter objects can be retained and released with
@@ -2055,7 +1982,7 @@ used for inference will be either originally noise-free or prefiltered
with separately trained auxiliary feature denoising models. If inference
will be done only with prefiltered features, even higher quality can be
achieved by training with prefiltered features instead of the reference
ones. This can be achieved by first training the auxiliary feature
onces. This can be achieved by first training the auxiliary feature
models and then specifying the list of these results when preprocessing
the dataset for the main feature (`--aux_results` or `-a` option).
@@ -2143,7 +2070,7 @@ by running the `visualize.py` script, e.g.:
./visualize.py --result rt_hdr_alb
Training is performed with mixed precision (FP16 and FP32) by default,
if it is supported by the hardware, which makes training faster and use
if it supported by the hardware, which makes training faster and use
less memory. However, in some rare cases this might cause some
convergence issues. The training precision can be manually set to FP32
if necessary (`--precision` or `-p` option).
+139 -245
View File
@@ -141,10 +141,6 @@ OIDN_API_NAMESPACE_BEGIN
}
}
// -----------------------------------------------------------------------------------------------
// Physical Device
// -----------------------------------------------------------------------------------------------
OIDN_API int oidnGetNumPhysicalDevices()
{
OIDN_TRY
@@ -197,10 +193,6 @@ OIDN_API_NAMESPACE_BEGIN
return nullptr;
}
// -----------------------------------------------------------------------------------------------
// Device
// -----------------------------------------------------------------------------------------------
OIDN_API bool oidnIsCPUDeviceSupported()
{
OIDN_TRY
@@ -463,45 +455,45 @@ OIDN_API_NAMESPACE_BEGIN
return reinterpret_cast<OIDNDevice>(device.detach());
}
OIDN_API void oidnRetainDevice(OIDNDevice deviceC)
OIDN_API void oidnRetainDevice(OIDNDevice hDevice)
{
Device* device = reinterpret_cast<Device*>(deviceC);
Device* device = reinterpret_cast<Device*>(hDevice);
retainObject(device);
}
OIDN_API void oidnReleaseDevice(OIDNDevice deviceC)
OIDN_API void oidnReleaseDevice(OIDNDevice hDevice)
{
Device* device = reinterpret_cast<Device*>(deviceC);
Device* device = reinterpret_cast<Device*>(hDevice);
releaseObject(device);
}
OIDN_API void oidnSetDeviceBool(OIDNDevice deviceC, const char* name, bool value)
OIDN_API void oidnSetDeviceBool(OIDNDevice hDevice, const char* name, bool value)
{
Device* device = reinterpret_cast<Device*>(deviceC);
Device* device = reinterpret_cast<Device*>(hDevice);
OIDN_TRY
checkHandle(deviceC);
checkHandle(hDevice);
OIDN_LOCK_DEVICE(device);
checkString(name);
device->setInt(name, value);
OIDN_CATCH_DEVICE(device)
}
OIDN_API void oidnSetDeviceInt(OIDNDevice deviceC, const char* name, int value)
OIDN_API void oidnSetDeviceInt(OIDNDevice hDevice, const char* name, int value)
{
Device* device = reinterpret_cast<Device*>(deviceC);
Device* device = reinterpret_cast<Device*>(hDevice);
OIDN_TRY
checkHandle(deviceC);
checkHandle(hDevice);
OIDN_LOCK_DEVICE(device);
checkString(name);
device->setInt(name, value);
OIDN_CATCH_DEVICE(device)
}
OIDN_API bool oidnGetDeviceBool(OIDNDevice deviceC, const char* name)
OIDN_API bool oidnGetDeviceBool(OIDNDevice hDevice, const char* name)
{
Device* device = reinterpret_cast<Device*>(deviceC);
Device* device = reinterpret_cast<Device*>(hDevice);
OIDN_TRY
checkHandle(deviceC);
checkHandle(hDevice);
OIDN_LOCK_DEVICE(device);
checkString(name);
return device->getInt(name);
@@ -509,11 +501,11 @@ OIDN_API_NAMESPACE_BEGIN
return false;
}
OIDN_API int oidnGetDeviceInt(OIDNDevice deviceC, const char* name)
OIDN_API int oidnGetDeviceInt(OIDNDevice hDevice, const char* name)
{
Device* device = reinterpret_cast<Device*>(deviceC);
Device* device = reinterpret_cast<Device*>(hDevice);
OIDN_TRY
checkHandle(deviceC);
checkHandle(hDevice);
OIDN_LOCK_DEVICE(device);
checkString(name);
return device->getInt(name);
@@ -521,19 +513,19 @@ OIDN_API_NAMESPACE_BEGIN
return 0;
}
OIDN_API void oidnSetDeviceErrorFunction(OIDNDevice deviceC, OIDNErrorFunction func, void* userPtr)
OIDN_API void oidnSetDeviceErrorFunction(OIDNDevice hDevice, OIDNErrorFunction func, void* userPtr)
{
Device* device = reinterpret_cast<Device*>(deviceC);
Device* device = reinterpret_cast<Device*>(hDevice);
OIDN_TRY
checkHandle(deviceC);
checkHandle(hDevice);
OIDN_LOCK_DEVICE(device);
device->setErrorFunction(reinterpret_cast<ErrorFunction>(func), userPtr);
OIDN_CATCH_DEVICE(device)
}
OIDN_API OIDNError oidnGetDeviceError(OIDNDevice deviceC, const char** outMessage)
OIDN_API OIDNError oidnGetDeviceError(OIDNDevice hDevice, const char** outMessage)
{
Device* device = reinterpret_cast<Device*>(deviceC);
Device* device = reinterpret_cast<Device*>(hDevice);
OIDN_TRY
return static_cast<OIDNError>(Device::getError(device, outMessage));
OIDN_CATCH_DEVICE(device)
@@ -541,36 +533,32 @@ OIDN_API_NAMESPACE_BEGIN
return OIDN_ERROR_UNKNOWN;
}
OIDN_API void oidnCommitDevice(OIDNDevice deviceC)
OIDN_API void oidnCommitDevice(OIDNDevice hDevice)
{
Device* device = reinterpret_cast<Device*>(deviceC);
Device* device = reinterpret_cast<Device*>(hDevice);
OIDN_TRY
checkHandle(deviceC);
checkHandle(hDevice);
OIDN_LOCK_DEVICE(device);
device->commit();
OIDN_CATCH_DEVICE(device)
}
OIDN_API void oidnSyncDevice(OIDNDevice deviceC)
OIDN_API void oidnSyncDevice(OIDNDevice hDevice)
{
Device* device = reinterpret_cast<Device*>(deviceC);
Device* device = reinterpret_cast<Device*>(hDevice);
OIDN_TRY
checkHandle(deviceC);
checkHandle(hDevice);
OIDN_LOCK_DEVICE(device);
device->checkCommitted();
device->waitAndThrow();
OIDN_CATCH_DEVICE(device)
}
// -----------------------------------------------------------------------------------------------
// Buffer
// -----------------------------------------------------------------------------------------------
OIDN_API OIDNBuffer oidnNewBuffer(OIDNDevice deviceC, size_t byteSize)
OIDN_API OIDNBuffer oidnNewBuffer(OIDNDevice hDevice, size_t byteSize)
{
Device* device = reinterpret_cast<Device*>(deviceC);
Device* device = reinterpret_cast<Device*>(hDevice);
OIDN_TRY
checkHandle(deviceC);
checkHandle(hDevice);
OIDN_LOCK_DEVICE(device);
device->checkCommitted();
Ref<Buffer> buffer = device->newUserBuffer(byteSize, Storage::Undefined);
@@ -579,11 +567,11 @@ OIDN_API_NAMESPACE_BEGIN
return nullptr;
}
OIDN_API OIDNBuffer oidnNewBufferWithStorage(OIDNDevice deviceC, size_t byteSize, OIDNStorage storage)
OIDN_API OIDNBuffer oidnNewBufferWithStorage(OIDNDevice hDevice, size_t byteSize, OIDNStorage storage)
{
Device* device = reinterpret_cast<Device*>(deviceC);
Device* device = reinterpret_cast<Device*>(hDevice);
OIDN_TRY
checkHandle(deviceC);
checkHandle(hDevice);
OIDN_LOCK_DEVICE(device);
device->checkCommitted();
Ref<Buffer> buffer = device->newUserBuffer(byteSize, static_cast<Storage>(storage));
@@ -592,11 +580,11 @@ OIDN_API_NAMESPACE_BEGIN
return nullptr;
}
OIDN_API OIDNBuffer oidnNewSharedBuffer(OIDNDevice deviceC, void* devPtr, size_t byteSize)
OIDN_API OIDNBuffer oidnNewSharedBuffer(OIDNDevice hDevice, void* devPtr, size_t byteSize)
{
Device* device = reinterpret_cast<Device*>(deviceC);
Device* device = reinterpret_cast<Device*>(hDevice);
OIDN_TRY
checkHandle(deviceC);
checkHandle(hDevice);
OIDN_LOCK_DEVICE(device);
device->checkCommitted();
Ref<Buffer> buffer = device->newUserBuffer(devPtr, byteSize);
@@ -605,49 +593,49 @@ OIDN_API_NAMESPACE_BEGIN
return nullptr;
}
OIDN_API OIDNBuffer oidnNewSharedBufferFromFD(OIDNDevice deviceC,
OIDNExternalMemoryTypeFlags fdTypeC,
OIDN_API OIDNBuffer oidnNewSharedBufferFromFD(OIDNDevice hDevice,
OIDNExternalMemoryTypeFlag fdType,
int fd, size_t byteSize)
{
Device* device = reinterpret_cast<Device*>(deviceC);
Device* device = reinterpret_cast<Device*>(hDevice);
OIDN_TRY
checkHandle(deviceC);
checkHandle(hDevice);
OIDN_LOCK_DEVICE(device);
device->checkCommitted();
const ExternalMemoryTypeFlags fdType = static_cast<ExternalMemoryTypeFlags>(fdTypeC);
if ((fdType & device->getExternalMemoryTypes()) != fdType)
if (!(static_cast<ExternalMemoryTypeFlag>(fdType) & device->getExternalMemoryTypes()))
throw Exception(Error::InvalidArgument, "external memory type not supported by the device");
Ref<Buffer> buffer = device->newExternalUserBuffer(fdType, fd, byteSize);
Ref<Buffer> buffer = device->newExternalUserBuffer(
static_cast<ExternalMemoryTypeFlag>(fdType), fd, byteSize);
return reinterpret_cast<OIDNBuffer>(buffer.detach());
OIDN_CATCH_DEVICE(device)
return nullptr;
}
OIDN_API OIDNBuffer oidnNewSharedBufferFromWin32Handle(OIDNDevice deviceC,
OIDNExternalMemoryTypeFlags handleTypeC,
OIDN_API OIDNBuffer oidnNewSharedBufferFromWin32Handle(OIDNDevice hDevice,
OIDNExternalMemoryTypeFlag handleType,
void* handle, const void* name, size_t byteSize)
{
Device* device = reinterpret_cast<Device*>(deviceC);
Device* device = reinterpret_cast<Device*>(hDevice);
OIDN_TRY
checkHandle(deviceC);
checkHandle(hDevice);
OIDN_LOCK_DEVICE(device);
device->checkCommitted();
const ExternalMemoryTypeFlags handleType = static_cast<ExternalMemoryTypeFlags>(handleTypeC);
if ((handleType & device->getExternalMemoryTypes()) != handleType)
if (!(static_cast<ExternalMemoryTypeFlag>(handleType) & device->getExternalMemoryTypes()))
throw Exception(Error::InvalidArgument, "external memory type not supported by the device");
if ((!handle && !name) || (handle && name))
throw Exception(Error::InvalidArgument, "exactly one of the external memory handle and name must be non-null");
Ref<Buffer> buffer = device->newExternalUserBuffer(handleType, handle, name, byteSize);
Ref<Buffer> buffer = device->newExternalUserBuffer(
static_cast<ExternalMemoryTypeFlag>(handleType), handle, name, byteSize);
return reinterpret_cast<OIDNBuffer>(buffer.detach());
OIDN_CATCH_DEVICE(device)
return nullptr;
}
OIDN_API OIDNBuffer oidnNewSharedBufferFromMetal(OIDNDevice deviceC, MTLBuffer_id mtlBuffer)
OIDN_API OIDNBuffer oidnNewSharedBufferFromMetal(OIDNDevice hDevice, MTLBuffer_id mtlBuffer)
{
Device* device = reinterpret_cast<Device*>(deviceC);
Device* device = reinterpret_cast<Device*>(hDevice);
OIDN_TRY
checkHandle(deviceC);
checkHandle(hDevice);
OIDN_LOCK_DEVICE(device);
device->checkCommitted();
Ref<Buffer> buffer = device->newNativeUserBuffer(mtlBuffer);
@@ -656,88 +644,88 @@ OIDN_API_NAMESPACE_BEGIN
return nullptr;
}
OIDN_API void oidnRetainBuffer(OIDNBuffer bufferC)
OIDN_API void oidnRetainBuffer(OIDNBuffer hBuffer)
{
Buffer* buffer = reinterpret_cast<Buffer*>(bufferC);
Buffer* buffer = reinterpret_cast<Buffer*>(hBuffer);
retainObject(buffer);
}
OIDN_API void oidnReleaseBuffer(OIDNBuffer bufferC)
OIDN_API void oidnReleaseBuffer(OIDNBuffer hBuffer)
{
Buffer* buffer = reinterpret_cast<Buffer*>(bufferC);
Buffer* buffer = reinterpret_cast<Buffer*>(hBuffer);
releaseObject(buffer);
}
OIDN_API void oidnReadBuffer(OIDNBuffer bufferC, size_t byteOffset, size_t byteSize, void* dstHostPtr)
OIDN_API void oidnReadBuffer(OIDNBuffer hBuffer, size_t byteOffset, size_t byteSize, void* dstHostPtr)
{
Buffer* buffer = reinterpret_cast<Buffer*>(bufferC);
Buffer* buffer = reinterpret_cast<Buffer*>(hBuffer);
OIDN_TRY
checkHandle(bufferC);
checkHandle(hBuffer);
OIDN_LOCK_DEVICE(buffer);
buffer->read(byteOffset, byteSize, dstHostPtr);
OIDN_CATCH_DEVICE(buffer);
}
OIDN_API void oidnReadBufferAsync(OIDNBuffer bufferC,
OIDN_API void oidnReadBufferAsync(OIDNBuffer hBuffer,
size_t byteOffset, size_t byteSize, void* dstHostPtr)
{
Buffer* buffer = reinterpret_cast<Buffer*>(bufferC);
Buffer* buffer = reinterpret_cast<Buffer*>(hBuffer);
OIDN_TRY
checkHandle(bufferC);
checkHandle(hBuffer);
OIDN_LOCK_DEVICE(buffer);
buffer->read(byteOffset, byteSize, dstHostPtr, SyncMode::Async);
OIDN_CATCH_DEVICE(buffer);
}
OIDN_API void oidnWriteBuffer(OIDNBuffer bufferC,
OIDN_API void oidnWriteBuffer(OIDNBuffer hBuffer,
size_t byteOffset, size_t byteSize, const void* srcHostPtr)
{
Buffer* buffer = reinterpret_cast<Buffer*>(bufferC);
Buffer* buffer = reinterpret_cast<Buffer*>(hBuffer);
OIDN_TRY
checkHandle(bufferC);
checkHandle(hBuffer);
OIDN_LOCK_DEVICE(buffer);
buffer->write(byteOffset, byteSize, srcHostPtr);
OIDN_CATCH_DEVICE(buffer);
}
OIDN_API void oidnWriteBufferAsync(OIDNBuffer bufferC,
OIDN_API void oidnWriteBufferAsync(OIDNBuffer hBuffer,
size_t byteOffset, size_t byteSize, const void* srcHostPtr)
{
Buffer* buffer = reinterpret_cast<Buffer*>(bufferC);
Buffer* buffer = reinterpret_cast<Buffer*>(hBuffer);
OIDN_TRY
checkHandle(bufferC);
checkHandle(hBuffer);
OIDN_LOCK_DEVICE(buffer);
buffer->write(byteOffset, byteSize, srcHostPtr, SyncMode::Async);
OIDN_CATCH_DEVICE(buffer);
}
OIDN_API size_t oidnGetBufferSize(OIDNBuffer bufferC)
OIDN_API size_t oidnGetBufferSize(OIDNBuffer hBuffer)
{
Buffer* buffer = reinterpret_cast<Buffer*>(bufferC);
Buffer* buffer = reinterpret_cast<Buffer*>(hBuffer);
OIDN_TRY
checkHandle(bufferC);
checkHandle(hBuffer);
OIDN_LOCK_DEVICE(buffer);
return buffer->getByteSize();
OIDN_CATCH_DEVICE(buffer)
return 0;
}
OIDN_API OIDNStorage oidnGetBufferStorage(OIDNBuffer bufferC)
OIDN_API OIDNStorage oidnGetBufferStorage(OIDNBuffer hBuffer)
{
Buffer* buffer = reinterpret_cast<Buffer*>(bufferC);
Buffer* buffer = reinterpret_cast<Buffer*>(hBuffer);
OIDN_TRY
checkHandle(bufferC);
checkHandle(hBuffer);
OIDN_LOCK_DEVICE(buffer);
return static_cast<OIDNStorage>(buffer->getStorage());
OIDN_CATCH_DEVICE(buffer)
return OIDN_STORAGE_UNDEFINED;
}
OIDN_API void* oidnGetBufferData(OIDNBuffer bufferC)
OIDN_API void* oidnGetBufferData(OIDNBuffer hBuffer)
{
Buffer* buffer = reinterpret_cast<Buffer*>(bufferC);
Buffer* buffer = reinterpret_cast<Buffer*>(hBuffer);
OIDN_TRY
checkHandle(bufferC);
checkHandle(hBuffer);
OIDN_LOCK_DEVICE(buffer);
if (void* ptr = buffer->getHostPtr())
return ptr;
@@ -747,105 +735,11 @@ OIDN_API_NAMESPACE_BEGIN
return nullptr;
}
// -----------------------------------------------------------------------------------------------
// Semaphore
// -----------------------------------------------------------------------------------------------
OIDN_API OIDNSemaphore oidnNewSharedSemaphoreFromFD(OIDNDevice deviceC,
OIDNExternalSemaphoreTypeFlags fdTypeC,
int fd)
OIDN_API OIDNFilter oidnNewFilter(OIDNDevice hDevice, const char* type)
{
Device* device = reinterpret_cast<Device*>(deviceC);
Device* device = reinterpret_cast<Device*>(hDevice);
OIDN_TRY
checkHandle(deviceC);
OIDN_LOCK_DEVICE(device);
device->checkCommitted();
const ExternalSemaphoreTypeFlags fdType = static_cast<ExternalSemaphoreTypeFlag>(fdTypeC);
if ((fdType & device->getExternalSemaphoreTypes()) != fdType)
throw Exception(Error::InvalidArgument, "external semaphore type not supported by the device");
Ref<Semaphore> semaphore = device->newExternalSemaphore(fdType, fd);
return reinterpret_cast<OIDNSemaphore>(semaphore.detach());
OIDN_CATCH_DEVICE(device)
return nullptr;
}
OIDN_API OIDNSemaphore oidnNewSharedSemaphoreFromWin32Handle(OIDNDevice deviceC,
OIDNExternalSemaphoreTypeFlags handleTypeC,
void* handle, const void* name)
{
Device* device = reinterpret_cast<Device*>(deviceC);
OIDN_TRY
checkHandle(deviceC);
OIDN_LOCK_DEVICE(device);
device->checkCommitted();
const ExternalSemaphoreTypeFlags handleType = static_cast<ExternalSemaphoreTypeFlag>(handleTypeC);
if ((handleType & device->getExternalSemaphoreTypes()) != handleType)
throw Exception(Error::InvalidArgument, "external semaphore type not supported by the device");
if ((!handle && !name) || (handle && name))
throw Exception(Error::InvalidArgument, "exactly one of the external semaphore handle and name must be non-null");
Ref<Semaphore> semaphore = device->newExternalSemaphore(handleType, handle, name);
return reinterpret_cast<OIDNSemaphore>(semaphore.detach());
OIDN_CATCH_DEVICE(device)
return nullptr;
}
OIDN_API void oidnSignalSemaphoresAsync(OIDNDevice deviceC,
const OIDNSemaphore* semaphoresC,
const uint64_t* values,
int numSemaphores)
{
Device* device = reinterpret_cast<Device*>(deviceC);
OIDN_TRY
checkHandle(deviceC);
OIDN_LOCK_DEVICE(device);
device->checkCommitted();
device->submitSignalSemaphores(
reinterpret_cast<Semaphore* const*>(semaphoresC),
values,
numSemaphores);
OIDN_CATCH_DEVICE(device)
}
OIDN_API void oidnWaitSemaphoresAsync(OIDNDevice deviceC,
const OIDNSemaphore* semaphoresC,
const uint64_t* values,
const uint32_t* timeoutsMs,
int numSemaphores)
{
Device* device = reinterpret_cast<Device*>(deviceC);
OIDN_TRY
checkHandle(deviceC);
OIDN_LOCK_DEVICE(device);
device->checkCommitted();
device->submitWaitSemaphores(
reinterpret_cast<Semaphore* const*>(semaphoresC),
values,
timeoutsMs,
numSemaphores);
OIDN_CATCH_DEVICE(device)
}
OIDN_API void oidnRetainSemaphore(OIDNSemaphore semaphoreC)
{
Semaphore* semaphore = reinterpret_cast<Semaphore*>(semaphoreC);
retainObject(semaphore);
}
OIDN_API void oidnReleaseSemaphore(OIDNSemaphore semaphoreC)
{
Semaphore* semaphore = reinterpret_cast<Semaphore*>(semaphoreC);
releaseObject(semaphore);
}
// -----------------------------------------------------------------------------------------------
// Filter
// -----------------------------------------------------------------------------------------------
OIDN_API OIDNFilter oidnNewFilter(OIDNDevice deviceC, const char* type)
{
Device* device = reinterpret_cast<Device*>(deviceC);
OIDN_TRY
checkHandle(deviceC);
checkHandle(hDevice);
OIDN_LOCK_DEVICE(device);
device->checkCommitted();
checkString(type);
@@ -855,31 +749,31 @@ OIDN_API_NAMESPACE_BEGIN
return nullptr;
}
OIDN_API void oidnRetainFilter(OIDNFilter filterC)
OIDN_API void oidnRetainFilter(OIDNFilter hFilter)
{
Filter* filter = reinterpret_cast<Filter*>(filterC);
Filter* filter = reinterpret_cast<Filter*>(hFilter);
retainObject(filter);
}
OIDN_API void oidnReleaseFilter(OIDNFilter filterC)
OIDN_API void oidnReleaseFilter(OIDNFilter hFilter)
{
Filter* filter = reinterpret_cast<Filter*>(filterC);
Filter* filter = reinterpret_cast<Filter*>(hFilter);
releaseObject(filter);
}
OIDN_API void oidnSetFilterImage(OIDNFilter filterC, const char* name,
OIDNBuffer bufferC, OIDNFormat format,
OIDN_API void oidnSetFilterImage(OIDNFilter hFilter, const char* name,
OIDNBuffer hBuffer, OIDNFormat format,
size_t width, size_t height,
size_t byteOffset,
size_t pixelByteStride, size_t rowByteStride)
{
Filter* filter = reinterpret_cast<Filter*>(filterC);
Filter* filter = reinterpret_cast<Filter*>(hFilter);
OIDN_TRY
checkHandle(filterC);
checkHandle(hFilter);
OIDN_LOCK_DEVICE(filter);
checkString(name);
checkHandle(bufferC);
Ref<Buffer> buffer = reinterpret_cast<Buffer*>(bufferC);
checkHandle(hBuffer);
Ref<Buffer> buffer = reinterpret_cast<Buffer*>(hBuffer);
if (buffer->getDevice() != filter->getDevice())
throw Exception(Error::InvalidArgument, "the specified objects are bound to different devices");
auto image = makeRef<Image>(buffer, static_cast<Format>(format),
@@ -889,15 +783,15 @@ OIDN_API_NAMESPACE_BEGIN
OIDN_CATCH_DEVICE(filter)
}
OIDN_API void oidnSetSharedFilterImage(OIDNFilter filterC, const char* name,
OIDN_API void oidnSetSharedFilterImage(OIDNFilter hFilter, const char* name,
void* devPtr, OIDNFormat format,
size_t width, size_t height,
size_t byteOffset,
size_t pixelByteStride, size_t rowByteStride)
{
Filter* filter = reinterpret_cast<Filter*>(filterC);
Filter* filter = reinterpret_cast<Filter*>(hFilter);
OIDN_TRY
checkHandle(filterC);
checkHandle(hFilter);
OIDN_LOCK_DEVICE(filter);
checkString(name);
auto image = makeRef<Image>(devPtr, static_cast<Format>(format),
@@ -907,23 +801,23 @@ OIDN_API_NAMESPACE_BEGIN
OIDN_CATCH_DEVICE(filter)
}
OIDN_API void oidnUnsetFilterImage(OIDNFilter filterC, const char* name)
OIDN_API void oidnUnsetFilterImage(OIDNFilter hFilter, const char* name)
{
Filter* filter = reinterpret_cast<Filter*>(filterC);
Filter* filter = reinterpret_cast<Filter*>(hFilter);
OIDN_TRY
checkHandle(filterC);
checkHandle(hFilter);
OIDN_LOCK_DEVICE(filter);
checkString(name);
filter->unsetImage(name);
OIDN_CATCH_DEVICE(filter)
}
OIDN_API void oidnSetSharedFilterData(OIDNFilter filterC, const char* name,
OIDN_API void oidnSetSharedFilterData(OIDNFilter hFilter, const char* name,
void* hostPtr, size_t byteSize)
{
Filter* filter = reinterpret_cast<Filter*>(filterC);
Filter* filter = reinterpret_cast<Filter*>(hFilter);
OIDN_TRY
checkHandle(filterC);
checkHandle(hFilter);
OIDN_LOCK_DEVICE(filter);
checkString(name);
Data data(hostPtr, byteSize);
@@ -931,44 +825,44 @@ OIDN_API_NAMESPACE_BEGIN
OIDN_CATCH_DEVICE(filter)
}
OIDN_API void oidnUpdateFilterData(OIDNFilter filterC, const char* name)
OIDN_API void oidnUpdateFilterData(OIDNFilter hFilter, const char* name)
{
Filter* filter = reinterpret_cast<Filter*>(filterC);
Filter* filter = reinterpret_cast<Filter*>(hFilter);
OIDN_TRY
checkHandle(filterC);
checkHandle(hFilter);
OIDN_LOCK_DEVICE(filter);
checkString(name);
filter->updateData(name);
OIDN_CATCH_DEVICE(filter)
}
OIDN_API void oidnUnsetFilterData(OIDNFilter filterC, const char* name)
OIDN_API void oidnUnsetFilterData(OIDNFilter hFilter, const char* name)
{
Filter* filter = reinterpret_cast<Filter*>(filterC);
Filter* filter = reinterpret_cast<Filter*>(hFilter);
OIDN_TRY
checkHandle(filterC);
checkHandle(hFilter);
OIDN_LOCK_DEVICE(filter);
checkString(name);
filter->unsetData(name);
OIDN_CATCH_DEVICE(filter)
}
OIDN_API void oidnSetFilterBool(OIDNFilter filterC, const char* name, bool value)
OIDN_API void oidnSetFilterBool(OIDNFilter hFilter, const char* name, bool value)
{
Filter* filter = reinterpret_cast<Filter*>(filterC);
Filter* filter = reinterpret_cast<Filter*>(hFilter);
OIDN_TRY
checkHandle(filterC);
checkHandle(hFilter);
OIDN_LOCK_DEVICE(filter);
checkString(name);
filter->setInt(name, int(value));
OIDN_CATCH_DEVICE(filter)
}
OIDN_API bool oidnGetFilterBool(OIDNFilter filterC, const char* name)
OIDN_API bool oidnGetFilterBool(OIDNFilter hFilter, const char* name)
{
Filter* filter = reinterpret_cast<Filter*>(filterC);
Filter* filter = reinterpret_cast<Filter*>(hFilter);
OIDN_TRY
checkHandle(filterC);
checkHandle(hFilter);
OIDN_LOCK_DEVICE(filter);
checkString(name);
return filter->getInt(name);
@@ -976,22 +870,22 @@ OIDN_API_NAMESPACE_BEGIN
return false;
}
OIDN_API void oidnSetFilterInt(OIDNFilter filterC, const char* name, int value)
OIDN_API void oidnSetFilterInt(OIDNFilter hFilter, const char* name, int value)
{
Filter* filter = reinterpret_cast<Filter*>(filterC);
Filter* filter = reinterpret_cast<Filter*>(hFilter);
OIDN_TRY
checkHandle(filterC);
checkHandle(hFilter);
OIDN_LOCK_DEVICE(filter);
checkString(name);
filter->setInt(name, value);
OIDN_CATCH_DEVICE(filter)
}
OIDN_API int oidnGetFilterInt(OIDNFilter filterC, const char* name)
OIDN_API int oidnGetFilterInt(OIDNFilter hFilter, const char* name)
{
Filter* filter = reinterpret_cast<Filter*>(filterC);
Filter* filter = reinterpret_cast<Filter*>(hFilter);
OIDN_TRY
checkHandle(filterC);
checkHandle(hFilter);
OIDN_LOCK_DEVICE(filter);
checkString(name);
return filter->getInt(name);
@@ -999,22 +893,22 @@ OIDN_API_NAMESPACE_BEGIN
return 0;
}
OIDN_API void oidnSetFilterFloat(OIDNFilter filterC, const char* name, float value)
OIDN_API void oidnSetFilterFloat(OIDNFilter hFilter, const char* name, float value)
{
Filter* filter = reinterpret_cast<Filter*>(filterC);
Filter* filter = reinterpret_cast<Filter*>(hFilter);
OIDN_TRY
checkHandle(filterC);
checkHandle(hFilter);
OIDN_LOCK_DEVICE(filter);
checkString(name);
filter->setFloat(name, value);
OIDN_CATCH_DEVICE(filter)
}
OIDN_API float oidnGetFilterFloat(OIDNFilter filterC, const char* name)
OIDN_API float oidnGetFilterFloat(OIDNFilter hFilter, const char* name)
{
Filter* filter = reinterpret_cast<Filter*>(filterC);
Filter* filter = reinterpret_cast<Filter*>(hFilter);
OIDN_TRY
checkHandle(filterC);
checkHandle(hFilter);
OIDN_LOCK_DEVICE(filter);
checkString(name);
return filter->getFloat(name);
@@ -1022,54 +916,54 @@ OIDN_API_NAMESPACE_BEGIN
return 0;
}
OIDN_API void oidnSetFilterProgressMonitorFunction(OIDNFilter filterC,
OIDN_API void oidnSetFilterProgressMonitorFunction(OIDNFilter hFilter,
OIDNProgressMonitorFunction func, void* userPtr)
{
Filter* filter = reinterpret_cast<Filter*>(filterC);
Filter* filter = reinterpret_cast<Filter*>(hFilter);
OIDN_TRY
checkHandle(filterC);
checkHandle(hFilter);
OIDN_LOCK_DEVICE(filter);
filter->setProgressMonitorFunction(func, userPtr);
OIDN_CATCH_DEVICE(filter)
}
OIDN_API void oidnCommitFilter(OIDNFilter filterC)
OIDN_API void oidnCommitFilter(OIDNFilter hFilter)
{
Filter* filter = reinterpret_cast<Filter*>(filterC);
Filter* filter = reinterpret_cast<Filter*>(hFilter);
OIDN_TRY
checkHandle(filterC);
checkHandle(hFilter);
OIDN_LOCK_DEVICE(filter);
filter->commit();
OIDN_CATCH_DEVICE(filter)
}
OIDN_API void oidnExecuteFilter(OIDNFilter filterC)
OIDN_API void oidnExecuteFilter(OIDNFilter hFilter)
{
Filter* filter = reinterpret_cast<Filter*>(filterC);
Filter* filter = reinterpret_cast<Filter*>(hFilter);
OIDN_TRY
checkHandle(filterC);
checkHandle(hFilter);
OIDN_LOCK_DEVICE(filter);
filter->execute();
OIDN_CATCH_DEVICE(filter)
}
OIDN_API void oidnExecuteFilterAsync(OIDNFilter filterC)
OIDN_API void oidnExecuteFilterAsync(OIDNFilter hFilter)
{
Filter* filter = reinterpret_cast<Filter*>(filterC);
Filter* filter = reinterpret_cast<Filter*>(hFilter);
OIDN_TRY
checkHandle(filterC);
checkHandle(hFilter);
OIDN_LOCK_DEVICE(filter);
filter->execute(SyncMode::Async);
OIDN_CATCH_DEVICE(filter)
}
OIDN_API void oidnExecuteSYCLFilterAsync(OIDNFilter filterC,
OIDN_API void oidnExecuteSYCLFilterAsync(OIDNFilter hFilter,
const sycl::event* depEvents, int numDepEvents,
sycl::event* doneEvent)
{
Filter* filter = reinterpret_cast<Filter*>(filterC);
Filter* filter = reinterpret_cast<Filter*>(hFilter);
OIDN_TRY
checkHandle(filterC);
checkHandle(hFilter);
OIDN_LOCK_DEVICE(filter);
// Check whether the filter belongs to a SYCL device
+36 -1
View File
@@ -11,4 +11,39 @@ endmacro()
oidn_add_app(oidnDenoise oidnDenoise.cpp)
oidn_add_app(oidnBenchmark oidnBenchmark.cpp)
oidn_add_app(oidnTest oidnTest.cpp "${PROJECT_SOURCE_DIR}/external/catch.hpp")
oidn_add_app(oidnTest oidnTest.cpp "${PROJECT_SOURCE_DIR}/external/catch.hpp")
if(OIDN_INTEROP_TEST)
find_package(OpenImageIO REQUIRED)
if(WIN32)
set(CMAKE_MSVC_RUNTIME_LIBRARY "MultiThreadedDLL")
add_executable(oidnInteropTest oidnInteropTest.cpp ${OIDN_RESOURCE_FILE})
set_property(TARGET oidnInteropTest PROPERTY CXX_STANDARD 17)
target_link_libraries(oidnInteropTest PRIVATE
OpenImageIO::OpenImageIO
OpenImageDenoise_common
#OpenImageDenoise_utils
OpenImageDenoise
d3d12
dxgi
dxguid
)
else()
find_package(Vulkan REQUIRED)
add_executable(oidnInteropTest oidnInteropTest.cpp ${OIDN_RESOURCE_FILE})
set_property(TARGET oidnInteropTest PROPERTY CXX_STANDARD 17)
target_link_libraries(oidnInteropTest PRIVATE
Vulkan::Vulkan
OpenImageIO::OpenImageIO
OpenImageDenoise_common
OpenImageDenoise_utils
OpenImageDenoise
)
endif()
endif()
+1 -1
View File
@@ -86,7 +86,7 @@ int main(int argc, char* argv[])
bool directional = false;
float inputScale = std::numeric_limits<float>::quiet_NaN();
bool cleanAux = false;
DataType dataType = DataType::Undefined;
DataType dataType = DataType::Void;
int numRuns = 1;
int numThreads = -1;
int setAffinity = -1;
+150
View File
@@ -0,0 +1,150 @@
#ifdef _WIN32
#include "oidnInteropTestDX.h"
#else
#include "oidnInteropTestVk.h"
#endif
std::string usage() {
return
"Usage:\n"
" oidnInteropTest - run oidn using directX/Vulkan interop\n"
"\n"
"SYNOPSIS\n"
" oidnInteropTest --json <path> [OPTION]\n"
"\n"
"OPTIONS\n"
" --scale <float> - scaling factor\n"
" --out <ext> - extension of the out files, i.e. <path>.out.<n>.exr, default = out\n"
" --help - print this help\n"
" --verbose - vebose console output\n"
;
}
bool has_arg(const std::string& pname, std::vector<std::string>& args) {
auto p = std::find(args.begin(), args.end(), pname);
return p != args.end();
}
bool has_param(const std::string& pname, std::vector<std::string>& args) {
auto p = std::find(args.begin(), args.end(), pname);
return p != args.end() && (p+1) != args.end();
}
std::string parse(const std::string& pname, std::vector<std::string>& args) {
auto p = std::find(args.begin(), args.end(), pname);
if (p == args.end()) {
throw std::runtime_error("could not parse parameter: " + pname + "\n" + usage());
}
++p;
if (p == args.end()) {
throw std::runtime_error("could not parse parameter: " + pname + "\n" + usage());
}
return *p;
}
std::string parse_opt(const std::string& pname, std::vector<std::string>& args, std::string def = "") {
return has_param(pname, args)? parse(pname, args) : def;
}
int main(int argc, char* argv[]) {
std::vector<std::string> args(argv + 1, argv + argc);
if (has_arg("--help", args)) {
std::cout << usage() << std::endl;
return 0;
}
auto path_in = parse("--in", args);
auto path_color = path_in + ".hdr.exr";
auto path_albedo = path_in + ".alb1.exr";
auto path_normal = path_in + ".nrm1.exr";
auto path_out = parse_opt("--out", args, "./out.exr");
auto scale = std::stof(parse_opt("--scale", args, "2.0f"));
std::cout << "oidn interop test" << std::endl;
std::cout << " color: " << path_color << std::endl;
std::cout << " albedo: " << path_albedo << std::endl;
std::cout << " normal: " << path_normal << std::endl;
std::cout << " output: " << path_out << std::endl;
std::cout << " scale: " << scale << std::endl;
#ifdef _WIN32
InteropTestDX test;
#else
InteropTestVk test;
#endif
std::cout << "Setting up OIDN device" << std::endl;
test.setup_oidn_device();
oidn::DeviceType deviceType = test.oidn_device.get<oidn::DeviceType>("type");
const int versionMajor = test.oidn_device.get<int>("versionMajor");
const int versionMinor = test.oidn_device.get<int>("versionMinor");
const int versionPatch = test.oidn_device.get<int>("versionPatch");
std::cout << std::setw(15) << std::left << "OIDN Device: ";
switch (deviceType)
{
case oidn::DeviceType::Default: std::cout << std::setw(30) << std::left << "default"; break;
case oidn::DeviceType::CPU: std::cout << std::setw(30) << std::left << "CPU"; break;
case oidn::DeviceType::SYCL: std::cout << std::setw(30) << std::left << "SYCL"; break;
case oidn::DeviceType::CUDA: std::cout << std::setw(30) << std::left << "CUDA"; break;
case oidn::DeviceType::HIP: std::cout << std::setw(30) << std::left << "HIP"; break;
case oidn::DeviceType::Metal: std::cout << std::setw(30) << std::left << "Metal"; break;
default:
throw std::invalid_argument("invalid device type");
}
std::cout << "Version: " << versionMajor << "." << versionMinor << "." << versionPatch << std::endl;
auto color = test.load_exr(path_color);
auto albedo = test.load_exr(path_albedo);
auto normal = test.load_exr(path_normal);
uint32_t render_width, render_height, target_width, target_height;
{
render_width = color.get_width();
render_height = color.get_height();
target_width = render_width * scale;
target_height = render_height * scale;
std::cout << render_height << " x " << render_width << " -> "
<< target_height << " x " << target_width << std::endl;
}
ImageBuffer out = test.create_buffer(target_width, target_height, oidn::Format::Half3);
// Initialize the denoising filter
std::cout << "Initializing filter" << std::endl;
oidn::FilterRef filter = test.oidn_device.newFilter("RT");
auto oidn_color = test.create_oidn_buffer(color);
auto oidn_albedo = test.create_oidn_buffer(albedo);
auto oidn_normal = test.create_oidn_buffer(normal);
auto oidn_out = test.create_oidn_buffer(out);
filter.setImage("color", oidn_color, color.get_format(), color.get_width(), color.get_height());
filter.setImage("albedo", oidn_albedo, albedo.get_format(), albedo.get_width(), albedo.get_height());
filter.setImage("normal", oidn_normal, normal.get_format(), normal.get_width(), normal.get_height());
filter.setImage("output", oidn_out, out.get_format(), out.get_width(), out.get_height());
filter.set("hdr", true);
filter.set("inputScale", scale);
filter.set("quality", oidn::Quality::Default);
filter.commit();
filter.execute();
test.save_exr(out, "./out.exr");
return 0;
}
+447
View File
@@ -0,0 +1,447 @@
#include <windows.h>
#include <wrl.h>
#include <d3d12.h>
#include <dxgi1_6.h>
#include <vector>
#include <optional>
#include <set>
#include <iostream>
#include <iomanip>
#include <algorithm>
#include <cstring>
#include "OpenImageIO/imageio.h"
#include "OpenImageDenoise/oidn.hpp"
using Microsoft::WRL::ComPtr;
class CmdList {
public:
CmdList(ComPtr<ID3D12Device> device) {
if (FAILED(device->CreateCommandAllocator(D3D12_COMMAND_LIST_TYPE_DIRECT, IID_PPV_ARGS(&commandAllocator)))) {
throw std::runtime_error("Failed to create command allocator.");
}
if (FAILED(device->CreateCommandList(0, D3D12_COMMAND_LIST_TYPE_DIRECT, commandAllocator.Get(), nullptr, IID_PPV_ARGS(&commandList)))) {
throw std::runtime_error("Failed to create command list.");
}
}
CmdList(const CmdList& other) = delete;
CmdList& operator=(const CmdList& other) = delete;
CmdList(CmdList&& other) noexcept : commandList(std::move(other.commandList)) {}
CmdList& operator=(CmdList&& other) noexcept {
if (this != &other) {
commandList = std::move(other.commandList);
}
return *this;
}
~CmdList() {
if (commandList) {
std::cout << "Releasing command list." << std::endl;
commandList.Reset();
commandAllocator.Reset();
}
}
public:
ID3D12GraphicsCommandList* operator->() { return commandList.Get(); }
ID3D12GraphicsCommandList& operator*() { return *commandList.Get(); }
void closeAndWait(ComPtr<ID3D12Device>& device, ComPtr<ID3D12CommandQueue>& commandQueue) {
// Close and execute command list
if (FAILED(commandList->Close())) {
throw std::runtime_error("Failed to close command list.");
}
ID3D12CommandList* ppCommandLists[] = { commandList.Get() };
commandQueue->ExecuteCommandLists(1, ppCommandLists);
// Create fence for synchronization
ComPtr<ID3D12Fence> fence;
if (FAILED(device->CreateFence(0, D3D12_FENCE_FLAG_NONE, IID_PPV_ARGS(&fence)))) {
throw std::runtime_error("Failed to create fence.");
}
HANDLE fenceEvent = CreateEvent(nullptr, FALSE, FALSE, nullptr);
if (!fenceEvent) {
throw std::runtime_error("Failed to create fence event.");
}
UINT64 fenceValue = 1;
if (FAILED(commandQueue->Signal(fence.Get(), fenceValue))) {
CloseHandle(fenceEvent);
throw std::runtime_error("Failed to signal fence.");
}
if (fence->GetCompletedValue() < fenceValue) {
if (FAILED(fence->SetEventOnCompletion(fenceValue, fenceEvent))) {
CloseHandle(fenceEvent);
throw std::runtime_error("Failed to set event on completion.");
}
WaitForSingleObject(fenceEvent, INFINITE);
}
CloseHandle(fenceEvent);
}
private:
ComPtr<ID3D12CommandAllocator> commandAllocator;
ComPtr<ID3D12GraphicsCommandList> commandList;
};
class Buffer {
public:
Buffer(ComPtr<ID3D12Device> device, size_t size, D3D12_HEAP_TYPE heapType, D3D12_HEAP_FLAGS heapFlags, D3D12_RESOURCE_FLAGS resourceFlags, D3D12_RESOURCE_STATES state): size(size) {
// Describe the buffer
D3D12_RESOURCE_DESC bufferDesc = {};
bufferDesc.Dimension = D3D12_RESOURCE_DIMENSION_BUFFER;
bufferDesc.Alignment = 0;
bufferDesc.Width = size;
bufferDesc.Height = 1;
bufferDesc.DepthOrArraySize = 1;
bufferDesc.MipLevels = 1;
bufferDesc.Format = DXGI_FORMAT_UNKNOWN;
bufferDesc.SampleDesc.Count = 1;
bufferDesc.SampleDesc.Quality = 0;
bufferDesc.Layout = D3D12_TEXTURE_LAYOUT_ROW_MAJOR;
bufferDesc.Flags = resourceFlags;
// Heap properties for shared resource
D3D12_HEAP_PROPERTIES heapProps = {};
heapProps.Type = heapType;
heapProps.CPUPageProperty = D3D12_CPU_PAGE_PROPERTY_UNKNOWN;
heapProps.MemoryPoolPreference = D3D12_MEMORY_POOL_UNKNOWN;
heapProps.CreationNodeMask = 1;
heapProps.VisibleNodeMask = 1;
// Create the buffer resource
HRESULT hr = device->CreateCommittedResource(
&heapProps,
heapFlags,
&bufferDesc,
state,
nullptr,
IID_PPV_ARGS(&buffer)
);
if (FAILED(hr)) {
throw std::runtime_error("Failed to create buffer resource.");
}
std::cout << "Buffer resource created successfully." << std::endl;
}
static Buffer make_image_buffer(ComPtr<ID3D12Device>& device, uint32_t width, uint32_t height, oidn::Format format) {
size_t size_bytes = 0;
switch (format) {
case oidn::Format::Undefined: size_bytes = 0; break;
case oidn::Format::Float: size_bytes = sizeof(float); break;
case oidn::Format::Float2: size_bytes = sizeof(float)*2; break;
case oidn::Format::Float3: size_bytes = sizeof(float)*3; break;
case oidn::Format::Float4: size_bytes = sizeof(float)*4; break;
case oidn::Format::Half: size_bytes = sizeof(int16_t); break;
case oidn::Format::Half2: size_bytes = sizeof(int16_t)*2; break;
case oidn::Format::Half3: size_bytes = sizeof(int16_t)*3; break;
case oidn::Format::Half4: size_bytes = sizeof(int16_t)*4; break;
default:
throw std::invalid_argument("invalid format");
}
size_bytes = size_bytes * width * height;
return Buffer(device, size_bytes, D3D12_HEAP_TYPE_DEFAULT, D3D12_HEAP_FLAG_SHARED, D3D12_RESOURCE_FLAG_ALLOW_UNORDERED_ACCESS, D3D12_RESOURCE_STATE_UNORDERED_ACCESS);
}
static Buffer make_staging_buffer(ComPtr<ID3D12Device> device, size_t size, char rw) {
D3D12_HEAP_TYPE heapType;
D3D12_RESOURCE_STATES state;
if (rw == 'w') {
heapType = D3D12_HEAP_TYPE_UPLOAD;
state = D3D12_RESOURCE_STATE_GENERIC_READ;
} else if (rw == 'r') {
heapType = D3D12_HEAP_TYPE_READBACK;
state = D3D12_RESOURCE_STATE_COPY_DEST;
} else {
throw std::invalid_argument("rw must be 'r' or 'w'");
}
return Buffer(device, size, heapType, D3D12_HEAP_FLAG_NONE, D3D12_RESOURCE_FLAG_NONE, state);
}
Buffer(const Buffer& other) = delete;
Buffer& operator=(const Buffer& other) = delete;
Buffer(Buffer&& other) noexcept
: size(other.size), buffer(std::move(other.buffer))
{
other.size = 0;
other.buffer = nullptr;
}
Buffer& operator=(Buffer&& other) noexcept
{
if (this != &other)
{
size = other.size;
buffer = std::move(other.buffer);
other.size = 0;
other.buffer = nullptr;
}
return *this;
}
~Buffer() {
if (buffer) {
std::cout << "Releasing buffer resource." << std::endl;
buffer.Reset();
}
}
public:
operator ComPtr<ID3D12Resource>() const { return get(); }
ComPtr<ID3D12Resource> get() const { return buffer; }
size_t get_size() const { return size; }
private:
size_t size;
ComPtr<ID3D12Resource> buffer = nullptr;
};
class ImageBuffer {
public:
ImageBuffer(ComPtr<ID3D12Device>& device, uint32_t width, uint32_t height, oidn::Format format)
: buffer(Buffer::make_image_buffer(device, width, height, format)),
width(width), height(height), format(format) {}
ImageBuffer(const ImageBuffer&) = delete;
ImageBuffer& operator=(const ImageBuffer&) = delete;
ImageBuffer(ImageBuffer&&) = default;
ImageBuffer& operator=(ImageBuffer&&) = default;
~ImageBuffer() = default;
public:
static ImageBuffer load_exr(std::string filename, ComPtr<ID3D12Device>& device, ComPtr<ID3D12CommandQueue>& commandQueue) {
auto in = OIIO::ImageInput::open(filename);
if (!in) {
throw std::runtime_error("failed to open image file");
}
OIIO::ImageSpec spec = in->spec();
if ((spec.channelformat(0) == OIIO::TypeDesc::FLOAT || spec.channelformat(0) == OIIO::TypeDesc::HALF) && (spec.nchannels < 1 || spec.nchannels > 4)) {
throw std::invalid_argument("unsupported format");
}
int width = spec.width;
int height = spec.height;
int nchannels = spec.nchannels;
std::vector<oidn::Format> formats = {oidn::Format::Float, oidn::Format::Float2, oidn::Format::Float3, oidn::Format::Float4,
oidn::Format::Half, oidn::Format::Half2, oidn::Format::Half3, oidn::Format::Half4};
oidn::Format format = formats[spec.channelformat(0) == OIIO::TypeDesc::FLOAT ? nchannels - 1 : nchannels + 3];
ImageBuffer buffer(device, width, height, format);
Buffer stagingBuffer = Buffer::make_staging_buffer(device, buffer.buffer.get_size(), 'w');
auto stagingResource = stagingBuffer.get();
void* mapped = nullptr;
stagingResource->Map(0, nullptr, &mapped);
in->read_image(0, 0, 0, nchannels, spec.channelformat(0), mapped);
D3D12_RANGE readingRange = {0, buffer.buffer.get_size()};
stagingResource->Unmap(0, &readingRange);
in->close();
CmdList cl(device);
// Transition buffer to COPY_DEST
D3D12_RESOURCE_BARRIER barrier = {};
barrier.Type = D3D12_RESOURCE_BARRIER_TYPE_TRANSITION;
barrier.Flags = D3D12_RESOURCE_BARRIER_FLAG_NONE;
barrier.Transition.pResource = buffer.get_buffer().get().Get();
barrier.Transition.StateBefore = D3D12_RESOURCE_STATE_UNORDERED_ACCESS;
barrier.Transition.StateAfter = D3D12_RESOURCE_STATE_COPY_DEST;
barrier.Transition.Subresource = D3D12_RESOURCE_BARRIER_ALL_SUBRESOURCES;
cl->ResourceBarrier(1, &barrier);
// Copy data from staging buffer to buffer
cl->CopyResource(buffer.get_buffer().get().Get(), stagingBuffer.get().Get());
// Transition buffer back to UNORDERED_ACCESS
std::swap(barrier.Transition.StateBefore, barrier.Transition.StateAfter);
cl->ResourceBarrier(1, &barrier);
cl.closeAndWait(device, commandQueue);
return buffer;
}
void save_exr(std::string filename, ComPtr<ID3D12Device>& device, ComPtr<ID3D12CommandQueue>& commandQueue) {
// Create a staging buffer to copy the data to
Buffer stagingBuffer = Buffer::make_staging_buffer(device, buffer.get_size(), 'r');
// Create command list for copy operation
CmdList cl(device);
// Transition buffer to COPY_SOURCE
D3D12_RESOURCE_BARRIER barrier = {};
barrier.Type = D3D12_RESOURCE_BARRIER_TYPE_TRANSITION;
barrier.Flags = D3D12_RESOURCE_BARRIER_FLAG_NONE;
barrier.Transition.pResource = buffer.get().Get();
barrier.Transition.StateBefore = D3D12_RESOURCE_STATE_UNORDERED_ACCESS;
barrier.Transition.StateAfter = D3D12_RESOURCE_STATE_COPY_SOURCE;
barrier.Transition.Subresource = D3D12_RESOURCE_BARRIER_ALL_SUBRESOURCES;
cl->ResourceBarrier(1, &barrier);
// Copy buffer to staging
cl->CopyResource(stagingBuffer.get().Get(), buffer.get().Get());
// Transition buffer back to UNORDERED_ACCESS
std::swap(barrier.Transition.StateBefore, barrier.Transition.StateAfter);
cl->ResourceBarrier(1, &barrier);
cl.closeAndWait(device, commandQueue);
// create output file
std::unique_ptr<OIIO::ImageOutput> out = OIIO::ImageOutput::create(filename);
if (!out) {
throw std::runtime_error("Could not create file: " + filename);
}
std::vector<oidn::Format> formats = {oidn::Format::Float, oidn::Format::Float2, oidn::Format::Float3, oidn::Format::Float4,
oidn::Format::Half, oidn::Format::Half2, oidn::Format::Half3, oidn::Format::Half4};
auto it = std::find(formats.begin(), formats.end(), format);
if (it == formats.end()) {
throw std::invalid_argument("unsupported format");
}
size_t nchannels = std::distance(formats.begin(), it) + 1;
OIIO::TypeDesc type = OIIO::TypeDesc::FLOAT;
if (nchannels > 4){
nchannels = nchannels - 4;
type = OIIO::TypeDesc::HALF;
}
OIIO::ImageSpec spec(width, height, nchannels, type);
out->open(filename, spec);
void* mapped = nullptr;
auto stagingResource = stagingBuffer.get();
stagingResource->Map(0, nullptr, &mapped);
out->write_image(type, mapped);
D3D12_RANGE readingRange = {0, buffer.get_size()};
stagingResource->Unmap(0, &readingRange);
out->close();
}
Buffer& get_buffer() {
return buffer;
}
uint32_t get_width() {
return width;
}
uint32_t get_height() {
return height;
}
oidn::Format get_format() {
return format;
}
private:
Buffer buffer;
uint32_t width;
uint32_t height;
oidn::Format format;
};
class InteropTestDX {
public:
InteropTestDX() {
#if defined(_DEBUG)
{
if (SUCCEEDED(D3D12GetDebugInterface(IID_PPV_ARGS(&debugController)))) {
debugController->EnableDebugLayer();
}
}
#endif
// Create DXGI factory
if (FAILED(CreateDXGIFactory1(IID_PPV_ARGS(&factory)))) {
throw std::runtime_error("Failed to create DXGI Factory.");
}
// Create D3D12 device
if (FAILED(D3D12CreateDevice(
nullptr, // default adapter
D3D_FEATURE_LEVEL_11_0,
IID_PPV_ARGS(&device)
))) {
throw std::runtime_error("Failed to create D3D12 Device.");
}
// Create command queue
D3D12_COMMAND_QUEUE_DESC queueDesc = {};
queueDesc.Type = D3D12_COMMAND_LIST_TYPE_DIRECT;
queueDesc.Flags = D3D12_COMMAND_QUEUE_FLAG_NONE;
if (FAILED(device->CreateCommandQueue(&queueDesc, IID_PPV_ARGS(&commandQueue)))) {
throw std::runtime_error("Failed to create command queue.");
}
}
~InteropTestDX() = default;
InteropTestDX(const InteropTestDX&) = delete;
InteropTestDX& operator=(const InteropTestDX&) = delete;
InteropTestDX(InteropTestDX&&) = default;
InteropTestDX& operator=(InteropTestDX&&) = default;
void setup_oidn_device() {
// Get the LUID of the adapter
LUID luid = device->GetAdapterLuid();
// Initialize the denoiser device
oidn_device = oidn::newDevice(oidn::LUID{luid.LowPart, luid.HighPart});
if (oidn_device.getError() != oidn::Error::None)
throw std::runtime_error("Failed to create OIDN device.");
oidn_device.commit();
if (oidn_device.getError() != oidn::Error::None)
throw std::runtime_error("Failed to commit OIDN device.");
// Find a compatible external memory handle type
const auto oidn_external_mem_types = oidn_device.get<oidn::ExternalMemoryTypeFlags>("externalMemoryTypes");
if (!(oidn_external_mem_types & oidn::ExternalMemoryTypeFlag::OpaqueWin32))
throw std::runtime_error("failed to find compatible external memory type");
}
ImageBuffer create_buffer(uint32_t width, uint32_t height, oidn::Format format) {
return ImageBuffer(device, width, height, format);
}
ImageBuffer load_exr(std::string& path) {
return ImageBuffer::load_exr(path, device, commandQueue);
}
oidn::BufferRef create_oidn_buffer(ImageBuffer& imageBuffer) {
// Create an OIDN buffer from the D3D12 resource
HANDLE sharedHandle = nullptr;
device->CreateSharedHandle(imageBuffer.get_buffer().get().Get(), nullptr, GENERIC_ALL, nullptr, &sharedHandle);
return oidn_device.newBuffer(oidn::ExternalMemoryTypeFlag::OpaqueWin32, &sharedHandle, nullptr, imageBuffer.get_buffer().get_size());
}
void save_exr(ImageBuffer& buf, const std::string& path) {
return buf.save_exr(path, device, commandQueue);
}
public:
ComPtr<ID3D12Debug> debugController;
ComPtr<IDXGIFactory6> factory;
ComPtr<ID3D12Device> device;
ComPtr<ID3D12CommandQueue> commandQueue;
oidn::DeviceRef oidn_device;
};
File diff suppressed because it is too large Load Diff
+1 -299
View File
@@ -8,9 +8,6 @@
#include <cassert>
#include <cmath>
#include <limits>
#include <atomic>
#include <cstdlib>
#include <string>
#define CATCH_CONFIG_RUNNER
#define CATCH_CONFIG_FAST_COMPILE
@@ -147,120 +144,6 @@ TEST_CASE("device creation", "[device]")
// -------------------------------------------------------------------------------------------------
TEST_CASE("device parameters", "[device_params]")
{
DeviceRef device = makeAndCommitDevice();
SECTION("version")
{
const int version = device.get<int>("version");
const int major = device.get<int>("versionMajor");
const int minor = device.get<int>("versionMinor");
const int patch = device.get<int>("versionPatch");
REQUIRE(device.getError() == Error::None);
REQUIRE(version == OIDN_VERSION);
REQUIRE(version == major * 10000 + minor * 100 + patch);
}
SECTION("type")
{
const DeviceType type = device.get<DeviceType>("type");
REQUIRE(device.getError() == Error::None);
REQUIRE(type != DeviceType::Default);
}
SECTION("capability queries")
{
// These must be queryable without error; the actual values are device-dependent
device.get<bool>("systemMemorySupported");
device.get<bool>("managedMemorySupported");
device.get<int>("externalMemoryTypes");
device.get<int>("externalSemaphoreTypes");
REQUIRE(device.getError() == Error::None);
}
SECTION("verbose round-trip")
{
// Setting verbose is ignored if the OIDN_VERBOSE environment variable is set, so skip then
if (std::getenv("OIDN_VERBOSE") == nullptr)
{
// Use a fresh, uncommitted device to avoid emitting any verbose output
DeviceRef verboseDevice = makeDevice();
verboseDevice.set("verbose", 0);
REQUIRE(verboseDevice.getError() == Error::None);
REQUIRE(verboseDevice.get<int>("verbose") == 0);
}
}
SECTION("invalid parameter name")
{
device.get<int>("qwertyuiop");
REQUIRE(device.getError() == Error::InvalidArgument);
device.get<int>(nullptr);
REQUIRE(device.getError() == Error::InvalidArgument);
// Setting an unknown parameter only triggers a warning, not an error
device.set("qwertyuiop", 1);
REQUIRE(device.getError() == Error::None);
}
}
// -------------------------------------------------------------------------------------------------
struct ErrorRecord
{
std::atomic<int> count;
Error lastError;
std::string lastMessage;
ErrorRecord() : count(0), lastError(Error::None) {}
};
void errorCallback(void* userPtr, Error error, const char* message)
{
ErrorRecord* record = static_cast<ErrorRecord*>(userPtr);
record->count++;
record->lastError = error;
record->lastMessage = message ? message : "";
}
TEST_CASE("error function", "[error_function]")
{
ErrorRecord record;
DeviceRef device = makeAndCommitDevice();
device.setErrorFunction(errorCallback, &record);
SECTION("error is reported via the callback")
{
// Trigger an error by querying an unknown device parameter
device.get<int>("nonexistentParameter");
REQUIRE(record.count == 1);
REQUIRE(record.lastError == Error::InvalidArgument);
REQUIRE(!record.lastMessage.empty());
// The same error is still retrievable via getError, including the message
const char* message = nullptr;
REQUIRE(device.getError(message) == Error::InvalidArgument);
REQUIRE(message != nullptr);
}
SECTION("callback not invoked on success")
{
device.get<int>("nonexistentParameter");
REQUIRE(record.count == 1);
device.getError(); // clear the stored error
// A successful operation must not invoke the error callback
device.get<int>("version");
REQUIRE(record.count == 1);
REQUIRE(device.getError() == Error::None);
}
}
// -------------------------------------------------------------------------------------------------
TEST_CASE("buffer creation", "[buffer]")
{
const size_t bufferSize = 1024*768*3*4; // Metal requires shared buffer size to be multiple of 16K
@@ -334,14 +217,11 @@ TEST_CASE("buffer creation", "[buffer]")
}
#endif
#if !defined(OIDN_SANITIZER)
// Passing an invalid enum value is a false positive for UBSan
SECTION("invalid buffer storage")
{
BufferRef buffer = device.newBuffer(bufferSize, static_cast<Storage>(-42));
REQUIRE(device.getError() == Error::InvalidArgument);
}
#endif
SECTION("device released before buffer")
{
@@ -434,86 +314,6 @@ TEST_CASE("buffer read/write", "[buffer_rw]")
// Verify the data
REQUIRE(memcmp(src.data(), dst.data(), bufferSize) == 0);
}
SECTION("partial offset")
{
// Fill the whole buffer first
buffer.write(0, bufferSize, src.data());
REQUIRE(device.getError() == Error::None);
// Read back a sub-range at a non-zero offset and verify it maps to the right bytes
const int offset = 64;
const int count = 1000;
std::vector<int> partial(count, 0);
buffer.read(offset * sizeof(int), count * sizeof(int), partial.data());
REQUIRE(device.getError() == Error::None);
REQUIRE(memcmp(partial.data(), src.data() + offset, count * sizeof(int)) == 0);
// Overwrite a sub-range and verify only that region changed
std::vector<int> patch(count);
for (int i = 0; i < count; ++i)
patch[i] = -(i + 1);
buffer.write(offset * sizeof(int), count * sizeof(int), patch.data());
REQUIRE(device.getError() == Error::None);
buffer.read(0, bufferSize, dst.data());
REQUIRE(device.getError() == Error::None);
REQUIRE(memcmp(dst.data() + offset, patch.data(), count * sizeof(int)) == 0);
// Bytes before the patched region must be untouched
REQUIRE(memcmp(dst.data(), src.data(), offset * sizeof(int)) == 0);
// Bytes after the patched region must be untouched
REQUIRE(memcmp(dst.data() + offset + count, src.data() + offset + count,
(N - offset - count) * sizeof(int)) == 0);
}
}
// -------------------------------------------------------------------------------------------------
TEST_CASE("buffer properties", "[buffer_props]")
{
const size_t bufferSize = 64 * 1024;
DeviceRef device = makeAndCommitDevice();
const bool managedMemorySupported = device.get<bool>("managedMemorySupported");
REQUIRE(device.getError() == Error::None);
SECTION("default buffer")
{
BufferRef buffer = device.newBuffer(bufferSize);
REQUIRE(device.getError() == Error::None);
REQUIRE(buffer.getSize() == bufferSize);
// A default buffer uses managed storage if supported, otherwise host storage; either way
// it is accessible by the host, so it must expose a valid data pointer
REQUIRE(buffer.getStorage() == (managedMemorySupported ? Storage::Managed : Storage::Host));
REQUIRE(buffer.getData() != nullptr);
}
SECTION("device buffer")
{
BufferRef buffer = device.newBuffer(bufferSize, Storage::Device);
REQUIRE(device.getError() == Error::None);
REQUIRE(buffer.getSize() == bufferSize);
REQUIRE(buffer.getStorage() == Storage::Device);
}
SECTION("managed buffer")
{
if (managedMemorySupported)
{
BufferRef buffer = device.newBuffer(bufferSize, Storage::Managed);
REQUIRE(device.getError() == Error::None);
REQUIRE(buffer.getSize() == bufferSize);
REQUIRE(buffer.getStorage() == Storage::Managed);
}
}
SECTION("empty buffer")
{
BufferRef buffer = device.newBuffer(0);
REQUIRE(device.getError() == Error::None);
REQUIRE(buffer.getSize() == 0);
}
}
// -------------------------------------------------------------------------------------------------
@@ -635,12 +435,9 @@ TEST_CASE("single filter", "[single_filter][minimal]")
setFilterImage(filter, nullptr, input);
REQUIRE(device.getError() == Error::InvalidArgument);
#if !defined(OIDN_SANITIZER)
// Try setting an image with invalid format (passing an invalid enum value is a false positive for
// UBSan)
// Try setting an image with invalid format
filter.setImage("color", input->getBuffer(), static_cast<Format>(-1), W, H);
REQUIRE(device.getError() == Error::InvalidArgument);
#endif
// Try setting an image with buffer overflow
filter.setImage("color", input->getBuffer(), input->getFormat(), W+1, H);
@@ -685,101 +482,6 @@ TEST_CASE("single filter", "[single_filter][minimal]")
// -------------------------------------------------------------------------------------------------
TEST_CASE("filter parameters", "[filter_params]")
{
const int W = 137;
const int H = 97;
DeviceRef device = makeAndCommitDevice();
FilterRef filter = device.newFilter("RT");
REQUIRE(bool(filter));
SECTION("boolean parameters round-trip")
{
filter.set("hdr", true);
REQUIRE(filter.get<bool>("hdr") == true);
filter.set("hdr", false);
REQUIRE(filter.get<bool>("hdr") == false);
filter.set("srgb", true);
REQUIRE(filter.get<bool>("srgb") == true);
filter.set("srgb", false);
REQUIRE(filter.get<bool>("srgb") == false);
filter.set("cleanAux", true);
REQUIRE(filter.get<bool>("cleanAux") == true);
filter.set("cleanAux", false);
REQUIRE(filter.get<bool>("cleanAux") == false);
REQUIRE(device.getError() == Error::None);
}
SECTION("quality parameter round-trip")
{
filter.set("quality", Quality::High);
REQUIRE(filter.get<Quality>("quality") == Quality::High);
filter.set("quality", Quality::Balanced);
REQUIRE(filter.get<Quality>("quality") == Quality::Balanced);
filter.set("quality", Quality::Fast);
REQUIRE(filter.get<Quality>("quality") == Quality::Fast);
REQUIRE(device.getError() == Error::None);
// Default quality must resolve to a concrete quality mode
filter.set("quality", Quality::Default);
REQUIRE(device.getError() == Error::None);
const Quality resolved = filter.get<Quality>("quality");
REQUIRE((resolved == Quality::High || resolved == Quality::Balanced ||
resolved == Quality::Fast));
// An invalid quality mode must be rejected
filter.set("quality", static_cast<Quality>(12345));
REQUIRE(device.getError() == Error::InvalidArgument);
}
SECTION("numeric parameters round-trip")
{
filter.set("maxMemoryMB", 1234);
REQUIRE(filter.get<int>("maxMemoryMB") == 1234);
filter.set("inputScale", 2.5f);
REQUIRE(filter.get<float>("inputScale") == 2.5f);
REQUIRE(device.getError() == Error::None);
}
SECTION("read-only output parameters")
{
auto color = makeConstImage(device, W, H);
auto output = makeImage(device, W, H);
setFilterImage(filter, "color", color);
setFilterImage(filter, "output", output);
filter.set("hdr", true);
filter.commit();
REQUIRE(device.getError() == Error::None);
// These are computed by the filter (device-dependent) and must obey their documented ranges
const int tileAlignment = filter.get<int>("tileAlignment");
const int tileOverlap = filter.get<int>("tileOverlap");
REQUIRE(device.getError() == Error::None);
REQUIRE(tileAlignment >= 1);
REQUIRE(tileOverlap >= 0);
}
SECTION("invalid parameter name")
{
// Getting an unknown parameter is an error
filter.get<int>("nonexistentParameter");
REQUIRE(device.getError() == Error::InvalidArgument);
// Setting an unknown parameter only produces a warning, not an error
filter.set("nonexistentParameter", 1);
REQUIRE(device.getError() == Error::None);
}
}
// -------------------------------------------------------------------------------------------------
void multiFilter1PerDeviceTest(DeviceRef& device, const std::vector<int>& sizes, bool reuseFilter)
{
FilterRef filter;
+1 -1
View File
@@ -25,7 +25,7 @@ target_link_libraries(OpenImageDenoise_utils PUBLIC OpenImageDenoise_common Open
# OpenImageIO
option(OIDN_APPS_OPENIMAGEIO "Enable OpenImageIO support in example and test apps." OFF)
if(OIDN_APPS_OPENIMAGEIO)
find_package(OpenImageIO CONFIG REQUIRED)
find_package(OpenImageIO REQUIRED)
if(OPENIMAGEIO_VERSION VERSION_GREATER_EQUAL "3")
# OpenImageIO 3.x requires C++17
set_target_properties(OpenImageDenoise_utils PROPERTIES CXX_STANDARD 17)
+2 -6
View File
@@ -13,10 +13,10 @@ OIDN_NAMESPACE_BEGIN
width(0),
height(0),
numChannels(0),
dataType(DataType::Undefined),
dataType(DataType::Void),
format(Format::Undefined) {}
ImageBuffer::ImageBuffer(const DeviceRef& device, size_t width, size_t height, size_t numChannels,
ImageBuffer::ImageBuffer(const DeviceRef& device, int width, int height, int numChannels,
DataType dataType, Storage storage, bool forceHostCopy)
: device(device),
numValues(size_t(width) * height * numChannels),
@@ -26,10 +26,6 @@ OIDN_NAMESPACE_BEGIN
dataType(dataType),
format(makeFormat(dataType, numChannels))
{
if (width > maxDim || height > maxDim || width * height * getC() > std::numeric_limits<int>::max())
throw std::runtime_error("image size is too large");
const size_t valueByteSize = getDataTypeSize(dataType);
byteSize = std::max(numValues * valueByteSize, size_t(1)); // avoid zero-sized buffer
buffer = device.newBuffer(byteSize, storage);
+8 -10
View File
@@ -15,19 +15,17 @@ OIDN_NAMESPACE_BEGIN
class ImageBuffer
{
public:
static constexpr size_t maxDim = 65536;
ImageBuffer();
ImageBuffer(const DeviceRef& device, size_t width, size_t height, size_t numChannels,
ImageBuffer(const DeviceRef& device, int width, int height, int numChannels,
DataType dataType = DataType::Float32,
Storage storage = Storage::Undefined,
bool forceHostCopy = false);
~ImageBuffer();
oidn_inline size_t getW() const { return width; }
oidn_inline size_t getH() const { return height; }
oidn_inline size_t getC() const { return numChannels; }
std::array<size_t, 3> getDims() const { return {width, height, numChannels}; }
oidn_inline int getW() const { return width; }
oidn_inline int getH() const { return height; }
oidn_inline int getC() const { return numChannels; }
std::array<int, 3> getDims() const { return {width, height, numChannels}; }
oidn_inline DataType getDataType() const { return dataType; }
oidn_inline Format getFormat() const { return format; }
@@ -100,9 +98,9 @@ OIDN_NAMESPACE_BEGIN
char* hostPtr;
size_t byteSize;
size_t numValues;
size_t width;
size_t height;
size_t numChannels;
int width;
int height;
int numChannels;
DataType dataType;
Format format;
};
+14 -12
View File
@@ -60,7 +60,7 @@ OIDN_NAMESPACE_BEGIN
// Read the header
std::string id;
file >> id;
size_t C;
int C;
if (id == "PF")
C = 3;
else if (id == "Pf")
@@ -70,10 +70,10 @@ OIDN_NAMESPACE_BEGIN
else
throw std::runtime_error("invalid PFM image");
if (dataType == DataType::Undefined)
if (dataType == DataType::Void)
dataType = DataType::Float32;
size_t H, W;
int H, W;
file >> W >> H;
float scale;
@@ -91,21 +91,23 @@ OIDN_NAMESPACE_BEGIN
// Read the pixels
auto image = std::make_shared<ImageBuffer>(device, W, H, C, dataType, storage);
for (size_t h = 0; h < H; ++h)
for (int h = 0; h < H; ++h)
{
for (size_t w = 0; w < W; ++w)
for (int w = 0; w < W; ++w)
{
for (size_t c = 0; c < C; ++c)
for (int c = 0; c < C; ++c)
{
float x;
file.read(reinterpret_cast<char*>(&x), sizeof(float));
if (file.fail())
throw std::runtime_error("invalid PFM image: error reading pixel data");
image->set((size_t(H-1-h)*W + w) * C + c, x * scale);
file.read((char*)&x, sizeof(float));
if (c < C)
image->set((size_t(H-1-h)*W + w) * C + c, x * scale);
}
}
}
if (file.fail())
throw std::runtime_error("invalid PFM image");
return image;
}
@@ -172,7 +174,7 @@ OIDN_NAMESPACE_BEGIN
else
throw std::runtime_error("invalid PHM image");
if (dataType == DataType::Undefined)
if (dataType == DataType::Void)
dataType = DataType::Float16;
int H, W;
@@ -309,7 +311,7 @@ OIDN_NAMESPACE_BEGIN
const OIIO::ImageSpec& spec = in->spec();
const int numChannels = std::min(spec.nchannels, 3);
if (dataType == DataType::Undefined)
if (dataType == DataType::Void)
dataType = (spec.channelformat(0) == OIIO::TypeDesc::HALF) ? DataType::Float16 : DataType::Float32;
auto image = std::make_shared<ImageBuffer>(device, spec.width, spec.height, numChannels,
+2 -2
View File
@@ -10,14 +10,14 @@ OIDN_NAMESPACE_BEGIN
// Loads an image with optionally specified number of channels and data type
std::shared_ptr<ImageBuffer> loadImage(const DeviceRef& device,
const std::string& filename,
DataType dataType = DataType::Undefined,
DataType dataType = DataType::Void,
Storage storage = Storage::Undefined);
// Loads an image with/without sRGB to linear conversion
std::shared_ptr<ImageBuffer> loadImage(const DeviceRef& device,
const std::string& filename,
bool srgb,
DataType dataType = DataType::Undefined,
DataType dataType = DataType::Void,
Storage storage = Storage::Undefined);
// Saves an image
+130
View File
@@ -0,0 +1,130 @@
###########################################################################
# OpenImageIO https://www.openimageio.org
# Copyright 2008-present Contributors to the OpenImageIO project.
# SPDX-License-Identifier: BSD-3-Clause
# https://github.com/OpenImageIO/oiio/blob/master/LICENSE.md
#
# For an up-to-date version of this file, see:
# https://github.com/OpenImageIO/oiio/blob/master/src/cmake/Modules/FindOpenImageIO.cmake
#
###########################################################################
#
# CMake module to find OpenImageIO
#
# This module will set
# OpenImageIO_FOUND True, if found
# OPENIMAGEIO_INCLUDES directory where headers are found
# OPENIMAGEIO_LIBRARIES libraries for OIIO
# OPENIMAGEIO_LIBRARY_DIRS library dirs for OIIO
# OPENIMAGEIO_VERSION Version ("major.minor.patch.tweak")
# OPENIMAGEIO_VERSION_MAJOR Version major number
# OPENIMAGEIO_VERSION_MINOR Version minor number
# OPENIMAGEIO_VERSION_PATCH Version minor patch
# OPENIMAGEIO_VERSION_TWEAK Version minor tweak
#
# Imported targets:
# OpenImageIO::OpenImageIO The libOpenImageIO library.
# OpenImageIO::oiiotool The oiiotool executable.
#
# Special inputs:
# OpenImageIO_ROOT - if using CMake >= 3.12, will automatically search
# this area for OIIO components.
# OPENIMAGEIO_ROOT_DIR - custom "prefix" location of OIIO installation
# (expecting bin, lib, include subdirectories)
# This is deprecated, but will work for a while.
# OpenImageIO_FIND_QUIETLY - if set, print minimal console output
# OIIO_LIBNAME_SUFFIX - if set, optional nonstandard library suffix
#
###########################################################################
#
# NOTE: This file is deprecated.
#
# In OIIO 2.1+, we generate OpenImageIOConfig.cmake files that are now the
# preferred way for downstream projects to find an installed OIIO. There
# should be no need to copy this FindOpenImageIO.cmake file into downstream
# projects, *unless* they need to work with a range of OIIO versions that
# may include <2.1, which would lack the generated config files.
#
###########################################################################
# If 'OPENIMAGE_HOME' not set, use the env variable of that name if available
if (NOT OPENIMAGEIO_ROOT_DIR AND NOT $ENV{OPENIMAGEIO_ROOT_DIR} STREQUAL "")
set (OPENIMAGEIO_ROOT_DIR $ENV{OPENIMAGEIO_ROOT_DIR})
endif ()
find_library ( OPENIMAGEIO_LIBRARY
NAMES OpenImageIO${OIIO_LIBNAME_SUFFIX}
HINTS ${OPENIMAGEIO_ROOT_DIR}
PATH_SUFFIXES lib64 lib )
find_library ( OPENIMAGEIO_UTIL_LIBRARY
NAMES OpenImageIO_Util${OIIO_LIBNAME_SUFFIX}
HINTS ${OPENIMAGEIO_ROOT_DIR}
PATH_SUFFIXES lib64 lib )
find_path ( OPENIMAGEIO_INCLUDE_DIR
NAMES OpenImageIO/imageio.h
HINTS ${OPENIMAGEIO_ROOT_DIR} )
# Try to figure out version number
set (OIIO_VERSION_HEADER "${OPENIMAGEIO_INCLUDE_DIR}/OpenImageIO/oiioversion.h")
if (EXISTS "${OIIO_VERSION_HEADER}")
file (STRINGS "${OIIO_VERSION_HEADER}" TMP REGEX "^#define OIIO_VERSION_MAJOR .*$")
string (REGEX MATCHALL "[0-9]+" OPENIMAGEIO_VERSION_MAJOR ${TMP})
file (STRINGS "${OIIO_VERSION_HEADER}" TMP REGEX "^#define OIIO_VERSION_MINOR .*$")
string (REGEX MATCHALL "[0-9]+" OPENIMAGEIO_VERSION_MINOR ${TMP})
file (STRINGS "${OIIO_VERSION_HEADER}" TMP REGEX "^#define OIIO_VERSION_PATCH .*$")
string (REGEX MATCHALL "[0-9]+" OPENIMAGEIO_VERSION_PATCH ${TMP})
file (STRINGS "${OIIO_VERSION_HEADER}" TMP REGEX "^#define OIIO_VERSION_TWEAK .*$")
if (TMP)
string (REGEX MATCHALL "[0-9]+" OPENIMAGEIO_VERSION_TWEAK ${TMP})
else ()
set (OPENIMAGEIO_VERSION_TWEAK 0)
endif ()
set (OPENIMAGEIO_VERSION "${OPENIMAGEIO_VERSION_MAJOR}.${OPENIMAGEIO_VERSION_MINOR}.${OPENIMAGEIO_VERSION_PATCH}.${OPENIMAGEIO_VERSION_TWEAK}")
endif ()
include (FindPackageHandleStandardArgs)
find_package_handle_standard_args (OpenImageIO
FOUND_VAR OpenImageIO_FOUND
REQUIRED_VARS OPENIMAGEIO_INCLUDE_DIR OPENIMAGEIO_LIBRARY
OPENIMAGEIO_VERSION
VERSION_VAR OPENIMAGEIO_VERSION
)
set (OPENIMAGEIO_FOUND ${OpenImageIO_FOUND}) # Old name
if (OpenImageIO_FOUND)
set (OPENIMAGEIO_INCLUDES ${OPENIMAGEIO_INCLUDE_DIR})
set (OPENIMAGEIO_LIBRARIES ${OPENIMAGEIO_LIBRARY})
get_filename_component (OPENIMAGEIO_LIBRARY_DIRS "${OPENIMAGEIO_LIBRARY}" DIRECTORY)
if (NOT OpenImageIO_FIND_QUIETLY)
message ( STATUS "OpenImageIO includes = ${OPENIMAGEIO_INCLUDE_DIR}" )
message ( STATUS "OpenImageIO libraries = ${OPENIMAGEIO_LIBRARIES}" )
message ( STATUS "OpenImageIO library_dirs = ${OPENIMAGEIO_LIBRARY_DIRS}" )
endif ()
if (NOT TARGET OpenImageIO::OpenImageIO)
add_library(OpenImageIO::OpenImageIO UNKNOWN IMPORTED)
set_target_properties(OpenImageIO::OpenImageIO PROPERTIES
INTERFACE_INCLUDE_DIRECTORIES "${OPENIMAGEIO_INCLUDES}")
set_property(TARGET OpenImageIO::OpenImageIO APPEND PROPERTY
IMPORTED_LOCATION "${OPENIMAGEIO_LIBRARIES}")
endif ()
# Starting with OIIO v2.3, some utility classes are now only declared in OpenImageIO_Util
# (and not in both libraries like in older versions).
if (${OPENIMAGEIO_VERSION} VERSION_GREATER_EQUAL "2.3" AND NOT TARGET OpenImageIO::OpenImageIO_Util)
add_library(OpenImageIO::OpenImageIO_Util UNKNOWN IMPORTED)
set_target_properties(OpenImageIO::OpenImageIO_Util PROPERTIES
IMPORTED_LOCATION "${OPENIMAGEIO_UTIL_LIBRARY}")
target_link_libraries(OpenImageIO::OpenImageIO INTERFACE OpenImageIO::OpenImageIO_Util)
endif ()
endif ()
mark_as_advanced (
OPENIMAGEIO_INCLUDE_DIR
OPENIMAGEIO_LIBRARY
OPENIMAGEIO_UTIL_LIBRARY
)
+166
View File
@@ -0,0 +1,166 @@
## Copyright 2018 Intel Corporation
## SPDX-License-Identifier: Apache-2.0
set(DNNL_SOURCE_DIR "${PROJECT_SOURCE_DIR}/external/mkl-dnn")
set(DNNL_BINARY_DIR "${PROJECT_BINARY_DIR}/external/mkl-dnn")
set(DNNL_VERSION_MAJOR 3)
set(DNNL_VERSION_MINOR 0)
set(DNNL_VERSION_PATCH 1)
set(DNNL_VERSION_HASH "N/A")
set(DNNL_CPU_RUNTIME "TBB")
set(DNNL_CPU_THREADING_RUNTIME "TBB")
set(DNNL_GPU_RUNTIME "NONE")
option(DNNL_ENABLE_JIT_PROFILING
"Enable registration of oneDNN kernels that are generated at runtime with
VTune Amplifier. Without the registrations, VTune Amplifier would report
data collected inside the kernels as `outside any known module`."
OFF)
mark_as_advanced(DNNL_ENABLE_JIT_PROFILING)
option(DNNL_ENABLE_ITT_TASKS
"Enable ITT Tasks tagging feature and tag all primitive execution. VTune
Amplifier can group profiling results based on those ITT tasks and show
corresponding timeline information."
OFF)
mark_as_advanced(DNNL_ENABLE_ITT_TASKS)
set(BUILD_INFERENCE TRUE)
set(BUILD_CONVOLUTION TRUE)
set(BUILD_PRIMITIVE_CPU_ISA_ALL TRUE)
configure_file(
"${DNNL_SOURCE_DIR}/include/oneapi/dnnl/dnnl_config.h.in"
"${DNNL_BINARY_DIR}/include/oneapi/dnnl/dnnl_config.h"
)
configure_file(
"${DNNL_SOURCE_DIR}/include/oneapi/dnnl/dnnl_version.h.in"
"${DNNL_BINARY_DIR}/include/oneapi/dnnl/dnnl_version.h"
)
file(GLOB DNNL_SOURCES
${DNNL_SOURCE_DIR}/src/common/*.[ch]pp
${DNNL_SOURCE_DIR}/src/cpu/jit_utils/*.[ch]pp
${DNNL_SOURCE_DIR}/src/cpu/x64/injectors/*.[ch]pp
${DNNL_SOURCE_DIR}/src/cpu/x64/xbyak/*.h
)
list(APPEND DNNL_SOURCES
${DNNL_SOURCE_DIR}/src/cpu/bfloat16.cpp
${DNNL_SOURCE_DIR}/src/cpu/binary_injector_utils.cpp
${DNNL_SOURCE_DIR}/src/cpu/binary_injector_utils.hpp
${DNNL_SOURCE_DIR}/src/cpu/cpu_concat.cpp
${DNNL_SOURCE_DIR}/src/cpu/cpu_convolution_list.cpp
${DNNL_SOURCE_DIR}/src/cpu/cpu_convolution_pd.hpp
${DNNL_SOURCE_DIR}/src/cpu/cpu_eltwise_pd.hpp
${DNNL_SOURCE_DIR}/src/cpu/cpu_engine.cpp
${DNNL_SOURCE_DIR}/src/cpu/cpu_engine.hpp
${DNNL_SOURCE_DIR}/src/cpu/cpu_memory_storage.hpp
${DNNL_SOURCE_DIR}/src/cpu/cpu_stream.hpp
${DNNL_SOURCE_DIR}/src/cpu/cpu_sum.cpp
${DNNL_SOURCE_DIR}/src/cpu/float16.cpp
${DNNL_SOURCE_DIR}/src/cpu/platform.cpp
${DNNL_SOURCE_DIR}/src/cpu/platform.hpp
${DNNL_SOURCE_DIR}/src/cpu/primitive_attr_postops.cpp
${DNNL_SOURCE_DIR}/src/cpu/primitive_attr_postops.hpp
${DNNL_SOURCE_DIR}/src/cpu/scale_utils.cpp
${DNNL_SOURCE_DIR}/src/cpu/scale_utils.hpp
${DNNL_SOURCE_DIR}/src/cpu/simple_q10n.hpp
${DNNL_SOURCE_DIR}/src/cpu/zero_point_utils.cpp
${DNNL_SOURCE_DIR}/src/cpu/zero_point_utils.hpp
${DNNL_SOURCE_DIR}/src/cpu/x64/amx_tile_configure.cpp
${DNNL_SOURCE_DIR}/src/cpu/x64/amx_tile_configure.hpp
${DNNL_SOURCE_DIR}/src/cpu/x64/cpu_barrier.hpp
${DNNL_SOURCE_DIR}/src/cpu/x64/cpu_barrier.cpp
${DNNL_SOURCE_DIR}/src/cpu/x64/cpu_isa_traits.cpp
${DNNL_SOURCE_DIR}/src/cpu/x64/cpu_isa_traits.hpp
${DNNL_SOURCE_DIR}/src/cpu/x64/cpu_reducer.cpp
${DNNL_SOURCE_DIR}/src/cpu/x64/cpu_reducer.hpp
${DNNL_SOURCE_DIR}/src/cpu/x64/jit_avx2_conv_kernel_f32.cpp
${DNNL_SOURCE_DIR}/src/cpu/x64/jit_avx2_conv_kernel_f32.hpp
${DNNL_SOURCE_DIR}/src/cpu/x64/jit_avx2_convolution.cpp
${DNNL_SOURCE_DIR}/src/cpu/x64/jit_avx2_convolution.hpp
${DNNL_SOURCE_DIR}/src/cpu/x64/jit_avx512_common_convolution.cpp
${DNNL_SOURCE_DIR}/src/cpu/x64/jit_avx512_common_convolution.hpp
${DNNL_SOURCE_DIR}/src/cpu/x64/jit_avx512_common_conv_kernel.cpp
${DNNL_SOURCE_DIR}/src/cpu/x64/jit_avx512_common_conv_kernel.hpp
${DNNL_SOURCE_DIR}/src/cpu/x64/jit_generator.hpp
${DNNL_SOURCE_DIR}/src/cpu/x64/jit_primitive_conf.hpp
${DNNL_SOURCE_DIR}/src/cpu/x64/jit_sse41_conv_kernel_f32.cpp
${DNNL_SOURCE_DIR}/src/cpu/x64/jit_sse41_conv_kernel_f32.hpp
${DNNL_SOURCE_DIR}/src/cpu/x64/jit_sse41_convolution.cpp
${DNNL_SOURCE_DIR}/src/cpu/x64/jit_sse41_convolution.hpp
${DNNL_SOURCE_DIR}/src/cpu/x64/jit_transpose_utils.cpp
${DNNL_SOURCE_DIR}/src/cpu/x64/jit_transpose_utils.hpp
)
if(DNNL_ENABLE_JIT_PROFILING OR DNNL_ENABLE_ITT_TASKS)
file(GLOB ITT_SOURCES ${DNNL_SOURCE_DIR}/src/common/ittnotify/*.[ch])
list(APPEND DNNL_SOURCES ${ITT_SOURCES})
endif()
if(DNNL_ENABLE_JIT_PROFILING)
list(APPEND DNNL_SOURCES
${DNNL_SOURCE_DIR}/src/cpu/jit_utils/linux_perf/linux_perf.cpp
${DNNL_SOURCE_DIR}/src/cpu/jit_utils/linux_perf/linux_perf.hpp
)
endif()
if(MSVC)
if(CMAKE_CXX_COMPILER_ID STREQUAL "MSVC" OR CMAKE_CXX_COMPILER_ID STREQUAL "Intel")
set_source_files_properties(${DNNL_SOURCES} PROPERTIES COMPILE_FLAGS "/bigobj")
endif()
endif()
add_library(dnnl STATIC ${DNNL_SOURCES})
target_include_directories(dnnl
PUBLIC
$<BUILD_INTERFACE:${DNNL_SOURCE_DIR}/include>
$<BUILD_INTERFACE:${DNNL_BINARY_DIR}/include>
$<BUILD_INTERFACE:${DNNL_SOURCE_DIR}/src>
$<BUILD_INTERFACE:${DNNL_SOURCE_DIR}/src/common>
$<BUILD_INTERFACE:${DNNL_SOURCE_DIR}/src/cpu>
$<BUILD_INTERFACE:${DNNL_SOURCE_DIR}/src/cpu/xbyak>
)
target_compile_definitions(dnnl
PUBLIC
-DDNNL_ENABLE_CONCURRENT_EXEC
)
set(DNNL_COMPILE_OPTIONS "")
if(WIN32 AND CMAKE_CXX_COMPILER_ID STREQUAL "Intel")
# Correct 'jnl' macro/jit issue
list(APPEND DNNL_COMPILE_OPTIONS "/Qlong-double")
endif()
target_compile_options(dnnl PRIVATE ${DNNL_COMPILE_OPTIONS})
target_link_libraries(dnnl
PUBLIC
${CMAKE_THREAD_LIBS_INIT}
TBB::tbb
)
if(DNNL_ENABLE_JIT_PROFILING OR DNNL_ENABLE_ITT_TASKS)
if(UNIX AND NOT APPLE)
# Not every compiler adds -ldl automatically
target_link_libraries(dnnl PUBLIC ${CMAKE_DL_LIBS})
endif()
endif()
if(DNNL_ENABLE_ITT_TASKS)
target_compile_definitions(dnnl PUBLIC -DDNNL_ENABLE_ITT_TASKS)
endif()
if(NOT DNNL_ENABLE_JIT_PROFILING)
target_compile_definitions(dnnl PUBLIC -DDNNL_ENABLE_JIT_PROFILING=0)
endif()
if(OIDN_STATIC_LIB)
install(TARGETS dnnl
EXPORT
OpenImageDenoise_Exports
ARCHIVE
DESTINATION ${CMAKE_INSTALL_LIBDIR} COMPONENT devel
)
endif()
+13 -10
View File
@@ -2,7 +2,10 @@
## SPDX-License-Identifier: Apache-2.0
# ISPC versions to look for, in descending order (newest first)
set(ISPC_VERSION_WORKING "1.30.0")
set(ISPC_VERSION_WORKING "1.26.0" "1.25.3" "1.25.2" "1.24.0" "1.23.0" "1.22.0" "1.21.1" "1.21.0")
if(WIN32 AND OIDN_ARCH STREQUAL "ARM64")
list(REMOVE_ITEM ISPC_VERSION_WORKING "1.22.0")
endif()
list(GET ISPC_VERSION_WORKING -1 ISPC_VERSION_REQUIRED)
list(GET ISPC_VERSION_WORKING 0 ISPC_VERSION_LATEST)
@@ -59,6 +62,10 @@ if(NOT ISPC_VERSION)
mark_as_advanced(ISPC_EXECUTABLE)
endif()
if(WIN32 AND OIDN_ARCH STREQUAL "ARM64" AND ${ISPC_VERSION} VERSION_EQUAL "1.22.0")
message(FATAL_ERROR "ISPC v1.22.0 is not supported on Windows ARM64. Please use v${ISPC_VERSION_LATEST}.")
endif()
get_filename_component(ISPC_DIR ${ISPC_EXECUTABLE} PATH)
## -----------------------------------------------------------------------------
@@ -84,8 +91,8 @@ endmacro()
## -----------------------------------------------------------------------------
macro(ispc_compile)
set(ISPC_ADDITIONAL_ARGS ${ISPC_ADDITIONAL_ARGS})
set(ISPC_TARGETS ${ISPC_TARGET_LIST})
set(ISPC_ADDITIONAL_ARGS "")
set(ISPC_TARGETS ${OIDN_ISPC_TARGET_LIST})
set(ISPC_TARGET_EXT ${CMAKE_CXX_OUTPUT_EXTENSION})
string(REPLACE ";" "," ISPC_TARGET_ARGS "${ISPC_TARGETS}")
@@ -165,7 +172,6 @@ macro(ispc_compile)
string(REPLACE "-i32x8" "" target ${target}) # strip (sse4|avx|avx2)-i32x8
string(REPLACE "-i32x16" "" target ${target}) # strip avx512(knl|skx)-i32x16
string(REPLACE "-x16" "" target ${target}) # strip avx512spr-x16
string(REPLACE "-x32" "" target ${target}) # strip avx512spr-x32
set(results ${results} "${outdir}/${fname}.dev_${target}${ISPC_TARGET_EXT}")
endforeach()
endif()
@@ -181,7 +187,7 @@ macro(ispc_compile)
-I ${CMAKE_CURRENT_SOURCE_DIR}
${ISPC_INCLUDE_DIR_PARMS}
--arch=${ISPC_ARCHITECTURE}
--addressing=${ISPC_ADDRESSING}
--addressing=${OIDN_ISPC_ADDRESSING}
${ISPC_OPT_FLAGS}
--target=${ISPC_TARGET_ARGS}
--woff
@@ -204,15 +210,12 @@ endmacro()
## -----------------------------------------------------------------------------
function(ispc_target_add_sources target)
cmake_parse_arguments(PARSE_ARGV 1 ISPC_ARG "" "" "TARGET_LIST;COMPILE_OPTIONS")
set(ISPC_TARGET_LIST ${ISPC_TARGET_LIST} ${ISPC_ARG_TARGET_LIST})
set(ISPC_ADDITIONAL_ARGS ${ISPC_ADDITIONAL_ARGS} ${ISPC_ARG_COMPILE_OPTIONS})
# Split-out C/C++ from ISPC files
set(ISPC_SOURCES "")
set(OTHER_SOURCES "")
foreach(src ${ISPC_ARG_UNPARSED_ARGUMENTS})
foreach(src ${ARGN})
get_filename_component(ext ${src} EXT)
if(ext STREQUAL ".ispc")
set(ISPC_SOURCES ${ISPC_SOURCES} ${src})
+1 -1
View File
@@ -80,7 +80,7 @@ else()
if(APPLE)
# macOS specific settings
set(CPACK_GENERATOR "TGZ;ZIP") # ZIP for notarization only
set(CPACK_GENERATOR TGZ)
set(CPACK_PACKAGE_FILE_NAME "${CPACK_PACKAGE_FILE_NAME}.${ARCH}.macos")
set(CPACK_MONOLITHIC_INSTALL 1)
else()
+2 -2
View File
@@ -2,8 +2,8 @@
## SPDX-License-Identifier: Apache-2.0
set(OIDN_VERSION_MAJOR 2)
set(OIDN_VERSION_MINOR 5)
set(OIDN_VERSION_PATCH 0)
set(OIDN_VERSION_MINOR 3)
set(OIDN_VERSION_PATCH 3)
set(OIDN_VERSION_NOTE "")
set(OIDN_VERSION ${OIDN_VERSION_MAJOR}.${OIDN_VERSION_MINOR}.${OIDN_VERSION_PATCH})
+2 -2
View File
@@ -22,7 +22,7 @@ OIDN_NAMESPACE_BEGIN
switch (format)
{
case Format::Undefined:
return DataType::Undefined;
return DataType::Void;
case Format::Float:
case Format::Float2:
case Format::Float3:
@@ -40,7 +40,7 @@ OIDN_NAMESPACE_BEGIN
Format makeFormat(DataType dataType, int numChannels)
{
if (dataType == DataType::Undefined)
if (dataType == DataType::Void)
return Format::Undefined;
Format baseFormat;
+2 -2
View File
@@ -18,8 +18,8 @@ OIDN_NAMESPACE_BEGIN
template<typename T>
struct DataTypeOf;
template<> struct DataTypeOf<void> { static constexpr DataType value = DataType::Undefined; };
template<> struct DataTypeOf<uint8_t> { static constexpr DataType value = DataType::UInt8; };
template<> struct DataTypeOf<void> { static constexpr DataType value = DataType::Void; };
template<> struct DataTypeOf<uint8_t> { static constexpr DataType value = DataType::UInt8; };
template<> struct DataTypeOf<half> { static constexpr DataType value = DataType::Float16; };
template<> struct DataTypeOf<float> { static constexpr DataType value = DataType::Float32; };
+1 -1
View File
@@ -24,7 +24,7 @@ FILESUBTYPE 0x0L
VALUE "FileDescription", "Intel Open Image Denoise Library"
VALUE "FileVersion", OIDN_VERSION_STRING
VALUE "ProductVersion", OIDN_VERSION_STRING
VALUE "LegalCopyright", "Copyright 2009-2026 Intel Corporation"
VALUE "LegalCopyright", "Copyright 2009-2025 Intel Corporation"
VALUE "InternalName", "Intel Open Image Denoise"
VALUE "ProductName", "Intel(R) Open Image Denoise"
}
+5 -5
View File
@@ -47,11 +47,11 @@ OIDN_NAMESPACE_BEGIN
{
switch (dataType)
{
case DataType::Undefined: sm << "x"; break;
case DataType::UInt8: sm << "u8"; break;
case DataType::Float16: sm << "f16"; break;
case DataType::Float32: sm << "f32"; break;
default: sm << "?"; break;
case DataType::Void: sm << "v"; break;
case DataType::UInt8: sm << "u8"; break;
case DataType::Float16: sm << "f16"; break;
case DataType::Float32: sm << "f32"; break;
default: sm << "?"; break;
}
return sm;
+1 -5
View File
@@ -164,8 +164,6 @@ OIDN_NAMESPACE_BEGIN
namespace esimdx = sycl::ext::intel::experimental::esimd;
using sycl::half;
using namespace esimd;
using namespace esimdx;
#endif
#if !defined(OIDN_COMPILE_METAL_DEVICE)
@@ -242,7 +240,7 @@ OIDN_NAMESPACE_BEGIN
// Data types sorted by precision in ascending order
enum class DataType
{
Undefined,
Void,
UInt8,
Float16,
Float32,
@@ -267,8 +265,6 @@ OIDN_NAMESPACE_BEGIN
void* alignedMalloc(size_t size, size_t alignment = memoryAlignment);
void alignedFree(void* ptr);
static_assert(sizeof(size_t) == 8, "size_t is not 64-bit!");
// -----------------------------------------------------------------------------------------------
// String functions
// -----------------------------------------------------------------------------------------------
+4 -2
View File
@@ -15,6 +15,10 @@ set(OIDN_CORE_SOURCES
color.cpp
concat_conv.h
concat_conv.cpp
concat_conv_chw.h
concat_conv_chw.cpp
concat_conv_hwc.h
concat_conv_hwc.cpp
context.h
context.cpp
conv.cpp
@@ -57,13 +61,11 @@ set(OIDN_CORE_SOURCES
rt_filter.cpp
rtlightmap_filter.h
rtlightmap_filter.cpp
semaphore.h
subdevice.h
subdevice.cpp
tensor.h
tensor.cpp
tensor_accessor.h
tensor_desc.h
tensor_layout.h
tensor_reorder.h
tensor_reorder.cpp
+1 -1
View File
@@ -10,7 +10,7 @@ OIDN_NAMESPACE_BEGIN
TransferFunction::TransferFunction(Type type)
: type(type)
{
const float xMax = math::reduce_max(forward(vec3f(yMax)));
const float xMax = math::reduce_max(forward(yMax));
normScale = 1./xMax;
rcpNormScale = xMax;
}
+8 -54
View File
@@ -50,26 +50,6 @@ OIDN_NAMESPACE_BEGIN
else
return math::pow((x - d) / b, 1.f/c);
}
#if defined(OIDN_COMPILE_SYCL)
template<int N>
static oidn_inline simd<float, N> forward(simd<float, N> y)
{
simd<float, N> v0 = a * y;
simd<float, N> v = b * pow(y, c) + d;
v.merge(v0, y <= y0);
return v;
}
template<int N>
static oidn_inline simd<float, N> inverse(simd<float, N> x)
{
simd<float, N> v0 = x / a;
simd<float, N> v = pow((x - d) / b, 1.f/c);
v.merge(v0, x <= x0);
return v;
}
#endif
};
struct PU
@@ -105,30 +85,6 @@ OIDN_NAMESPACE_BEGIN
else
return math::exp((x - g) / e) - f;
}
#if defined(OIDN_COMPILE_SYCL)
template<int N>
static oidn_inline simd<float, N> forward(simd<float, N> y)
{
simd<float, N> v0 = a * y;
simd<float, N> v1 = b * pow(y, c) + d;
simd<float, N> v = e * log(y + f) + g;
v.merge(v1, y <= y1);
v.merge(v0, y <= y0);
return v;
}
template<int N>
static oidn_inline simd<float, N> inverse(simd<float, N> x)
{
simd<float, N> v0 = x / a;
simd<float, N> v1 = pow((x - d) / b, 1.f/c);
simd<float, N> v = exp((x - g) / e) - f;
v.merge(v1, x <= x1);
v.merge(v0, x <= x0);
return v;
}
#endif
};
#if !defined(OIDN_COMPILE_METAL_DEVICE)
@@ -166,8 +122,7 @@ OIDN_NAMESPACE_BEGIN
return outputScale;
}
template<typename T>
oidn_host_device_inline vec3<T> forward(vec3<T> y) const
oidn_host_device_inline vec3f forward(vec3f y) const
{
switch (type)
{
@@ -175,21 +130,20 @@ OIDN_NAMESPACE_BEGIN
return y;
case Type::SRGB:
return vec3<T>(SRGB::forward(y.x), SRGB::forward(y.y), SRGB::forward(y.z));
return vec3f(SRGB::forward(y.x), SRGB::forward(y.y), SRGB::forward(y.z));
case Type::PU:
return vec3<T>(PU::forward(y.x), PU::forward(y.y), PU::forward(y.z)) * normScale;
return vec3f(PU::forward(y.x), PU::forward(y.y), PU::forward(y.z)) * normScale;
case Type::Log:
return math::log(y + 1.f) * normScale;
default:
return T(0.f);
return 0;
}
}
template<typename T>
oidn_host_device_inline vec3<T> inverse(vec3<T> x) const
oidn_host_device_inline vec3f inverse(vec3f x) const
{
switch (type)
{
@@ -197,16 +151,16 @@ OIDN_NAMESPACE_BEGIN
return x;
case Type::SRGB:
return vec3<T>(SRGB::inverse(x.x), SRGB::inverse(x.y), SRGB::inverse(x.z));
return vec3f(SRGB::inverse(x.x), SRGB::inverse(x.y), SRGB::inverse(x.z));
case Type::PU:
return vec3<T>(PU::inverse(x.x * rcpNormScale), PU::inverse(x.y * rcpNormScale), PU::inverse(x.z * rcpNormScale));
return vec3f(PU::inverse(x.x * rcpNormScale), PU::inverse(x.y * rcpNormScale), PU::inverse(x.z * rcpNormScale));
case Type::Log:
return math::exp(x * rcpNormScale) - 1.f;
default:
return T(0.f);
return 0;
}
}
};
+14 -171
View File
@@ -5,35 +5,18 @@
OIDN_NAMESPACE_BEGIN
ConcatConvBase::ConcatConvBase(const ConcatConvDesc& desc)
ConcatConv::ConcatConv(const ConcatConvDesc& desc)
: ConcatConvDesc(desc)
{
if (src0Desc.getRank() != 3 ||
src1Desc.getRank() != 3 ||
src0Desc.layout != src1Desc.layout ||
src0Desc.dataType != src1Desc.dataType)
if (src1Desc.getRank() != 3 ||
src2Desc.getRank() != 3 ||
src1Desc.getH() != src2Desc.getH() ||
src1Desc.getW() != src2Desc.getW() ||
src1Desc.layout != src2Desc.layout ||
src1Desc.dataType != src2Desc.dataType)
throw std::invalid_argument("invalid concat+conv source descriptor");
switch (fusion)
{
case Fusion::None:
if (src0Desc.getH() != src1Desc.getH() ||
src0Desc.getW() != src1Desc.getW())
throw std::invalid_argument("invalid concat+conv source descriptor");
break;
case Fusion::UpsampleSrc0: // src0 is upsampled by 2x
if (src0Desc.getH() * 2 != src1Desc.getH() ||
src0Desc.getW() * 2 != src1Desc.getW())
throw std::invalid_argument("invalid concat+conv source descriptor");
break;
default:
throw std::invalid_argument("unsupported concat+conv fusion");
}
if (weightDesc.getRank() != 4 || weightDesc.getI() != (src0Desc.getC() + src1Desc.getC()) ||
weightDesc.getPaddedI() != (src0Desc.getPaddedC() + src1Desc.getPaddedC()))
if (weightDesc.getRank() != 4 || weightDesc.getI() != (src1Desc.getC() + src2Desc.getC()) ||
weightDesc.getPaddedI() != (src1Desc.getPaddedC() + src2Desc.getPaddedC()))
throw std::invalid_argument("invalid concat+conv weight shape");
TensorDims dstDims{weightDesc.getO(), src1Desc.getH(), src1Desc.getW()};
@@ -41,17 +24,17 @@ OIDN_NAMESPACE_BEGIN
dstDesc = {dstDims, dstPaddedDims, src1Desc.layout, src1Desc.dataType};
}
void ConcatConvBase::setSrc(const Ref<Tensor>& src0, const Ref<Tensor>& src1)
void ConcatConv::setSrc(const Ref<Tensor>& src1, const Ref<Tensor>& src2)
{
if (!src0 || src0->getDesc() != src0Desc || !src1 || src1->getDesc() != src1Desc)
if (!src1 || src1->getDesc() != src1Desc || !src2 || src2->getDesc() != src2Desc)
throw std::invalid_argument("invalid concat+conv source");
this->src0 = src0;
this->src1 = src1;
this->src2 = src2;
updateSrc();
}
void ConcatConvBase::setBias(const Ref<Tensor>& bias)
void ConcatConv::setBias(const Ref<Tensor>& bias)
{
if (!bias || bias->getDesc() != biasDesc)
throw std::invalid_argument("invalid concat+conv bias");
@@ -60,7 +43,7 @@ OIDN_NAMESPACE_BEGIN
updateBias();
}
void ConcatConvBase::setDst(const Ref<Tensor>& dst)
void ConcatConv::setDst(const Ref<Tensor>& dst)
{
if (!dst || dst->getDesc() != dstDesc)
throw std::invalid_argument("invalid concat+conv destination");
@@ -69,144 +52,4 @@ OIDN_NAMESPACE_BEGIN
updateDst();
}
// -----------------------------------------------------------------------------------------------
void ConcatConv::setWeight(const Ref<Tensor>& weight)
{
if (!weight || weight->getDesc() != weightDesc)
throw std::invalid_argument("invalid concat+conv weight");
this->weight = weight;
updateWeight();
}
// -----------------------------------------------------------------------------------------------
ConcatConv2::ConcatConv2(const ConcatConvDesc& desc)
: ConcatConvBase(desc)
{
if (src0Desc.layout != TensorLayout::hwc)
throw std::logic_error("unsupported concat+conv source layout");
if (fusion != Fusion::None && fusion != Fusion::UpsampleSrc0) // only pre-ops supported
throw std::invalid_argument("unsupported concat+conv fusion");
// Split the convolution into two smaller convolutions
weight0Desc = {{dstDesc.getC(), src0Desc.getC(), weightDesc.getH(), weightDesc.getW()},
{dstDesc.getPaddedC(), src0Desc.getPaddedC(), weightDesc.getH(), weightDesc.getW()},
weightDesc.layout,
weightDesc.dataType};
weight1Desc = {{dstDesc.getC(), src1Desc.getC(), weightDesc.getH(), weightDesc.getW()},
{dstDesc.getPaddedC(), src1Desc.getPaddedC(), weightDesc.getH(), weightDesc.getW()},
weightDesc.layout,
weightDesc.dataType};
}
void ConcatConv2::setWeight(const Ref<Tensor>& weight0, const Ref<Tensor>& weight1)
{
if (!weight0 || weight0->getDesc() != weight0Desc)
throw std::invalid_argument("invalid concat+conv weight0");
if (!weight1 || weight1->getDesc() != weight1Desc)
throw std::invalid_argument("invalid concat+conv weight1");
this->weight0 = weight0;
this->weight1 = weight1;
updateWeight();
}
// -----------------------------------------------------------------------------------------------
PreConcatConvCHW::PreConcatConvCHW(Engine* engine, const ConcatConvDesc& desc)
: ConcatConv(desc)
{
if (src0Desc.layout == TensorLayout::hwc)
throw std::invalid_argument("unsupported concat+conv source layout");
if (fusion != Fusion::None && fusion != Fusion::PoolDst) // only post-ops supported
throw std::invalid_argument("unsupported concat+conv fusion");
TensorDims srcDims{src0Desc.getC() + src1Desc.getC(), src0Desc.getH(), src0Desc.getW()};
TensorDims srcPaddedDims{src0Desc.getPaddedC() + src1Desc.getPaddedC(), src0Desc.getH(), src0Desc.getW()};
srcDesc = {srcDims, srcPaddedDims, src0Desc.layout, src0Desc.dataType};
conv = engine->newConv({srcDesc, weightDesc, biasDesc, activation, fusion, fastMath});
}
void PreConcatConvCHW::updateSrc()
{
if (!src0->getBuffer() || !src1->getBuffer())
throw std::invalid_argument("concat+conv sources must be backed by buffers");
if (src0->getBuffer() != src1->getBuffer() ||
(static_cast<char*>(src0->getPtr()) + src0->getByteSize()) != static_cast<char*>(src1->getPtr()))
throw std::invalid_argument("concat+conv sources are not pre-concatenated in memory");
auto src = src0->getBuffer()->newTensor(srcDesc, src0->getByteOffset());
conv->setSrc(src);
}
// -----------------------------------------------------------------------------------------------
InplaceConcatConv2::InplaceConcatConv2(Engine* engine, const ConcatConvDesc& desc)
: ConcatConv2(desc)
{
// Convolution 0: dst = conv(src0, weight0) + bias
conv0 = engine->newConv({src0Desc, weight0Desc, biasDesc, Activation::None, fusion, fastMath});
// Convolution 1: dst = activation(conv(src1, weight1) + dst)
// We use dst as bias
conv1 = engine->newConv({src1Desc, weight1Desc, dstDesc, activation, Fusion::None, fastMath});
}
bool InplaceConcatConv2::isSupported() const
{
return conv0->isSupported() && conv1->isSupported();
}
size_t InplaceConcatConv2::getScratchByteSize()
{
return max(conv0->getScratchByteSize(), conv1->getScratchByteSize());
}
void InplaceConcatConv2::setScratch(const Ref<Buffer>& scratch)
{
conv0->setScratch(scratch);
conv1->setScratch(scratch);
}
void InplaceConcatConv2::updateSrc()
{
conv0->setSrc(src0);
conv1->setSrc(src1);
}
void InplaceConcatConv2::updateWeight()
{
conv0->setWeight(weight0);
conv1->setWeight(weight1);
}
void InplaceConcatConv2::updateBias()
{
conv0->setBias(bias);
}
void InplaceConcatConv2::updateDst()
{
conv0->setDst(dst);
conv1->setBias(dst);
conv1->setDst(dst);
}
void InplaceConcatConv2::finalize()
{
conv0->finalize();
conv1->finalize();
}
void InplaceConcatConv2::submitKernels(const Ref<CancellationToken>& ct)
{
conv0->submitKernels(ct);
conv1->submitKernels(ct);
}
OIDN_NAMESPACE_END
+5 -92
View File
@@ -10,24 +10,23 @@ OIDN_NAMESPACE_BEGIN
// Concatenation + convolution descriptor
struct ConcatConvDesc
{
TensorDesc src0Desc;
TensorDesc src1Desc;
TensorDesc src2Desc;
TensorDesc weightDesc;
TensorDesc biasDesc;
Activation activation;
Fusion fusion;
bool fastMath; // prefer performance over accuracy
};
class ConcatConvBase : public BaseOp, protected ConcatConvDesc
class ConcatConv : public BaseOp, protected ConcatConvDesc
{
public:
ConcatConvBase(const ConcatConvDesc& desc);
ConcatConv(const ConcatConvDesc& desc);
TensorDesc getDstDesc() const { return dstDesc; }
Ref<Tensor> getDst() const { return dst; }
void setSrc(const Ref<Tensor>& src0, const Ref<Tensor>& src1);
void setSrc(const Ref<Tensor>& src1, const Ref<Tensor>& src2);
void setBias(const Ref<Tensor>& bias);
void setDst(const Ref<Tensor>& dst);
@@ -38,96 +37,10 @@ OIDN_NAMESPACE_BEGIN
TensorDesc dstDesc;
Ref<Tensor> src0;
Ref<Tensor> src1;
Ref<Tensor> src2;
Ref<Tensor> bias;
Ref<Tensor> dst;
};
class ConcatConv : public ConcatConvBase
{
public:
ConcatConv(const ConcatConvDesc& desc) : ConcatConvBase(desc) {}
// Is pre-concatenation of the source tensors required?
virtual bool isPreConcatConv() const { return false; }
void setWeight(const Ref<Tensor>& weight);
protected:
virtual void updateWeight() {}
Ref<Tensor> weight;
};
// Concatenation + convolution implemented as two separate convolutions
class ConcatConv2 : public ConcatConvBase
{
public:
ConcatConv2(const ConcatConvDesc& desc);
TensorDesc getWeight0Desc() const { return weight0Desc; }
TensorDesc getWeight1Desc() const { return weight1Desc; }
void setWeight(const Ref<Tensor>& weight0, const Ref<Tensor>& weight1);
protected:
virtual void updateWeight() {}
TensorDesc weight0Desc;
TensorDesc weight1Desc;
Ref<Tensor> weight0;
Ref<Tensor> weight1;
};
// Concatenation + convolution for pre-concatenated CHW tensors (stored consecutively in memory)
class PreConcatConvCHW final : public ConcatConv
{
public:
PreConcatConvCHW(Engine* engine, const ConcatConvDesc& desc);
Engine* getEngine() const override { return conv->getEngine(); }
bool isPreConcatConv() const override { return true; }
size_t getScratchByteSize() override { return conv->getScratchByteSize(); }
void setScratch(const Ref<Buffer>& scratch) override { conv->setScratch(scratch); }
void finalize() override { conv->finalize(); }
void submitKernels(const Ref<CancellationToken>& ct) override { conv->submitKernels(ct); }
private:
void updateSrc() override;
void updateWeight() override { conv->setWeight(weight); }
void updateBias() override { conv->setBias(bias); }
void updateDst() override { conv->setDst(dst); }
TensorDesc srcDesc; // pre-concatenated source
Ref<Conv> conv;
};
// Executes the two convolutions of ConcatConv2 in-place without an intermediate buffer for the
// first convolution's output
class InplaceConcatConv2 final : public ConcatConv2
{
public:
InplaceConcatConv2(Engine* engine, const ConcatConvDesc& desc);
Engine* getEngine() const override { return conv0->getEngine(); }
bool isSupported() const override;
size_t getScratchByteSize() override;
void setScratch(const Ref<Buffer>& scratch) override;
void finalize() override;
void submitKernels(const Ref<CancellationToken>& ct) override;
private:
void updateSrc() override;
void updateWeight() override;
void updateBias() override;
void updateDst() override;
Ref<Conv> conv0;
Ref<Conv> conv1;
};
OIDN_NAMESPACE_END
+34
View File
@@ -0,0 +1,34 @@
// Copyright 2023 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
#include "concat_conv_chw.h"
#include "engine.h"
OIDN_NAMESPACE_BEGIN
ConcatConvCHW::ConcatConvCHW(Engine* engine, const ConcatConvDesc& desc)
: ConcatConv(desc)
{
if (src1Desc.layout == TensorLayout::hwc)
throw std::invalid_argument("unsupported concat+conv source layout");
TensorDims srcDims{src1Desc.getC() + src2Desc.getC(), src1Desc.getH(), src1Desc.getW()};
TensorDims srcPaddedDims{src1Desc.getPaddedC() + src2Desc.getPaddedC(), src1Desc.getH(), src1Desc.getW()};
srcDesc = {srcDims, srcPaddedDims, src1Desc.layout, src1Desc.dataType};
conv = engine->newConv({srcDesc, weightDesc, biasDesc, activation, PostOp::None, fastMath});
}
void ConcatConvCHW::updateSrc()
{
if (!src1->getBuffer() || !src2->getBuffer())
throw std::invalid_argument("concat+conv sources must be backed by buffers");
if (src1->getBuffer() != src2->getBuffer() ||
(static_cast<char*>(src1->getPtr()) + src1->getByteSize()) != static_cast<char*>(src2->getPtr()))
throw std::invalid_argument("concat+conv sources are not pre-concatenated in memory");
auto src = src1->getBuffer()->newTensor(srcDesc, src1->getByteOffset());
conv->setSrc(src);
}
OIDN_NAMESPACE_END
+36
View File
@@ -0,0 +1,36 @@
// Copyright 2023 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
#pragma once
#include "concat_conv.h"
OIDN_NAMESPACE_BEGIN
// Concatenation + convolution for CHW tensors (including blocked) stored consecutively in memory
// Since the tensors are pre-concatenated in memory, only the convolution needs to be executed
class ConcatConvCHW final : public ConcatConv
{
public:
ConcatConvCHW(Engine* engine, const ConcatConvDesc& desc);
Engine* getEngine() const override { return conv->getEngine(); }
size_t getScratchByteSize() override { return conv->getScratchByteSize(); }
void setScratch(const Ref<Buffer>& scratch) override { conv->setScratch(scratch); }
void setWeight(const Ref<Tensor>& weight) { conv->setWeight(weight); }
void finalize() override { conv->finalize(); }
void submitKernels(const Ref<CancellationToken>& ct) override { conv->submitKernels(ct); }
private:
void updateSrc() override;
void updateBias() override { conv->setBias(bias); }
void updateDst() override { conv->setDst(dst); }
TensorDesc srcDesc; // pre-concatenated source
Ref<Conv> conv;
};
OIDN_NAMESPACE_END
+82
View File
@@ -0,0 +1,82 @@
// Copyright 2022 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
#include "concat_conv_hwc.h"
#include "engine.h"
OIDN_NAMESPACE_BEGIN
ConcatConvHWC::ConcatConvHWC(Engine* engine, const ConcatConvDesc& desc)
: ConcatConv(desc)
{
if (src1Desc.layout != TensorLayout::hwc)
throw std::logic_error("unsupported concat+conv source layout");
// Split the convolution into two smaller convolutions
weight1Desc = {{dstDesc.getC(), src1Desc.getC(), weightDesc.getH(), weightDesc.getW()},
{dstDesc.getPaddedC(), src1Desc.getPaddedC(), weightDesc.getH(), weightDesc.getW()},
weightDesc.layout,
weightDesc.dataType};
weight2Desc = {{dstDesc.getC(), src2Desc.getC(), weightDesc.getH(), weightDesc.getW()},
{dstDesc.getPaddedC(), src2Desc.getPaddedC(), weightDesc.getH(), weightDesc.getW()},
weightDesc.layout,
weightDesc.dataType};
// Convolution 1: dst = conv(src1, weight1) + bias
conv1 = engine->newConv({src1Desc, weight1Desc, biasDesc, Activation::None, PostOp::None, fastMath});
// Convolution 2: dst = activation(conv(src2, weight2) + dst)
// We use dst as bias
conv2 = engine->newConv({src2Desc, weight2Desc, dstDesc, activation, PostOp::None, fastMath});
}
bool ConcatConvHWC::isSupported() const
{
return conv1->isSupported() && conv2->isSupported();
}
size_t ConcatConvHWC::getScratchByteSize()
{
return max(conv1->getScratchByteSize(), conv2->getScratchByteSize());
}
void ConcatConvHWC::setScratch(const Ref<Buffer>& scratch)
{
conv1->setScratch(scratch);
conv2->setScratch(scratch);
}
void ConcatConvHWC::setWeight(const Ref<Tensor>& weight1, const Ref<Tensor>& weight2)
{
conv1->setWeight(weight1);
conv2->setWeight(weight2);
}
void ConcatConvHWC::updateSrc()
{
conv1->setSrc(src1);
conv2->setSrc(src2);
}
void ConcatConvHWC::updateDst()
{
conv1->setDst(dst);
conv2->setBias(dst);
conv2->setDst(dst);
}
void ConcatConvHWC::finalize()
{
conv1->finalize();
conv2->finalize();
}
void ConcatConvHWC::submitKernels(const Ref<CancellationToken>& ct)
{
conv1->submitKernels(ct);
conv2->submitKernels(ct);
}
OIDN_NAMESPACE_END
+43
View File
@@ -0,0 +1,43 @@
// Copyright 2022 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
#pragma once
#include "concat_conv.h"
OIDN_NAMESPACE_BEGIN
// Concatenation + convolution for HWC tensors
// The convolution is split into two smaller convolutions, one for each input tensor
// The weights for each convolution must be set separately
class ConcatConvHWC final : public ConcatConv
{
public:
ConcatConvHWC(Engine* engine, const ConcatConvDesc& desc);
Engine* getEngine() const override { return conv1->getEngine(); }
bool isSupported() const override;
size_t getScratchByteSize() override;
void setScratch(const Ref<Buffer>& scratch) override;
TensorDesc getWeight1Desc() const { return weight1Desc; }
TensorDesc getWeight2Desc() const { return weight2Desc; }
void setWeight(const Ref<Tensor>& weight1, const Ref<Tensor>& weight2);
void finalize() override;
void submitKernels(const Ref<CancellationToken>& ct) override;
private:
void updateSrc() override;
void updateBias() override { conv1->setBias(bias); }
void updateDst() override;
TensorDesc weight1Desc;
TensorDesc weight2Desc;
Ref<Conv> conv1;
Ref<Conv> conv2;
};
OIDN_NAMESPACE_END
+8 -8
View File
@@ -16,24 +16,24 @@ OIDN_NAMESPACE_BEGIN
throw std::invalid_argument("invalid convolution weight shape");
TensorDims dstDims;
switch (fusion)
switch (postOp)
{
case Fusion::None:
case PostOp::None:
dstDims = {weightDesc.getO(), srcDesc.getH(), srcDesc.getW()};
break;
case Fusion::UpsampleSrc0:
dstDims = {weightDesc.getO(), srcDesc.getH() * 2, srcDesc.getW() * 2};
break;
case Fusion::PoolDst:
case PostOp::Pool:
if (srcDesc.getH() % 2 != 0 || srcDesc.getW() % 2 != 0)
throw std::invalid_argument("invalid pooling source shape");
dstDims = {weightDesc.getO(), srcDesc.getH() / 2, srcDesc.getW() / 2};
break;
case PostOp::Upsample:
dstDims = {weightDesc.getO(), srcDesc.getH() * 2, srcDesc.getW() * 2};
break;
default:
throw std::invalid_argument("unsupported convolution fusion");
throw std::invalid_argument("unsupported convolution postop");
}
TensorDims dstPaddedDims = dstDims;
+4 -4
View File
@@ -15,11 +15,11 @@ OIDN_NAMESPACE_BEGIN
ReLU
};
enum class Fusion
enum class PostOp
{
None,
UpsampleSrc0, // upsample first source
PoolDst, // pool destination
Pool,
Upsample
};
// Convolution descriptor
@@ -29,7 +29,7 @@ OIDN_NAMESPACE_BEGIN
TensorDesc weightDesc;
TensorDesc biasDesc;
Activation activation;
Fusion fusion;
PostOp postOp;
bool fastMath; // prefer performance over accuracy
};
+3 -31
View File
@@ -196,9 +196,7 @@ OIDN_NAMESPACE_BEGIN
else if (name == "managedMemorySupported")
return managedMemorySupported;
else if (name == "externalMemoryTypes")
return static_cast<OIDNFlags>(externalMemoryTypes);
else if (name == "externalSemaphoreTypes")
return static_cast<OIDNFlags>(externalSemaphoreTypes);
return static_cast<int>(externalMemoryTypes);
else
throw Exception(Error::InvalidArgument, "unknown device parameter or type mismatch: '" + name + "'");
}
@@ -270,44 +268,18 @@ OIDN_NAMESPACE_BEGIN
return getEngine()->newNativeBuffer(handle)->toUser();
}
Ref<Buffer> Device::newExternalUserBuffer(ExternalMemoryTypeFlags fdType,
Ref<Buffer> Device::newExternalUserBuffer(ExternalMemoryTypeFlag fdType,
int fd, size_t byteSize)
{
return getEngine()->newExternalBuffer(fdType, fd, byteSize)->toUser();
}
Ref<Buffer> Device::newExternalUserBuffer(ExternalMemoryTypeFlags handleType,
Ref<Buffer> Device::newExternalUserBuffer(ExternalMemoryTypeFlag handleType,
void* handle, const void* name, size_t byteSize)
{
return getEngine()->newExternalBuffer(handleType, handle, name, byteSize)->toUser();
}
Ref<Semaphore> Device::newExternalSemaphore(ExternalSemaphoreTypeFlags fdType, int fd)
{
return getEngine()->newExternalSemaphore(fdType, fd);
}
Ref<Semaphore> Device::newExternalSemaphore(ExternalSemaphoreTypeFlags handleType,
void* handle, const void* name)
{
return getEngine()->newExternalSemaphore(handleType, handle, name);
}
void Device::submitSignalSemaphores(Semaphore* const* semaphores,
const uint64_t* values,
int numSemaphores)
{
getEngine()->submitSignalSemaphores(semaphores, values, numSemaphores);
}
void Device::submitWaitSemaphores(Semaphore* const* semaphores,
const uint64_t* values,
const uint32_t* timeoutsMs,
int numSemaphores)
{
getEngine()->submitWaitSemaphores(semaphores, values, timeoutsMs, numSemaphores);
}
Ref<Filter> Device::newFilter(const std::string& type)
{
if (isVerbose(2))
+2 -21
View File
@@ -17,7 +17,6 @@ OIDN_NAMESPACE_BEGIN
class Subdevice;
class Engine;
class Buffer;
class Semaphore;
class Filter;
class PhysicalDevice : public RefCount
@@ -78,10 +77,10 @@ OIDN_NAMESPACE_BEGIN
Ref<Buffer> newUserBuffer(void* ptr, size_t byteSize);
Ref<Buffer> newNativeUserBuffer(void* handle);
Ref<Buffer> newExternalUserBuffer(ExternalMemoryTypeFlags fdType,
Ref<Buffer> newExternalUserBuffer(ExternalMemoryTypeFlag fdType,
int fd, size_t byteSize);
Ref<Buffer> newExternalUserBuffer(ExternalMemoryTypeFlags handleType,
Ref<Buffer> newExternalUserBuffer(ExternalMemoryTypeFlag handleType,
void* handle, const void* name, size_t byteSize);
// Filter
@@ -111,22 +110,6 @@ OIDN_NAMESPACE_BEGIN
ExternalMemoryTypeFlags getExternalMemoryTypes() const { return externalMemoryTypes; }
void trimScratch();
// Semaphore
Ref<Semaphore> newExternalSemaphore(ExternalSemaphoreTypeFlags fdType, int fd);
Ref<Semaphore> newExternalSemaphore(ExternalSemaphoreTypeFlags handleType,
void* handle, const void* name);
void submitSignalSemaphores(Semaphore* const* semaphores,
const uint64_t* values,
int numSemaphores);
void submitWaitSemaphores(Semaphore* const* semaphores,
const uint64_t* values,
const uint32_t* timeoutsMs,
int numSemaphores);
ExternalSemaphoreTypeFlags getExternalSemaphoreTypes() const { return externalSemaphoreTypes; }
// Executes operations on the device, making sure to wait/flush and release temporary
// allocations (e.g. from ObjC) at the end, even if an exception is thrown
virtual void execute(std::function<void()>&& f, SyncMode sync = SyncMode::Blocking);
@@ -164,8 +147,6 @@ OIDN_NAMESPACE_BEGIN
bool managedMemorySupported = false;
ExternalMemoryTypeFlags externalMemoryTypes;
ExternalSemaphoreTypeFlags externalSemaphoreTypes;
// State
bool dirty = true;
bool committed = false;
+4 -45
View File
@@ -45,51 +45,20 @@ OIDN_NAMESPACE_BEGIN
"creating a shared buffer from a native handle is not supported by the device");
}
Ref<Buffer> Engine::newExternalBuffer(ExternalMemoryTypeFlags fdType,
Ref<Buffer> Engine::newExternalBuffer(ExternalMemoryTypeFlag fdType,
int fd, size_t byteSize)
{
throw Exception(Error::InvalidOperation,
"creating a shared buffer from a POSIX file descriptor is not supported by the device");
}
Ref<Buffer> Engine::newExternalBuffer(ExternalMemoryTypeFlags handleType,
Ref<Buffer> Engine::newExternalBuffer(ExternalMemoryTypeFlag handleType,
void* handle, const void* name, size_t byteSize)
{
throw Exception(Error::InvalidOperation,
"creating a shared buffer from a Win32 handle is not supported by the device");
}
Ref<Semaphore> Engine::newExternalSemaphore(ExternalSemaphoreTypeFlags fdType,
int fd)
{
throw Exception(Error::InvalidOperation,
"creating a shared semaphore from a POSIX file descriptor is not supported by the device");
}
Ref<Semaphore> Engine::newExternalSemaphore(ExternalSemaphoreTypeFlags handleType,
void* handle, const void* name)
{
throw Exception(Error::InvalidOperation,
"creating a shared semaphore from a Win32 handle is not supported by the device");
}
void Engine::submitSignalSemaphores(Semaphore* const* semaphores,
const uint64_t* values,
int numSemaphores)
{
throw Exception(Error::InvalidOperation,
"signaling semaphores is not supported by the device");
}
void Engine::submitWaitSemaphores(Semaphore* const* semaphores,
const uint64_t* values,
const uint32_t* timeoutsMs,
int numSemaphores)
{
throw Exception(Error::InvalidOperation,
"waiting on semaphores is not supported by the device");
}
bool Engine::isSupported(const TensorDesc& desc) const
{
// We store tensor byte offsets in 32-bit unsigned integers
@@ -114,19 +83,9 @@ OIDN_NAMESPACE_BEGIN
return makeRef<DeviceTensor>(buffer, desc, byteOffset);
}
bool Engine::isConvSupported(Fusion fusion) const
bool Engine::isConvSupported(PostOp postOp)
{
return fusion == Fusion::None;
}
Ref<ConcatConv> Engine::newConcatConv(const ConcatConvDesc& desc)
{
throw std::logic_error("concat+conv operation is not implemented");
}
Ref<ConcatConv2> Engine::newConcatConv2(const ConcatConvDesc& desc)
{
throw std::logic_error("concat+conv operation is not implemented");
return postOp == PostOp::None;
}
void* Engine::usmAlloc(size_t byteSize, Storage storage)
+5 -27
View File
@@ -9,7 +9,6 @@
#include "heap.h"
#include "buffer.h"
#include "image.h"
#include "semaphore.h"
#include "progress.h"
OIDN_NAMESPACE_BEGIN
@@ -21,10 +20,9 @@ OIDN_NAMESPACE_BEGIN
struct InputProcessDesc;
struct OutputProcessDesc;
enum class PostOp;
class Conv;
enum class Fusion;
class ConcatConv;
class ConcatConv2;
class Pool;
class Upsample;
class Autoexposure;
@@ -55,40 +53,20 @@ OIDN_NAMESPACE_BEGIN
virtual Ref<Buffer> newNativeBuffer(void* handle);
virtual Ref<Buffer> newExternalBuffer(ExternalMemoryTypeFlags fdType,
virtual Ref<Buffer> newExternalBuffer(ExternalMemoryTypeFlag fdType,
int fd, size_t byteSize);
virtual Ref<Buffer> newExternalBuffer(ExternalMemoryTypeFlags handleType,
virtual Ref<Buffer> newExternalBuffer(ExternalMemoryTypeFlag handleType,
void* handle, const void* name, size_t byteSize);
// Tensor
virtual bool isSupported(const TensorDesc& desc) const;
virtual Ref<Tensor> newTensor(const TensorDesc& desc, Storage storage = Storage::Device);
virtual Ref<Tensor> newTensor(const Ref<Buffer>& buffer, const TensorDesc& desc, size_t byteOffset = 0);
// Semaphore
virtual Ref<Semaphore> newExternalSemaphore(ExternalSemaphoreTypeFlags fdType,
int fd);
virtual Ref<Semaphore> newExternalSemaphore(ExternalSemaphoreTypeFlags handleType,
void* handle, const void* name);
virtual void submitSignalSemaphores(Semaphore* const* semaphores,
const uint64_t* values,
int numSemaphores);
virtual void submitWaitSemaphores(Semaphore* const* semaphores,
const uint64_t* values,
const uint32_t* timeoutsMs,
int numSemaphores);
// Ops
virtual bool isConvSupported(Fusion fusion) const;
virtual bool isConcatConvSupported(Fusion fusion) const { return false; }
virtual bool isConcatConv2Supported(Fusion fusion) const { return false; }
virtual bool isConvSupported(PostOp postOp);
virtual Ref<Conv> newConv(const ConvDesc& desc) = 0;
virtual Ref<ConcatConv> newConcatConv(const ConcatConvDesc& desc);
virtual Ref<ConcatConv2> newConcatConv2(const ConcatConvDesc& desc);
virtual Ref<Pool> newPool(const PoolDesc& desc) = 0;
virtual Ref<Upsample> newUpsample(const UpsampleDesc& desc) = 0;
virtual Ref<Autoexposure> newAutoexposure(const ImageDesc& srcDesc) = 0;
+82 -117
View File
@@ -4,6 +4,8 @@
//#define OIDN_MICROBENCH 1000 // number of microbenchmark iterations
#include "graph.h"
#include "concat_conv_chw.h"
#include "concat_conv_hwc.h"
#include "tensor_reorder.h"
#if defined(OIDN_MICROBENCH)
#include "common/timer.h"
@@ -60,27 +62,20 @@ OIDN_NAMESPACE_BEGIN
Ref<Op> Graph::addConv(const std::string& name,
const Ref<Op>& srcOp,
Activation activation,
Fusion fusion)
PostOp postOp)
{
if (fusion != Fusion::None && !engine->isConvSupported(fusion))
if (postOp != PostOp::None && !engine->isConvSupported(postOp))
{
// If the engine does not support the specified fused convolution, split it into two ops
switch (fusion)
auto conv = addConv(name, srcOp, activation, PostOp::None);
switch (postOp)
{
case Fusion::UpsampleSrc0:
{
auto srcUpsample = addUpsample(name + "_upsample", srcOp);
return addConv(name, srcUpsample, activation);
}
case Fusion::PoolDst:
{
auto conv = addConv(name, srcOp, activation);
return addPool(name + "_pool", conv);
}
case PostOp::Pool:
return addPool(name + "_pool", conv);
case PostOp::Upsample:
return addUpsample(name + "_upsample", conv);
default:
throw std::invalid_argument("unsupported convolution fusion");
throw std::invalid_argument("cannot split fused convolution");
}
}
@@ -111,7 +106,7 @@ OIDN_NAMESPACE_BEGIN
device->getTensorDataType()};
auto srcAlloc = tensorAllocs[srcOp.get()];
auto conv = engine->newConv({srcAlloc->desc, finalWeightDesc, finalBiasDesc, activation, fusion, fastMath});
auto conv = engine->newConv({srcAlloc->desc, finalWeightDesc, finalBiasDesc, activation, postOp, fastMath});
conv->setName(name);
auto dstAlloc = addOp(conv, {srcOp}, conv->getDstDesc());
@@ -148,28 +143,11 @@ OIDN_NAMESPACE_BEGIN
return conv;
}
Ref<Op> Graph::addConvPool(const std::string& name,
const Ref<Op>& srcOp,
Activation activation)
{
return addConv(name, srcOp, activation, Fusion::PoolDst);
}
Ref<Op> Graph::addConcatConv(const std::string& name,
const Ref<Op>& src0Op,
const Ref<Op>& src1Op,
Activation activation,
Fusion fusion)
const Ref<Op>& src2Op,
Activation activation)
{
Device* device = engine->getDevice();
if (!engine->isConcatConvSupported(fusion) && !engine->isConcatConv2Supported(fusion))
{
// Need to split into separate upsample and concat+conv
auto src0Upsample = addUpsample(name + "_upsample0", src0Op);
return addConcatConv(name, src0Upsample, src1Op, activation);
}
const std::string weightName = name + ".weight";
const std::string biasName = name + ".bias";
Ref<Tensor> weight = (*constTensors)[weightName];
@@ -178,15 +156,16 @@ OIDN_NAMESPACE_BEGIN
if (weight->getRank() != 4 || bias->getRank() != 1)
throw std::invalid_argument("invalid convolution weight/bias");
Device* device = engine->getDevice();
const int blockC = device->getTensorBlockC();
auto src0Alloc = tensorAllocs[src0Op.get()];
auto src1Alloc = tensorAllocs[src1Op.get()];
const TensorDesc src0Desc = src0Alloc->desc;
auto src2Alloc = tensorAllocs[src2Op.get()];
const TensorDesc src1Desc = src1Alloc->desc;
const TensorDesc src2Desc = src2Alloc->desc;
TensorDims finalWeightDims{round_up(weight->getO(), blockC),
src0Desc.getPaddedC() + src1Desc.getPaddedC(),
src1Desc.getPaddedC() + src2Desc.getPaddedC(),
weight->getH(),
weight->getW()};
@@ -200,18 +179,72 @@ OIDN_NAMESPACE_BEGIN
TensorLayout::x,
device->getTensorDataType()};
ConcatConvDesc concatConvDesc{src0Desc, src1Desc, finalWeightDesc, finalBiasDesc, activation, fusion, fastMath};
ConcatConvDesc concatConvDesc{src1Desc, src2Desc, finalWeightDesc, finalBiasDesc, activation, fastMath};
if (engine->isConcatConvSupported(fusion))
if (device->getTensorLayout() == TensorLayout::hwc)
{
auto concatConv = engine->newConcatConv(concatConvDesc);
auto concatConv = makeRef<ConcatConvHWC>(engine, concatConvDesc);
concatConv->setName(name);
auto dstAlloc = addOp(concatConv, {src0Op, src1Op}, concatConv->getDstDesc(),
concatConv->isPreConcatConv());
auto dstAlloc = addOp(concatConv, {src1Op, src2Op}, concatConv->getDstDesc());
lazyInits.push_back([=]()
{
concatConv->setSrc(src0Alloc->tensor, src1Alloc->tensor);
concatConv->setSrc(src1Alloc->tensor, src2Alloc->tensor);
concatConv->setDst(dstAlloc->tensor);
const std::string weight1Name = weightName + "1";
const std::string weight2Name = weightName + "2";
Ref<Tensor> finalWeight1 = getCachedConstTensor(weight1Name, concatConv->getWeight1Desc());
Ref<Tensor> finalWeight2 = getCachedConstTensor(weight2Name, concatConv->getWeight2Desc());
if (!finalWeight1 || !finalWeight2)
{
finalWeight1 = makeRef<HostTensor>(concatConv->getWeight1Desc());
finalWeight2 = makeRef<HostTensor>(concatConv->getWeight2Desc());
reorderWeight(*weight, 0, src1Desc.getC(),
*finalWeight1, 0, src1Desc.getPaddedC());
reorderWeight(*weight, src1Desc.getC(), src2Desc.getC(),
*finalWeight2, 0, src2Desc.getPaddedC());
if (device->needWeightAndBiasOnDevice())
{
finalWeight1 = finalWeight1->toDevice(engine);
finalWeight2 = finalWeight2->toDevice(engine);
}
setCachedConstTensor(weight1Name, finalWeight1);
setCachedConstTensor(weight2Name, finalWeight2);
}
Ref<Tensor> finalBias = getCachedConstTensor(biasName, finalBiasDesc);
if (!finalBias)
{
finalBias = makeRef<HostTensor>(finalBiasDesc);
reorderBias(*bias, *finalBias);
if (device->needWeightAndBiasOnDevice())
finalBias = finalBias->toDevice(engine);
setCachedConstTensor(biasName, finalBias);
}
concatConv->setWeight(finalWeight1, finalWeight2);
concatConv->setBias(finalBias);
});
privateByteSize += concatConv->getWeight1Desc().getByteSize() +
concatConv->getWeight2Desc().getByteSize() +
finalBiasDesc.getByteSize();
return concatConv;
}
else
{
auto concatConv = makeRef<ConcatConvCHW>(engine, concatConvDesc);
concatConv->setName(name);
auto dstAlloc = addOp(concatConv, {src1Op, src2Op}, concatConv->getDstDesc(), true);
lazyInits.push_back([=]()
{
concatConv->setSrc(src1Alloc->tensor, src2Alloc->tensor);
concatConv->setDst(dstAlloc->tensor);
Ref<Tensor> finalWeight = getCachedConstTensor(weightName, finalWeightDesc);
@@ -219,10 +252,10 @@ OIDN_NAMESPACE_BEGIN
{
finalWeight = makeRef<HostTensor>(finalWeightDesc);
reorderWeight(*weight, 0, src0Desc.getC(),
*finalWeight, 0, src0Desc.getPaddedC());
reorderWeight(*weight, src0Desc.getC(), src1Desc.getC(),
*finalWeight, src0Desc.getPaddedC(), src1Desc.getPaddedC());
reorderWeight(*weight, 0, src1Desc.getC(),
*finalWeight, 0, src1Desc.getPaddedC());
reorderWeight(*weight, src1Desc.getC(), src2Desc.getC(),
*finalWeight, src1Desc.getPaddedC(), src2Desc.getPaddedC());
if (device->needWeightAndBiasOnDevice())
finalWeight = finalWeight->toDevice(engine);
@@ -247,74 +280,6 @@ OIDN_NAMESPACE_BEGIN
privateByteSize += finalWeightDesc.getByteSize() + finalBiasDesc.getByteSize();
return concatConv;
}
else if (engine->isConcatConv2Supported(fusion))
{
// Concatenation + convolution is implemented as two separate convolutions
auto concatConv = engine->newConcatConv2(concatConvDesc);
concatConv->setName(name);
auto dstAlloc = addOp(concatConv, {src0Op, src1Op}, concatConv->getDstDesc());
lazyInits.push_back([=]()
{
concatConv->setSrc(src0Alloc->tensor, src1Alloc->tensor);
concatConv->setDst(dstAlloc->tensor);
const std::string weight0Name = weightName + "0";
const std::string weight1Name = weightName + "1";
Ref<Tensor> finalWeight0 = getCachedConstTensor(weight0Name, concatConv->getWeight0Desc());
Ref<Tensor> finalWeight1 = getCachedConstTensor(weight1Name, concatConv->getWeight1Desc());
if (!finalWeight0 || !finalWeight1)
{
finalWeight0 = makeRef<HostTensor>(concatConv->getWeight0Desc());
finalWeight1 = makeRef<HostTensor>(concatConv->getWeight1Desc());
reorderWeight(*weight, 0, src0Desc.getC(),
*finalWeight0, 0, src0Desc.getPaddedC());
reorderWeight(*weight, src0Desc.getC(), src1Desc.getC(),
*finalWeight1, 0, src1Desc.getPaddedC());
if (device->needWeightAndBiasOnDevice())
{
finalWeight0 = finalWeight0->toDevice(engine);
finalWeight1 = finalWeight1->toDevice(engine);
}
setCachedConstTensor(weight0Name, finalWeight0);
setCachedConstTensor(weight1Name, finalWeight1);
}
Ref<Tensor> finalBias = getCachedConstTensor(biasName, finalBiasDesc);
if (!finalBias)
{
finalBias = makeRef<HostTensor>(finalBiasDesc);
reorderBias(*bias, *finalBias);
if (device->needWeightAndBiasOnDevice())
finalBias = finalBias->toDevice(engine);
setCachedConstTensor(biasName, finalBias);
}
concatConv->setWeight(finalWeight0, finalWeight1);
concatConv->setBias(finalBias);
});
privateByteSize += concatConv->getWeight0Desc().getByteSize() +
concatConv->getWeight1Desc().getByteSize() +
finalBiasDesc.getByteSize();
return concatConv;
}
else
{
throw std::logic_error("concat+conv operation is not implemented");
}
}
Ref<Op> Graph::addUpsampleConcatConv(const std::string& name,
const Ref<Op>& src0Op,
const Ref<Op>& src1Op,
Activation activation)
{
return addConcatConv(name, src0Op, src1Op, activation, Fusion::UpsampleSrc0);
}
Ref<Op> Graph::addPool(const std::string& name,
+3 -13
View File
@@ -39,22 +39,12 @@ OIDN_NAMESPACE_BEGIN
Ref<Op> addConv(const std::string& name,
const Ref<Op>& srcOp,
Activation activation,
Fusion fusion = Fusion::None);
Ref<Op> addConvPool(const std::string& name,
const Ref<Op>& srcOp,
Activation activation);
PostOp postOp = PostOp::None);
Ref<Op> addConcatConv(const std::string& name,
const Ref<Op>& src0Op,
const Ref<Op>& src1Op,
Activation activation,
Fusion fusion = Fusion::None);
Ref<Op> addUpsampleConcatConv(const std::string& name,
const Ref<Op>& src0Op,
const Ref<Op>& src1Op,
Activation activation);
const Ref<Op>& src2Op,
Activation activation);
Ref<Op> addPool(const std::string& name,
const Ref<Op>& srcOp);
-153
View File
@@ -26,14 +26,6 @@ OIDN_NAMESPACE_BEGIN
return size_t(h) * hByteStride + size_t(w) * wByteStride;
}
#if defined(OIDN_COMPILE_SYCL)
template<int N>
oidn_inline simd<uint32_t, N> getByteOffset(int h, simd<int, N> w) const
{
return uint32_t(h) * uint32_t(hByteStride) + simd<uint32_t, N>(w) * uint32_t(wByteStride);
}
#endif
template<typename T = float>
oidn_host_device_inline vec3<T> get3(int h, int w) const
{
@@ -99,151 +91,6 @@ OIDN_NAMESPACE_BEGIN
pixel[0] = value.x;
}
}
#if defined(OIDN_COMPILE_SYCL)
template<typename T = float, int N>
oidn_inline vec3<simd<T, N>> get3(int h, simd<int, N> w, simd_mask<N> mask) const
{
vec3<simd<T, N>> c;
const simd<uint32_t, N> offset = getByteOffset(h, w);
if (dataType == DataType::Float32)
{
const float* base = reinterpret_cast<const float*>(ptr);
#if defined(OIDN_ARCH_XEHPG) || defined(OIDN_ARCH_XEHPC) || defined(OIDN_ARCH_XE2)
if (C == 3)
{
simd<float, N*3> data = lsc_gather<float, 3>(base, offset, mask);
c.x = data.template select<N, 1>(0);
c.y = data.template select<N, 1>(N);
c.z = data.template select<N, 1>(N*2);
}
else if (C == 2)
{
simd<float, N*2> data = lsc_gather<float, 2>(base, offset, mask);
c.x = data.template select<N, 1>(0);
c.y = data.template select<N, 1>(N);
c.z = c.y;
}
else // C == 1
{
c.x = lsc_gather<float>(base, offset, mask);
c.y = c.z = c.x;
}
#else
c.x = gather<float, N>(base, offset, mask);
if (C >= 2)
{
c.y = gather<float, N>(base, offset + uint32_t(sizeof(float)), mask);
if (C >= 3)
c.z = gather<float, N>(base, offset + uint32_t(2 * sizeof(float)), mask);
else
c.z = c.y;
}
else
c.y = c.z = c.x;
#endif
}
else // Float16
{
const half* base = reinterpret_cast<const half*>(ptr);
#if defined(OIDN_ARCH_XEHPG) || defined(OIDN_ARCH_XEHPC) || defined(OIDN_ARCH_XE2)
// d16u32 does not support vector lsc_gather (x2/x3)
c.x = lsc_gather<half>(base, offset, mask);
if (C >= 2)
{
c.y = lsc_gather<half>(base, offset + uint32_t(sizeof(half)), mask);
if (C >= 3)
c.z = lsc_gather<half>(base, offset + uint32_t(2 * sizeof(half)), mask);
else
c.z = c.y;
}
else
c.y = c.z = c.x;
#else
c.x = gather<half, N>(base, offset, mask);
if (C >= 2)
{
c.y = gather<half, N>(base, offset + uint32_t(sizeof(half)), mask);
if (C >= 3)
c.z = gather<half, N>(base, offset + uint32_t(2 * sizeof(half)), mask);
else
c.z = c.y;
}
else
c.y = c.z = c.x;
#endif
}
return c;
}
template<typename T, int N>
oidn_inline void set3(int h, simd<int, N> w, vec3<simd<T, N>> value, simd_mask<N> mask) const
{
const simd<uint32_t, N> offset = getByteOffset(h, w);
if (dataType == DataType::Float32)
{
float* base = reinterpret_cast<float*>(ptr);
#if defined(OIDN_ARCH_XEHPG) || defined(OIDN_ARCH_XEHPC) || defined(OIDN_ARCH_XE2)
if (C == 3)
{
simd<float, N*3> data;
data.template select<N, 1>(0) = simd<float, N>(value.x);
data.template select<N, 1>(N) = simd<float, N>(value.y);
data.template select<N, 1>(N*2) = simd<float, N>(value.z);
lsc_scatter<float, 3>(base, offset, data, mask);
}
else if (C == 2)
{
simd<float, N*2> data;
data.template select<N, 1>(0) = simd<float, N>(value.x);
data.template select<N, 1>(N) = simd<float, N>(value.y);
lsc_scatter<float, 2>(base, offset, data, mask);
}
else // C == 1
{
lsc_scatter<float>(base, offset, simd<float, N>(value.x), mask);
}
#else
scatter<float, N>(base, offset, value.x, mask);
if (C >= 2)
{
scatter<float, N>(base, offset + uint32_t(sizeof(float)), value.y, mask);
if (C >= 3)
scatter<float, N>(base, offset + uint32_t(sizeof(float)*2), value.z, mask);
}
#endif
}
else // Float16
{
half* base = reinterpret_cast<half*>(ptr);
#if defined(OIDN_ARCH_XEHPG) || defined(OIDN_ARCH_XEHPC) || defined(OIDN_ARCH_XE2)
// d16u32 does not support vector lsc_gather (x2/x3)
lsc_scatter<half>(base, offset, simd<half, N>(value.x), mask);
if (C >= 2)
{
lsc_scatter<half>(base, offset + uint32_t(sizeof(half)), simd<half, N>(value.y), mask);
if (C >= 3)
lsc_scatter<half>(base, offset + uint32_t(sizeof(half)*2), simd<half, N>(value.z), mask);
}
#else
scatter<half, N>(base, offset, simd<half, N>(value.x), mask);
if (C >= 2)
{
scatter<half, N>(base, offset + uint32_t(sizeof(half)), simd<half, N>(value.y), mask);
if (C >= 3)
scatter<half, N>(base, offset + uint32_t(sizeof(half)*2), simd<half, N>(value.z), mask);
}
#endif
}
}
#endif
};
OIDN_NAMESPACE_END
+1 -1
View File
@@ -190,7 +190,7 @@ OIDN_NAMESPACE_BEGIN
template<typename T>
oidn_device_inline void subgroupStore(GlobalPtr<T> dst, const T& x) const
{
sycl::ext::oneapi::experimental::group_store(item.get_sub_group(), x, dst.get());
item.get_sub_group().store(dst, x);
}
oidn_device_inline void subgroupBarrier() const
-9
View File
@@ -79,14 +79,5 @@ namespace math {
return isnan(x) ? 0.f : x;
}
#if defined(OIDN_COMPILE_SYCL)
template<int N>
oidn_inline simd<float, N> nan_to_zero(simd<float, N> x)
{
x.merge(0.f, x != x);
return x;
}
#endif
} // namespace math
OIDN_NAMESPACE_END
+1 -1
View File
@@ -34,7 +34,7 @@ OIDN_NAMESPACE_BEGIN
filename += ".dll";
#else
#if defined(OIDN_LIBRARY_VERSIONED)
const std::string versionStr = "." OIDN_VERSION_STRING_SHORT;
const std::string versionStr = "." OIDN_VERSION_STRING;
#else
const std::string versionStr = "";
#endif
-21
View File
@@ -1,21 +0,0 @@
// Copyright 2026 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
#pragma once
#include "device.h"
OIDN_NAMESPACE_BEGIN
class Semaphore : public RefCount
{
public:
explicit Semaphore(const Ref<Device>& device) : device(device) {}
Device* getDevice() const { return device.get(); }
protected:
Ref<Device> device;
};
OIDN_NAMESPACE_END
+2 -4
View File
@@ -27,16 +27,14 @@ OIDN_NAMESPACE_BEGIN
Tensor::Tensor(const TensorDesc& desc)
: TensorDesc(desc)
{
if (!desc.isValid())
throw std::invalid_argument("invalid tensor descriptor");
assert(desc.isValid());
}
Tensor::Tensor(const Ref<Buffer>& buffer, const TensorDesc& desc, size_t byteOffset)
: Memory(buffer, byteOffset),
TensorDesc(desc)
{
if (!desc.isValid())
throw std::invalid_argument("invalid tensor descriptor");
assert(desc.isValid());
}
Ref<Tensor> Tensor::toDevice(Engine* engine, Storage storage)
+149 -5
View File
@@ -4,8 +4,8 @@
#pragma once
#include "buffer.h"
#include "tensor_desc.h"
#include "tensor_accessor.h"
#include <vector>
#include <unordered_map>
#include <iostream>
@@ -13,6 +13,152 @@ OIDN_NAMESPACE_BEGIN
class Engine;
// Tensor dimensions
// Canonical order: CHW / OIHW
using TensorDims = std::vector<int>;
std::ostream& operator <<(std::ostream& sm, const TensorDims& dims);
// Tensor descriptor
struct TensorDesc
{
TensorDims dims; // logical dimensions
TensorDims paddedDims; // storage dimensions with zero-padding
TensorLayout layout; // storage layout
DataType dataType; // element data type
TensorDesc() = default;
TensorDesc(TensorDims dims, TensorDims paddedDims, TensorLayout layout, DataType dataType)
: dims(dims), paddedDims(paddedDims), layout(layout), dataType(dataType)
{
assert(isValid());
}
TensorDesc(TensorDims dims, TensorLayout layout, DataType dataType)
: dims(dims), paddedDims(dims), layout(layout), dataType(dataType)
{
assert(isValid());
}
bool isValid() const
{
const auto info = getTensorLayoutInfo(layout);
return getRank() == info.rank &&
dims.size() == paddedDims.size() &&
std::mismatch(dims.begin(), dims.end(), paddedDims.begin(),
[](int a, int b) { return a <= b; }).first == dims.end() &&
(info.blockC == 1 ||
(getRank() == 3 && getPaddedC() % info.blockC == 0) ||
(getRank() == 4 && getPaddedO() % info.blockC == 0 && getPaddedI() % info.blockC == 0));
}
// Returns the number of dimensions
oidn_inline int getRank() const { return int(dims.size()); }
// Returns the number of elements in a 1D tensor
oidn_inline int getX() const
{
assert(dims.size() == 1);
return dims[0];
}
oidn_inline int getPaddedX() const
{
assert(paddedDims.size() == 1);
return paddedDims[0];
}
// Returns the number of output channels in the tensor
oidn_inline int getO() const
{
assert(dims.size() >= 4);
return dims[dims.size()-4];
}
oidn_inline int getPaddedO() const
{
assert(paddedDims.size() >= 4);
return paddedDims[paddedDims.size()-4];
}
// Returns the number of input channels in the tensor
oidn_inline int getI() const
{
assert(dims.size() >= 3);
return dims[dims.size()-3];
}
oidn_inline int getPaddedI() const
{
assert(paddedDims.size() >= 3);
return paddedDims[paddedDims.size()-3];
}
// Returns the number of channels in the tensor
oidn_inline int getC() const
{
assert(dims.size() >= 3);
return dims[dims.size()-3];
}
oidn_inline int getPaddedC() const
{
assert(paddedDims.size() >= 3);
return paddedDims[paddedDims.size()-3];
}
// Returns the height of the tensor
oidn_inline int getH() const
{
assert(dims.size() >= 2);
return dims[dims.size()-2];
}
// Returns the width of the tensor
oidn_inline int getW() const
{
assert(dims.size() >= 2);
return dims[dims.size()-1];
}
// Returns the number of elements in the tensor
oidn_inline size_t getNumElements() const
{
if (dims.empty())
return 0;
size_t num = 1;
for (size_t i = 0; i < dims.size(); ++i)
num *= size_t(dims[i]);
return num;
}
// Returns the size in bytes of the tensor
oidn_inline size_t getByteSize() const
{
if (paddedDims.empty())
return 0;
size_t num = 1;
for (size_t i = 0; i < paddedDims.size(); ++i)
num *= size_t(paddedDims[i]);
return num * getDataTypeSize(dataType);
}
bool operator ==(const TensorDesc& other) const
{
return (dims == other.dims) && (paddedDims == other.paddedDims) &&
(layout == other.layout) && (dataType == other.dataType);
}
bool operator !=(const TensorDesc& other) const
{
return (dims != other.dims) || (paddedDims != other.paddedDims) ||
(layout != other.layout) || (dataType != other.dataType);
}
};
// Tensor
class Tensor : public Memory, protected TensorDesc
{
public:
@@ -62,10 +208,8 @@ OIDN_NAMESPACE_BEGIN
}
operator ispc::TensorAccessor1D();
operator ispc::TensorAccessor3D_chw();
operator ispc::TensorAccessor3D_ChwBc();
operator ispc::TensorAccessor4D_IOhwBiBo();
operator ispc::TensorAccessor4D_OIhwPoQiRoSi();
operator ispc::TensorAccessor3D();
operator ispc::TensorAccessor4D();
virtual Ref<Tensor> toDevice(Engine* engine, Storage storage = Storage::Device);
+2 -4
View File
@@ -10,10 +10,8 @@
namespace ispc
{
struct TensorAccessor1D;
struct TensorAccessor3D_chw;
struct TensorAccessor3D_ChwBc;
struct TensorAccessor4D_IOhwBiBo;
struct TensorAccessor4D_OIhwPoQiRoSi;
struct TensorAccessor3D;
struct TensorAccessor4D;
};
OIDN_NAMESPACE_BEGIN
-174
View File
@@ -1,174 +0,0 @@
// Copyright 2018 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
#pragma once
#include "tensor_layout.h"
#include <vector>
OIDN_NAMESPACE_BEGIN
// Tensor dimensions
// Canonical order: CHW / OIHW
using TensorDims = std::vector<int>;
std::ostream& operator <<(std::ostream& sm, const TensorDims& dims);
// Tensor descriptor
struct TensorDesc
{
TensorDims dims; // logical dimensions
TensorDims paddedDims; // storage dimensions with zero-padding
TensorLayout layout = TensorLayout::Undefined; // storage layout
DataType dataType = DataType::Undefined; // element data type
TensorDesc() = default;
TensorDesc(TensorDims dims, TensorDims paddedDims, TensorLayout layout, DataType dataType)
: dims(dims), paddedDims(paddedDims), layout(layout), dataType(dataType)
{
assert(isValid());
}
TensorDesc(TensorDims dims, TensorLayout layout, DataType dataType)
: dims(dims), paddedDims(dims), layout(layout), dataType(dataType)
{
assert(isValid());
}
bool isValid() const
{
const auto info = getTensorLayoutInfo(layout);
return getRank() == info.rank &&
dims.size() == paddedDims.size() &&
std::mismatch(dims.begin(), dims.end(), paddedDims.begin(),
[](int a, int b) { return a <= b; }).first == dims.end() &&
(info.blockC == 1 ||
(getRank() == 3 && getPaddedC() % info.blockC == 0) ||
(getRank() == 4 && getPaddedO() % info.blockC == 0 && getPaddedI() % info.blockC == 0));
}
// Returns the number of dimensions
oidn_inline int getRank() const { return int(dims.size()); }
// Returns the number of elements in a 1D tensor
oidn_inline int getX() const
{
assert(dims.size() == 1);
return dims[0];
}
oidn_inline int getPaddedX() const
{
assert(paddedDims.size() == 1);
return paddedDims[0];
}
// Returns the number of output channels in the tensor
oidn_inline int getO() const
{
assert(dims.size() >= 4);
return dims[dims.size()-4];
}
oidn_inline int getPaddedO() const
{
assert(paddedDims.size() >= 4);
return paddedDims[paddedDims.size()-4];
}
// Returns the number of input channels in the tensor
oidn_inline int getI() const
{
assert(dims.size() >= 3);
return dims[dims.size()-3];
}
oidn_inline int getPaddedI() const
{
assert(paddedDims.size() >= 3);
return paddedDims[paddedDims.size()-3];
}
// Returns the number of channels in the tensor
oidn_inline int getC() const
{
assert(dims.size() >= 3);
return dims[dims.size()-3];
}
oidn_inline int getPaddedC() const
{
assert(paddedDims.size() >= 3);
return paddedDims[paddedDims.size()-3];
}
// Returns the height of the tensor
oidn_inline int getH() const
{
assert(dims.size() >= 2);
return dims[dims.size()-2];
}
// Returns the width of the tensor
oidn_inline int getW() const
{
assert(dims.size() >= 2);
return dims[dims.size()-1];
}
// Returns the number of elements in the tensor
oidn_inline size_t getNumElements() const
{
if (dims.empty())
return 0;
size_t num = 1;
for (size_t i = 0; i < dims.size(); ++i)
num *= size_t(dims[i]);
return num;
}
// Returns the size in bytes of the tensor
oidn_inline size_t getByteSize() const
{
if (paddedDims.empty())
return 0;
const size_t elementSize = getDataTypeSize(dataType);
if (layout == TensorLayout::Chw8c ||
layout == TensorLayout::Chw16c ||
layout == TensorLayout::Chw32c)
{
// For blocked CHW layouts, the C planes need to be aligned
const size_t B = getTensorLayoutInfo(layout).blockC;
const size_t cByteStride = elementSize;
const size_t wByteStride = B * cByteStride;
const size_t hByteStride = size_t(getW()) * wByteStride;
const size_t CByteStride = round_up(size_t(getH()) * hByteStride, TensorLayoutTraitsChwBc::CByteAlignment);
return size_t(getPaddedC() / B) * CByteStride;
}
else
{
size_t num = 1;
for (size_t i = 0; i < paddedDims.size(); ++i)
num *= size_t(paddedDims[i]);
return num * elementSize;
}
}
bool operator ==(const TensorDesc& other) const
{
return (dims == other.dims) && (paddedDims == other.paddedDims) &&
(layout == other.layout) && (dataType == other.dataType);
}
bool operator !=(const TensorDesc& other) const
{
return (dims != other.dims) || (paddedDims != other.paddedDims) ||
(layout != other.layout) || (dataType != other.dataType);
}
};
OIDN_NAMESPACE_END
+21 -52
View File
@@ -10,23 +10,18 @@ OIDN_NAMESPACE_BEGIN
// Tensor memory layout
enum class TensorLayout
{
Undefined,
x,
chw,
Chw8c, // blocked
Chw16c, // blocked
Chw32c, // blocked
oihw,
OIhw8i8o, // blocked
OIhw16i16o, // blocked
OIhw2o8i8o2i, // blocked (Xe-HPG DPAS)
OIhw8i16o2i, // blocked (Xe-HPC DPAS)
OIhw2o16i16o2i, // blocked (AMX-FP16)
IOhw8i8o, // blocked
IOhw16i16o, // blocked
OIhw8i8o, // blocked
OIhw16i16o, // blocked
OIhw2o8i8o2i, // blocked (Xe-HPG DPAS)
OIhw8i16o2i, // blocked (Xe-HPC DPAS)
IOhw8i8o, // blocked
IOhw16i16o, // blocked
hwc,
ohwi,
@@ -108,12 +103,6 @@ OIDN_NAMESPACE_BEGIN
};
};
struct TensorLayoutTraitsChwBc
{
// Xe 2D block load/store requires C planes to be aligned
static constexpr oidn_constant uint32_t CByteAlignment = 64;
};
template<typename T, int B>
struct TensorByteOffsetChwBc
{
@@ -129,7 +118,7 @@ OIDN_NAMESPACE_BEGIN
oidn_host_device_inline TensorByteOffsetChwBc(int C, int H, int W)
{
hByteStride = uint32_t(W) * wByteStride;
CByteStride = round_up(uint32_t(H) * hByteStride, TensorLayoutTraitsChwBc::CByteAlignment);
CByteStride = uint32_t(H) * hByteStride;
}
oidn_host_device_inline uint32_t operator ()(int c, int h, int w) const
@@ -155,13 +144,6 @@ OIDN_NAMESPACE_BEGIN
using ByteOffset = TensorByteOffsetChwBc<T, 16>;
};
template<>
struct TensorLayoutTraits<TensorLayout::Chw32c>
{
template<typename T>
using ByteOffset = TensorByteOffsetChwBc<T, 32>;
};
template<>
struct TensorLayoutTraits<TensorLayout::oihw>
{
@@ -337,13 +319,6 @@ OIDN_NAMESPACE_BEGIN
using ByteOffset = TensorByteOffsetOIhwPoQiRoSi<T, 1, 8, 16, 2>;
};
template<>
struct TensorLayoutTraits<TensorLayout::OIhw2o16i16o2i>
{
template<typename T>
using ByteOffset = TensorByteOffsetOIhwPoQiRoSi<T, 2, 16, 16, 2>;
};
template<>
struct TensorLayoutTraits<TensorLayout::ohwi>
{
@@ -400,8 +375,6 @@ OIDN_NAMESPACE_BEGIN
return {3, 8};
case TensorLayout::Chw16c:
return {3, 16};
case TensorLayout::Chw32c:
return {3, 32};
case TensorLayout::oihw:
case TensorLayout::ohwi:
return {4, 1};
@@ -413,8 +386,6 @@ OIDN_NAMESPACE_BEGIN
case TensorLayout::OIhw2o8i8o2i:
case TensorLayout::OIhw8i16o2i:
return {4, 16};
case TensorLayout::OIhw2o16i16o2i:
return {4, 32};
default:
throw std::invalid_argument("invalid tensor layout");
}
@@ -424,22 +395,20 @@ OIDN_NAMESPACE_BEGIN
{
switch (layout)
{
case TensorLayout::x: sm << "x"; break;
case TensorLayout::chw: sm << "chw"; break;
case TensorLayout::Chw8c: sm << "Chw8c"; break;
case TensorLayout::Chw16c: sm << "Chw16c"; break;
case TensorLayout::Chw32c: sm << "Chw32c"; break;
case TensorLayout::oihw: sm << "oihw"; break;
case TensorLayout::OIhw8i8o: sm << "OIhw8i8o"; break;
case TensorLayout::OIhw16i16o: sm << "OIhw16i16o"; break;
case TensorLayout::OIhw2o8i8o2i: sm << "OIhw2o8i8o2i"; break;
case TensorLayout::OIhw8i16o2i: sm << "OIhw8i16o2i"; break;
case TensorLayout::OIhw2o16i16o2i: sm << "OIhw2o16i16o2i"; break;
case TensorLayout::IOhw8i8o: sm << "IOhw8i8o"; break;
case TensorLayout::IOhw16i16o: sm << "IOhw16i16o"; break;
case TensorLayout::hwc: sm << "hwc"; break;
case TensorLayout::ohwi: sm << "ohwi"; break;
default: sm << "?"; break;
case TensorLayout::x: sm << "x"; break;
case TensorLayout::chw: sm << "chw"; break;
case TensorLayout::Chw8c: sm << "Chw8c"; break;
case TensorLayout::Chw16c: sm << "Chw16c"; break;
case TensorLayout::oihw: sm << "oihw"; break;
case TensorLayout::OIhw8i8o: sm << "OIhw8i8o"; break;
case TensorLayout::OIhw16i16o: sm << "OIhw16i16o"; break;
case TensorLayout::OIhw2o8i8o2i: sm << "OIhw2o8i8o2i"; break;
case TensorLayout::OIhw8i16o2i: sm << "OIhw8i16o2i"; break;
case TensorLayout::IOhw8i8o: sm << "IOhw8i8o"; break;
case TensorLayout::IOhw16i16o: sm << "IOhw16i16o"; break;
case TensorLayout::hwc: sm << "hwc"; break;
case TensorLayout::ohwi: sm << "ohwi"; break;
default: sm << "?"; break;
}
return sm;
+12 -13
View File
@@ -44,19 +44,18 @@ OIDN_NAMESPACE_BEGIN
void reorderWeight(Tensor& src, int srcBeginI, int srcI, Tensor& dst, int dstBeginI, int dstI)
{
bool ok =
tryReorderWeight<half, half, TensorLayout::oihw, TensorLayout::oihw> (src, srcBeginI, srcI, dst, dstBeginI, dstI) ||
tryReorderWeight<half, float, TensorLayout::oihw, TensorLayout::oihw> (src, srcBeginI, srcI, dst, dstBeginI, dstI) ||
tryReorderWeight<half, half, TensorLayout::oihw, TensorLayout::OIhw8i8o> (src, srcBeginI, srcI, dst, dstBeginI, dstI) ||
tryReorderWeight<half, float, TensorLayout::oihw, TensorLayout::OIhw8i8o> (src, srcBeginI, srcI, dst, dstBeginI, dstI) ||
tryReorderWeight<half, half, TensorLayout::oihw, TensorLayout::OIhw16i16o> (src, srcBeginI, srcI, dst, dstBeginI, dstI) ||
tryReorderWeight<half, float, TensorLayout::oihw, TensorLayout::OIhw16i16o> (src, srcBeginI, srcI, dst, dstBeginI, dstI) ||
tryReorderWeight<half, half, TensorLayout::oihw, TensorLayout::OIhw2o8i8o2i> (src, srcBeginI, srcI, dst, dstBeginI, dstI) ||
tryReorderWeight<half, half, TensorLayout::oihw, TensorLayout::OIhw8i16o2i> (src, srcBeginI, srcI, dst, dstBeginI, dstI) ||
tryReorderWeight<half, half, TensorLayout::oihw, TensorLayout::OIhw2o16i16o2i>(src, srcBeginI, srcI, dst, dstBeginI, dstI) ||
tryReorderWeight<half, float, TensorLayout::oihw, TensorLayout::IOhw8i8o> (src, srcBeginI, srcI, dst, dstBeginI, dstI) ||
tryReorderWeight<half, float, TensorLayout::oihw, TensorLayout::IOhw16i16o> (src, srcBeginI, srcI, dst, dstBeginI, dstI) ||
tryReorderWeight<half, half, TensorLayout::oihw, TensorLayout::ohwi> (src, srcBeginI, srcI, dst, dstBeginI, dstI) ||
tryReorderWeight<half, float, TensorLayout::oihw, TensorLayout::ohwi> (src, srcBeginI, srcI, dst, dstBeginI, dstI);
tryReorderWeight<half, half, TensorLayout::oihw, TensorLayout::oihw> (src, srcBeginI, srcI, dst, dstBeginI, dstI) ||
tryReorderWeight<half, float, TensorLayout::oihw, TensorLayout::oihw> (src, srcBeginI, srcI, dst, dstBeginI, dstI) ||
tryReorderWeight<half, half, TensorLayout::oihw, TensorLayout::OIhw8i8o> (src, srcBeginI, srcI, dst, dstBeginI, dstI) ||
tryReorderWeight<half, float, TensorLayout::oihw, TensorLayout::OIhw8i8o> (src, srcBeginI, srcI, dst, dstBeginI, dstI) ||
tryReorderWeight<half, half, TensorLayout::oihw, TensorLayout::OIhw16i16o> (src, srcBeginI, srcI, dst, dstBeginI, dstI) ||
tryReorderWeight<half, float, TensorLayout::oihw, TensorLayout::OIhw16i16o> (src, srcBeginI, srcI, dst, dstBeginI, dstI) ||
tryReorderWeight<half, half, TensorLayout::oihw, TensorLayout::OIhw2o8i8o2i>(src, srcBeginI, srcI, dst, dstBeginI, dstI) ||
tryReorderWeight<half, half, TensorLayout::oihw, TensorLayout::OIhw8i16o2i> (src, srcBeginI, srcI, dst, dstBeginI, dstI) ||
tryReorderWeight<half, float, TensorLayout::oihw, TensorLayout::IOhw8i8o> (src, srcBeginI, srcI, dst, dstBeginI, dstI) ||
tryReorderWeight<half, float, TensorLayout::oihw, TensorLayout::IOhw16i16o> (src, srcBeginI, srcI, dst, dstBeginI, dstI) ||
tryReorderWeight<half, half, TensorLayout::oihw, TensorLayout::ohwi> (src, srcBeginI, srcI, dst, dstBeginI, dstI) ||
tryReorderWeight<half, float, TensorLayout::oihw, TensorLayout::ohwi> (src, srcBeginI, srcI, dst, dstBeginI, dstI);
if (!ok)
throw std::logic_error("unsupported weight layout or data type");
+24 -24
View File
@@ -469,27 +469,27 @@ OIDN_NAMESPACE_BEGIN
{
auto x = graph->addConv("enc_conv0", inputProcess, Activation::ReLU);
auto pool1 = x = graph->addConvPool("enc_conv1", x, Activation::ReLU);
auto pool1 = x = graph->addConv("enc_conv1", x, Activation::ReLU, PostOp::Pool);
auto pool2 = x = graph->addConvPool("enc_conv2", x, Activation::ReLU);
auto pool2 = x = graph->addConv("enc_conv2", x, Activation::ReLU, PostOp::Pool);
auto pool3 = x = graph->addConvPool("enc_conv3", x, Activation::ReLU);
auto pool3 = x = graph->addConv("enc_conv3", x, Activation::ReLU, PostOp::Pool);
auto pool4 = x = graph->addConvPool("enc_conv4", x, Activation::ReLU);
auto pool4 = x = graph->addConv("enc_conv4", x, Activation::ReLU, PostOp::Pool);
x = graph->addConv("enc_conv5a", pool4, Activation::ReLU);
x = graph->addConv("enc_conv5b", x, Activation::ReLU);
x = graph->addConv("enc_conv5b", x, Activation::ReLU, PostOp::Upsample);
x = graph->addUpsampleConcatConv("dec_conv4a", x, pool3, Activation::ReLU);
x = graph->addConv("dec_conv4b", x, Activation::ReLU);
x = graph->addConcatConv("dec_conv4a", x, pool3, Activation::ReLU);
x = graph->addConv("dec_conv4b", x, Activation::ReLU, PostOp::Upsample);
x = graph->addUpsampleConcatConv("dec_conv3a", x, pool2, Activation::ReLU);
x = graph->addConv("dec_conv3b", x, Activation::ReLU);
x = graph->addConcatConv("dec_conv3a", x, pool2, Activation::ReLU);
x = graph->addConv("dec_conv3b", x, Activation::ReLU, PostOp::Upsample);
x = graph->addUpsampleConcatConv("dec_conv2a", x, pool1, Activation::ReLU);
x = graph->addConv("dec_conv2b", x, Activation::ReLU);
x = graph->addConcatConv("dec_conv2a", x, pool1, Activation::ReLU);
x = graph->addConv("dec_conv2b", x, Activation::ReLU, PostOp::Upsample);
x = graph->addUpsampleConcatConv("dec_conv1a", x, inputProcess, Activation::ReLU);
x = graph->addConcatConv("dec_conv1a", x, inputProcess, Activation::ReLU);
x = graph->addConv("dec_conv1b", x, Activation::ReLU);
x = graph->addConv("dec_conv0", x, Activation::ReLU);
@@ -500,30 +500,30 @@ OIDN_NAMESPACE_BEGIN
Ref<Op> UNetFilter::addUNetLarge(const Ref<Graph>& graph, const Ref<Op>& inputProcess)
{
auto x = graph->addConv("enc_conv1a", inputProcess, Activation::ReLU);
auto pool1 = x = graph->addConvPool("enc_conv1b", x, Activation::ReLU);
auto pool1 = x = graph->addConv("enc_conv1b", x, Activation::ReLU, PostOp::Pool);
x = graph->addConv("enc_conv2a", x, Activation::ReLU);
auto pool2 = x = graph->addConvPool("enc_conv2b", x, Activation::ReLU);
auto pool2 = x = graph->addConv("enc_conv2b", x, Activation::ReLU, PostOp::Pool);
x = graph->addConv("enc_conv3a", x, Activation::ReLU);
auto pool3 = x = graph->addConvPool("enc_conv3b", x, Activation::ReLU);
auto pool3 = x = graph->addConv("enc_conv3b", x, Activation::ReLU, PostOp::Pool);
x = graph->addConv("enc_conv4a", x, Activation::ReLU);
auto pool4 = x = graph->addConvPool("enc_conv4b", x, Activation::ReLU);
auto pool4 = x = graph->addConv("enc_conv4b", x, Activation::ReLU, PostOp::Pool);
x = graph->addConv("enc_conv5a", pool4, Activation::ReLU);
x = graph->addConv("enc_conv5b", x, Activation::ReLU);
x = graph->addConv("enc_conv5b", x, Activation::ReLU, PostOp::Upsample);
x = graph->addUpsampleConcatConv("dec_conv4a", x, pool3, Activation::ReLU);
x = graph->addConv("dec_conv4b", x, Activation::ReLU);
x = graph->addConcatConv("dec_conv4a", x, pool3, Activation::ReLU);
x = graph->addConv("dec_conv4b", x, Activation::ReLU, PostOp::Upsample);
x = graph->addUpsampleConcatConv("dec_conv3a", x, pool2, Activation::ReLU);
x = graph->addConv("dec_conv3b", x, Activation::ReLU);
x = graph->addConcatConv("dec_conv3a", x, pool2, Activation::ReLU);
x = graph->addConv("dec_conv3b", x, Activation::ReLU, PostOp::Upsample);
x = graph->addUpsampleConcatConv("dec_conv2a", x, pool1, Activation::ReLU);
x = graph->addConv("dec_conv2b", x, Activation::ReLU);
x = graph->addConcatConv("dec_conv2a", x, pool1, Activation::ReLU);
x = graph->addConv("dec_conv2b", x, Activation::ReLU, PostOp::Upsample);
x = graph->addUpsampleConcatConv("dec_conv1a", x, inputProcess, Activation::ReLU);
x = graph->addConcatConv("dec_conv1a", x, inputProcess, Activation::ReLU);
x = graph->addConv("dec_conv1b", x, Activation::ReLU);
x = graph->addConv("dec_conv1c", x, Activation::ReLU);
+4 -123
View File
@@ -19,22 +19,6 @@ namespace math {
template<typename U>
oidn_host_device_inline vec2(vec2<U> v) : x(v.x), y(v.y) {}
#if defined(OIDN_COMPILE_SYCL)
template<typename U = T, enable_if_t<esimd::detail::is_simd_type_v<U>, bool> = true>
oidn_inline void merge(const vec2& val, simd_mask<U::length> mask)
{
x.merge(val.x, mask);
y.merge(val.y, mask);
}
template<typename U = T, enable_if_t<esimd::detail::is_simd_type_v<U>, bool> = true>
oidn_inline void merge(typename U::element_type val, simd_mask<U::length> mask)
{
x.merge(val, mask);
y.merge(val, mask);
}
#endif
};
template<typename T>
@@ -48,24 +32,6 @@ namespace math {
template<typename U>
oidn_host_device_inline vec3(vec3<U> v) : x(v.x), y(v.y), z(v.z) {}
#if defined(OIDN_COMPILE_SYCL)
template<typename U = T, enable_if_t<esimd::detail::is_simd_type_v<U>, bool> = true>
oidn_inline void merge(const vec3& val, simd_mask<U::length> mask)
{
x.merge(val.x, mask);
y.merge(val.y, mask);
z.merge(val.z, mask);
}
template<typename U = T, enable_if_t<esimd::detail::is_simd_type_v<U>, bool> = true>
oidn_inline void merge(typename U::element_type val, simd_mask<U::length> mask)
{
x.merge(val, mask);
y.merge(val, mask);
z.merge(val, mask);
}
#endif
};
using vec2f = vec2<float>;
@@ -116,10 +82,10 @@ namespace math {
return vec3<T>(f(v.x), f(v.y), f(v.z)); \
}
define_vec_unary_func(pow)
define_vec_unary_func(powr)
define_vec_unary_func(log)
define_vec_unary_func(log2)
define_vec_unary_func(exp)
define_vec_unary_func(exp2)
define_vec_unary_func(nan_to_zero)
#undef define_vec_unary_func
@@ -152,7 +118,6 @@ namespace math {
define_vec_binary_func(min)
define_vec_binary_func(max)
define_vec_binary_func(pow)
#undef define_vec_binary_func
@@ -174,99 +139,15 @@ namespace math {
template<typename T>
oidn_host_device_inline vec2<T> clamp(vec2<T> v, T minVal, T maxVal)
{
return vec2<T>(clamp(v.x, minVal, maxVal),
clamp(v.y, minVal, maxVal));
return vec2<T>(clamp(v.x, minVal, maxVal), clamp(v.y, minVal, maxVal));
}
template<typename T>
oidn_host_device_inline vec3<T> clamp(vec3<T> v, T minVal, T maxVal)
{
return vec3<T>(clamp(v.x, minVal, maxVal),
clamp(v.y, minVal, maxVal),
clamp(v.z, minVal, maxVal));
return vec3<T>(clamp(v.x, minVal, maxVal), clamp(v.y, minVal, maxVal), clamp(v.z, minVal, maxVal));
}
#if defined(OIDN_COMPILE_SYCL)
#define define_vec_simd_binary_op(name, op) \
template<typename T, int N> \
oidn_inline vec2<simd<T, N>> name(vec2<simd<T, N>> a, vec2<simd<T, N>> b) { \
return vec2<simd<T, N>>(a.x op b.x, a.y op b.y); \
} \
template<typename T, int N> \
oidn_inline vec2<simd<T, N>> name(vec2<simd<T, N>> a, T b) { \
return vec2<simd<T, N>>(a.x op b, a.y op b); \
} \
template<typename T, int N> \
oidn_inline vec2<simd<T, N>> name(T a, vec2<simd<T, N>> b) { \
return vec2<simd<T, N>>(a op b.x, a op b.y); \
} \
template<typename T, int N> \
oidn_inline vec3<simd<T, N>> name(vec3<simd<T, N>> a, vec3<simd<T, N>> b) { \
return {a.x op b.x, a.y op b.y, a.z op b.z}; \
} \
template<typename T, int N> \
oidn_inline vec3<simd<T, N>> name(vec3<simd<T, N>> a, T b) { \
return {a.x op b, a.y op b, a.z op b}; \
} \
template<typename T, int N> \
oidn_inline vec3<simd<T, N>> name(T a, vec3<simd<T, N>> b) { \
return {a op b.x, a op b.y, a op b.z}; \
}
define_vec_simd_binary_op(operator+, +)
define_vec_simd_binary_op(operator-, -)
define_vec_simd_binary_op(operator*, *)
define_vec_simd_binary_op(operator/, /)
#undef define_vec_simd_binary_op
#define define_vec_simd_binary_func(f) \
template<typename T, int N> \
oidn_inline vec2<simd<T, N>> f(vec2<simd<T, N>> a, vec2<simd<T, N>> b) { \
return vec2<simd<T, N>>(f(a.x, b.x), f(a.y, b.y)); \
} \
template<typename T, int N> \
oidn_inline vec2<simd<T, N>> f(vec2<simd<T, N>> a, T b) { \
return vec2<simd<T, N>>(f(a.x, b), f(a.y, b)); \
} \
template<typename T, int N> \
oidn_inline vec2<simd<T, N>> f(T a, vec2<simd<T, N>> b) { \
return vec2<simd<T, N>>(f(a, b.x), f(a, b.y)); \
} \
template<typename T, int N> \
oidn_inline vec3<simd<T, N>> f(vec3<simd<T, N>> a, vec3<simd<T, N>> b) { \
return {f(a.x, b.x), f(a.y, b.y), f(a.z, b.z)}; \
} \
template<typename T, int N> \
oidn_inline vec3<simd<T, N>> f(vec3<simd<T, N>> a, T b) { \
return {f(a.x, b), f(a.y, b), f(a.z, b)}; \
} \
template<typename T, int N> \
oidn_inline vec3<simd<T, N>> f(T a, vec3<simd<T, N>> b) { \
return {f(a, b.x), f(a, b.y), f(a, b.z)}; \
}
define_vec_simd_binary_func(min)
define_vec_simd_binary_func(max)
#undef define_vec_simd_binary_func
template<typename T, int N>
oidn_inline vec2<simd<T, N>> clamp(vec2<simd<T, N>> v, T minVal, T maxVal)
{
return {clamp(v.x, minVal, maxVal),
clamp(v.y, minVal, maxVal)};
}
template<typename T, int N>
oidn_inline vec3<simd<T, N>> clamp(vec3<simd<T, N>> v, T minVal, T maxVal)
{
return {clamp(v.x, minVal, maxVal),
clamp(v.y, minVal, maxVal),
clamp(v.z, minVal, maxVal)};
}
#endif
} // namespace math
using math::vec2;
+19 -31
View File
@@ -42,38 +42,31 @@ if(OIDN_DEVICE_CUDA)
list(JOIN CMAKE_PREFIX_PATH "|" _prefix_path_str)
set(_oidn_cuda_cache_args
-DCMAKE_PREFIX_PATH:STRING=${_prefix_path_str}
-DCMAKE_CXX_COMPILER:FILEPATH=${_host_compiler}
-DCMAKE_TOOLCHAIN_FILE:FILEPATH=${CMAKE_TOOLCHAIN_FILE}
-DCMAKE_BUILD_TYPE:STRING=${CMAKE_BUILD_TYPE}
-DCMAKE_INSTALL_PREFIX:PATH=${CMAKE_CURRENT_BINARY_DIR}/cuda/preinstall
-DCMAKE_INSTALL_BINDIR:PATH=${CMAKE_INSTALL_BINDIR}
-DCMAKE_INSTALL_LIBDIR:PATH=${CMAKE_INSTALL_LIBDIR}
-DOIDN_ROOT_BINARY_DIR:PATH=${OIDN_ROOT_BINARY_DIR}
-DOIDN_INSTALL_RPATH_PREFIX:PATH=${CMAKE_INSTALL_PREFIX}
-DOIDN_INSTALL_DEPENDENCIES:BOOL=${OIDN_INSTALL_DEPENDENCIES}
-DOIDN_ZIP_MODE:BOOL=${OIDN_ZIP_MODE}
-DOIDN_LIBRARY_NAME:STRING=${OIDN_LIBRARY_NAME}
-DOIDN_LIBRARY_VERSIONED:STRING=${OIDN_LIBRARY_VERSIONED}
-DOIDN_API_NAMESPACE:STRING=${OIDN_API_NAMESPACE}
-DOIDN_WARN_AS_ERRORS:BOOL=${OIDN_WARN_AS_ERRORS}
-DOIDN_SANITIZER:STRING=${OIDN_SANITIZER}
-DOIDN_DEVICE_CUDA_API:STRING=${OIDN_DEVICE_CUDA_API}
)
if(DEFINED CUDAToolkit_ROOT)
list(APPEND _oidn_cuda_cache_args
-DCUDAToolkit_ROOT:PATH=${CUDAToolkit_ROOT}
)
endif()
ExternalProject_Add(OpenImageDenoise_device_cuda
PREFIX cuda
SOURCE_DIR ${PROJECT_SOURCE_DIR}/devices/cuda
BINARY_DIR cuda/build
STAMP_DIR cuda/stamp
LIST_SEPARATOR |
CMAKE_CACHE_ARGS ${_oidn_cuda_cache_args}
CMAKE_CACHE_ARGS
-DCMAKE_PREFIX_PATH:STRING=${_prefix_path_str}
-DCMAKE_CXX_COMPILER:FILEPATH=${_host_compiler}
-DCMAKE_TOOLCHAIN_FILE:FILEPATH=${CMAKE_TOOLCHAIN_FILE}
-DCMAKE_BUILD_TYPE:STRING=${CMAKE_BUILD_TYPE}
-DCMAKE_INSTALL_PREFIX:PATH=${CMAKE_CURRENT_BINARY_DIR}/cuda/preinstall
-DCMAKE_INSTALL_BINDIR:PATH=${CMAKE_INSTALL_BINDIR}
-DCMAKE_INSTALL_LIBDIR:PATH=${CMAKE_INSTALL_LIBDIR}
-DCUDAToolkit_ROOT:PATH=${CUDAToolkit_ROOT}
-DOIDN_ROOT_BINARY_DIR:PATH=${OIDN_ROOT_BINARY_DIR}
-DOIDN_INSTALL_RPATH_PREFIX:PATH=${CMAKE_INSTALL_PREFIX}
-DOIDN_INSTALL_DEPENDENCIES:BOOL=${OIDN_INSTALL_DEPENDENCIES}
-DOIDN_ZIP_MODE:BOOL=${OIDN_ZIP_MODE}
-DOIDN_LIBRARY_NAME:STRING=${OIDN_LIBRARY_NAME}
-DOIDN_LIBRARY_VERSIONED:STRING=${OIDN_LIBRARY_VERSIONED}
-DOIDN_API_NAMESPACE:STRING=${OIDN_API_NAMESPACE}
-DOIDN_WARN_AS_ERRORS:BOOL=${OIDN_WARN_AS_ERRORS}
-DOIDN_SANITIZER:STRING=${OIDN_SANITIZER}
-DOIDN_DEVICE_CUDA_API:STRING=${OIDN_DEVICE_CUDA_API}
BUILD_ALWAYS TRUE
DEPENDS
OpenImageDenoise_core
@@ -105,13 +98,8 @@ if(OIDN_DEVICE_HIP)
find_path(ROCM_PATH
NAMES
bin/hipconfig
bin/hipconfig.bat
bin/hipconfig.exe
bin/hipconfig.pl
HINTS
$ENV{ROCM_PATH}
$ENV{HIP_PATH}
$ENV{HIP_PATH_64}
PATHS
${_rocm_search_paths}
NO_DEFAULT_PATH
+40 -81
View File
@@ -16,7 +16,15 @@ if(NOT TBB_FOUND)
find_package(TBB ${OIDN_TBB_MIN_VERSION} REQUIRED tbb)
endif()
if(APPLE AND OIDN_ARCH STREQUAL "ARM64")
if(OIDN_ARCH STREQUAL "X64")
# oneDNN
option(OIDN_DEVICE_CPU_ONEDNN "Enable oneDNN for CPU device." OFF)
mark_as_advanced(OIDN_DEVICE_CPU_ONEDNN)
if(OIDN_DEVICE_CPU_ONEDNN)
set(OIDN_DNNL ON)
include(oidn_dnnl)
endif()
elseif(APPLE AND OIDN_ARCH STREQUAL "ARM64")
# BNNS
option(OIDN_DEVICE_CPU_BNNS "Enable BNNS for CPU device." ON)
mark_as_advanced(OIDN_DEVICE_CPU_BNNS)
@@ -28,15 +36,11 @@ endif()
# ISPC
if(OIDN_ARCH STREQUAL "X64")
if(APPLE)
set(OIDN_ISPC_TARGET_LIST sse4.1-i32x8 avx2-i32x8 avx512skx-x16)
else()
set(OIDN_ISPC_TARGET_LIST sse4.1-i32x8 avx2-i32x8 avx512skx-x16 avx512gnr-x16)
endif()
set(OIDN_ISPC_TARGET_LIST sse4.1-i32x8;avx2-i32x8;avx512skx-i32x16)
elseif(OIDN_ARCH STREQUAL "ARM64")
set(OIDN_ISPC_TARGET_LIST neon-i32x8)
endif()
set(ISPC_ADDRESSING 64)
set(OIDN_ISPC_ADDRESSING 64)
include(oidn_ispc)
set(OIDN_CPU_SOURCES
@@ -63,28 +67,36 @@ set(OIDN_CPU_SOURCES
tasking.cpp
)
set(OIDN_ISPC_SOURCES_COMMON
set(OIDN_CPU_SOURCES_ISPC
cpu_autoexposure.ispc
cpu_image_copy.ispc
cpu_input_process.isph
cpu_input_process_f32.ispc
cpu_output_process.isph
cpu_output_process_f32.ispc
cpu_pool.isph
cpu_pool_f32.ispc
cpu_upsample.isph
cpu_upsample_f32.ispc
cpu_input_process.ispc
cpu_output_process.ispc
cpu_pool.ispc
cpu_upsample.ispc
color.isph
color.ispc
image_accessor.isph
math.isph
platform.isph
#platform.ispc
tensor_accessor.isph
tile.isph
vec.isph
)
if(OIDN_BNNS)
if(OIDN_DNNL)
list(APPEND OIDN_CPU_SOURCES
dnnl/dnnl_conv.h
dnnl/dnnl_conv.cpp
dnnl/dnnl_common.h
dnnl/dnnl_common.cpp
dnnl/dnnl_engine.h
dnnl/dnnl_engine.cpp
dnnl/dnnl_tensor.h
dnnl/dnnl_tensor.cpp
)
elseif(OIDN_BNNS)
list(APPEND OIDN_CPU_SOURCES
bnns/bnns_common.h
bnns/bnns_common.cpp
@@ -101,15 +113,10 @@ else()
cpu_conv.cpp
)
if(OIDN_ARCH STREQUAL "X64" AND NOT APPLE)
list(APPEND OIDN_CPU_SOURCES
cpu_conv_amx.h
cpu_conv_amx.cpp
)
endif()
list(APPEND OIDN_ISPC_SOURCES_COMMON
list(APPEND OIDN_CPU_SOURCES_ISPC
cpu_conv.ispc
cpu_conv_compute.isph
cpu_conv_compute_block.isph
)
endif()
@@ -122,69 +129,21 @@ if(OIDN_LIBRARY_VERSIONED)
set_target_properties(OpenImageDenoise_device_cpu PROPERTIES VERSION ${PROJECT_VERSION})
endif()
if(OIDN_BNNS)
if(OIDN_DNNL)
target_compile_definitions(OpenImageDenoise_device_cpu PRIVATE OIDN_DNNL)
ispc_add_definitions(-DOIDN_DNNL)
target_link_libraries(OpenImageDenoise_device_cpu PRIVATE dnnl)
elseif(OIDN_BNNS)
target_compile_definitions(OpenImageDenoise_device_cpu PRIVATE OIDN_BNNS)
ispc_add_definitions(-DOIDN_BNNS)
target_link_libraries(OpenImageDenoise_device_cpu PRIVATE "-framework Accelerate")
endif()
ispc_target_add_sources(OpenImageDenoise_device_cpu ${OIDN_CPU_SOURCES_ISPC})
if(OIDN_ARCH STREQUAL "X64")
ispc_target_add_sources(OpenImageDenoise_device_cpu
platform.ispc
TARGET_LIST sse2-i32x4 ${OIDN_ISPC_TARGET_LIST} # SSE2 added just for ISA detection
)
else()
list(APPEND OIDN_ISPC_SOURCES_COMMON platform.ispc)
endif()
ispc_target_add_sources(OpenImageDenoise_device_cpu
${OIDN_ISPC_SOURCES_COMMON}
TARGET_LIST ${OIDN_ISPC_TARGET_LIST}
)
if(OIDN_ARCH STREQUAL "X64" AND NOT APPLE)
ispc_target_add_sources(OpenImageDenoise_device_cpu
cpu_input_process_f16.ispc
cpu_output_process_f16.ispc
TARGET_LIST avx512gnr-x16
)
ispc_target_add_sources(OpenImageDenoise_device_cpu
cpu_conv_amx.ispc
TARGET_LIST avx512gnr-x16
COMPILE_OPTIONS --enable-llvm-intrinsics # AMX requires LLVM intrinsics
)
ispc_target_add_sources(OpenImageDenoise_device_cpu
cpu_pool_f16.ispc
cpu_upsample_f16.ispc
TARGET_LIST avx512gnr-x32
)
elseif(OIDN_ARCH STREQUAL "ARM64")
ispc_target_add_sources(OpenImageDenoise_device_cpu
cpu_input_process_f16.ispc
cpu_output_process_f16.ispc
TARGET_LIST neon-i32x8
)
ispc_target_add_sources(OpenImageDenoise_device_cpu
cpu_pool_f16.ispc
cpu_upsample_f16.ispc
TARGET_LIST neon-i16x16
)
else()
ispc_target_add_sources(OpenImageDenoise_device_cpu
cpu_input_process_f16.ispc
cpu_output_process_f16.ispc
TARGET_LIST avx512skx-x16
)
ispc_target_add_sources(OpenImageDenoise_device_cpu
cpu_pool_f16.ispc
cpu_upsample_f16.ispc
TARGET_LIST avx512skx-x32
)
set(OIDN_ISPC_TARGET_LIST sse2-i32x4;${OIDN_ISPC_TARGET_LIST}) # required just for ISA detection
endif()
ispc_target_add_sources(OpenImageDenoise_device_cpu platform.ispc)
target_link_libraries(OpenImageDenoise_device_cpu PRIVATE OpenImageDenoise_core ${CMAKE_THREAD_LIBS_INIT} TBB::tbb)
+23 -76
View File
@@ -2,11 +2,8 @@
// SPDX-License-Identifier: Apache-2.0
#include "cpu_common.h"
#if !defined(OIDN_BNNS)
#if !defined(OIDN_DNNL) && !defined(OIDN_BNNS)
#include "cpu_conv_ispc.h"
#if defined(OIDN_ARCH_X64) && !defined(__APPLE__)
#include "cpu_conv_amx_ispc.h"
#endif
#endif
OIDN_NAMESPACE_BEGIN
@@ -14,14 +11,14 @@ OIDN_NAMESPACE_BEGIN
Image::operator ispc::ImageAccessor()
{
ispc::ImageAccessor acc;
acc.ptr = reinterpret_cast<int8_t*>(ptr);
acc.ptr = reinterpret_cast<uint8_t*>(ptr);
acc.hByteStride = hByteStride;
acc.wByteStride = wByteStride;
switch (getDataType())
{
case DataType::Undefined:
acc.dataType = ispc::DataType_Undefined;
case DataType::Void:
acc.dataType = ispc::DataType_Void;
break;
/*
case DataType::UInt8:
@@ -47,67 +44,45 @@ OIDN_NAMESPACE_BEGIN
return acc;
}
#if !defined(OIDN_BNNS)
Tensor::operator ispc::TensorAccessor3D_ChwBc()
Tensor::operator ispc::TensorAccessor3D()
{
if (getRank() != 3 || (layout != TensorLayout::Chw8c &&
layout != TensorLayout::Chw16c &&
layout != TensorLayout::Chw32c))
if (getRank() != 3 || layout == TensorLayout::hwc)
throw std::logic_error("incompatible tensor accessor");
ispc::TensorAccessor3D_ChwBc acc;
acc.ptr = static_cast<int8_t*>(getPtr());
ispc::TensorAccessor3D acc;
acc.ptr = static_cast<uint8_t*>(getPtr());
acc.C = getPaddedC();
acc.H = getH();
acc.W = getW();
const int B = getTensorLayoutInfo(layout).blockC;
const uint32_t cByteStride = uint32_t(getDataTypeSize(dataType));
const uint32_t wByteStride = uint32_t(B) * cByteStride;
acc.hByteStride = uint32_t(getW()) * wByteStride;
acc.CByteStride = round_up(uint32_t(getH()) * acc.hByteStride, TensorLayoutTraitsChwBc::CByteAlignment);
const size_t cByteStride = getDataTypeSize(dataType);
const size_t wByteStride = B * cByteStride;
acc.hByteStride = size_t(acc.W) * wByteStride;
acc.CByteStride = size_t(acc.H) * acc.hByteStride;
return acc;
}
#else
Tensor::operator ispc::TensorAccessor3D_chw()
{
if (getRank() != 3 || layout != TensorLayout::chw)
throw std::logic_error("incompatible tensor accessor");
ispc::TensorAccessor3D_chw acc;
acc.ptr = static_cast<int8_t*>(getPtr());
acc.C = getPaddedC();
acc.H = getH();
acc.W = getW();
const uint32_t wByteStride = uint32_t(getDataTypeSize(dataType));
acc.hByteStride = uint32_t(getW()) * wByteStride;
acc.cByteStride = uint32_t(getH()) * acc.hByteStride;
return acc;
}
#endif
#if !defined(OIDN_BNNS)
#if !defined(OIDN_DNNL) && !defined(OIDN_BNNS)
Tensor::operator ispc::TensorAccessor1D()
{
if (layout != TensorLayout::x)
throw std::logic_error("incompatible tensor accessor");
ispc::TensorAccessor1D acc;
acc.ptr = static_cast<int8_t*>(getPtr());
acc.ptr = static_cast<uint8_t*>(getPtr());
acc.X = getPaddedX();
return acc;
}
Tensor::operator ispc::TensorAccessor4D_IOhwBiBo()
Tensor::operator ispc::TensorAccessor4D()
{
if (getRank() != 4 || (layout != TensorLayout::IOhw8i8o && layout != TensorLayout::IOhw16i16o))
throw std::logic_error("incompatible tensor accessor");
ispc::TensorAccessor4D_IOhwBiBo acc;
acc.ptr = static_cast<int8_t*>(getPtr());
ispc::TensorAccessor4D acc;
acc.ptr = static_cast<uint8_t*>(getPtr());
acc.O = getPaddedO();
acc.I = getPaddedI();
@@ -115,43 +90,15 @@ OIDN_NAMESPACE_BEGIN
acc.W = getW();
const int B = getTensorLayoutInfo(layout).blockC;
const uint32_t BoByteStride = uint32_t(getDataTypeSize(dataType));
const uint32_t BiByteStride = uint32_t(B) * BoByteStride;
const uint32_t wByteStride = uint32_t(B) * BiByteStride;
acc.hByteStride = uint32_t(getW()) * wByteStride;
acc.OByteStride = uint32_t(getH()) * acc.hByteStride;
acc.IByteStride = uint32_t(getPaddedO() / B) * acc.OByteStride;
const size_t BoByteStride = getDataTypeSize(dataType);
const size_t BiByteStride = B * BoByteStride;
const size_t wByteStride = B * BiByteStride;
acc.hByteStride = size_t(acc.W) * wByteStride;
acc.OByteStride = size_t(acc.H) * acc.hByteStride;
acc.IByteStride = size_t(acc.O / B) * acc.OByteStride;
return acc;
}
#if defined(OIDN_ARCH_X64) && !defined(__APPLE__)
Tensor::operator ispc::TensorAccessor4D_OIhwPoQiRoSi()
{
if (getRank() != 4 || layout != TensorLayout::OIhw2o16i16o2i)
throw std::logic_error("incompatible tensor accessor");
ispc::TensorAccessor4D_OIhwPoQiRoSi acc;
acc.ptr = static_cast<int8_t*>(getPtr());
const int B = getTensorLayoutInfo(layout).blockC;
const int P = 2;
const int Q = 16;
const int R = 16;
const int S = 2;
const uint32_t SiByteStride = uint32_t(getDataTypeSize(dataType));
const uint32_t RoByteStride = uint32_t(S * SiByteStride);
const uint32_t QiByteStride = uint32_t(R * RoByteStride);
const uint32_t PoByteStride = uint32_t(Q * QiByteStride);
const uint32_t wByteStride = uint32_t(P * PoByteStride);
const uint32_t hByteStride = uint32_t(getW()) * wByteStride;
const uint32_t IByteStride = uint32_t(getH()) * hByteStride;
acc.OByteStride = uint32_t(getPaddedI() / B) * IByteStride;
return acc;
}
#endif
#endif
ispc::Tile toISPC(const Tile& tile)
+1 -1
View File
@@ -7,7 +7,7 @@
#include "core/tensor.h"
#include "core/tile.h"
#include "core/color.h"
#include "cpu_input_process_f32_ispc.h"
#include "cpu_input_process_ispc.h"
#include "color_ispc.h"
OIDN_NAMESPACE_BEGIN
+26 -26
View File
@@ -23,41 +23,41 @@ OIDN_NAMESPACE_BEGIN
throw std::invalid_argument("unsupported convolution bias layout/data type");
const int blockC = getTensorLayoutInfo(dstDesc.layout).blockC;
const int IC = srcDesc.getPaddedC(); // input channels
const int OC = dstDesc.getPaddedC(); // output channels
const int OH = dstDesc.getH(); // output height
const int OW = dstDesc.getW(); // output width
const int IC = srcDesc.getPaddedC();
const int OC = dstDesc.getPaddedC();
const int OH = dstDesc.getH();
const int OW = dstDesc.getW();
const int OCB = OC / blockC; // number of output channel blocks
const int OCB = OC / blockC;
blockOCB = min(OCB, ispc::CPUConvKernel_getMaxBlockOCB());
while (OCB % blockOCB != 0)
blockOCB--;
OCBB = OCB / blockOCB; // number of output channel block blocks
OCBB = OCB / blockOCB;
blockOW = ispc::CPUConvKernel_getBlockOW(blockOCB);
// Split the output width into chunks to fit into the L2 cache
// Split the output width into tiles to fit into the L2 cache
const size_t cacheSize = 512 * 1024; // FIXME: query actual size but this also works well
const int chunkOW = max(
const int tileOW = max(
static_cast<int>((cacheSize - weightDesc.getByteSize() - biasDesc.getByteSize()) /
(IC * getDataTypeSize(srcDesc.dataType) * weightDesc.getH() +
OC * getDataTypeSize(dstDesc.dataType)) - weightDesc.getW() + 1),
1);
const int maxOWC = max(OW / (2*blockOW), 1); // max number of output width chunks
OWC = min(ceil_div(OW, chunkOW), maxOWC); // number of output width chunks
const int maxOWT = max(OW / (2*blockOW), 1); // max number of OW tiles
OWT = min(ceil_div(OW, tileOW), maxOWT); // number of OW tiles
// Tweak the number of output width chunks to maximize threading efficiency
// Tweak the number of OW tiles to maximize threading efficiency
const int numThreads = engine->getNumThreads();
double bestThreadEff = 0;
for (int curOWC = OWC; curOWC < maxOWC; ++curOWC)
for (int curOWT = OWT; curOWT < maxOWT; ++curOWT)
{
const size_t N = size_t(OCBB) * OH * curOWC; // work amount
const size_t N = size_t(OCBB) * OH * curOWT; // work amount
const double threadEff = 1. - double(N % numThreads) / N;
if (threadEff > bestThreadEff)
{
OWC = curOWC;
OWT = curOWT;
bestThreadEff = threadEff;
if (threadEff >= 0.99)
break;
@@ -67,8 +67,8 @@ OIDN_NAMESPACE_BEGIN
void CPUConv::submitKernels(const Ref<CancellationToken>& ct)
{
if (!src || !dst || !weight || !bias)
throw std::logic_error("convolution argument not set");
if (!src || !dst)
throw std::logic_error("conving source/destination not set");
ispc::CPUConvKernel kernel;
kernel.src = *src;
@@ -79,23 +79,23 @@ OIDN_NAMESPACE_BEGIN
engine->submitFunc([=]
{
const int OH = kernel.dst.H; // output height
const int OW = kernel.dst.W; // output width
const size_t N = size_t(OCBB) * OH * OWC; // total number of work items
const int OH = kernel.dst.H;
const int OW = kernel.dst.W;
const size_t N = size_t(OCBB) * OH * OWT;
parallel_for(N, [&](size_t i)
{
const size_t j = i / OCBB;
const int ocbb = int(i % OCBB); // output channel block block index
const int oh = int(j % OH); // output height index
const int owc = int(j / OH); // output width chunk index
const int ocbb = int(i % OCBB);
const int oh = int(j % OH);
const int owt = int(j / OH);
constexpr int PW = 1; // KW = 3
const int owr = OWC * (blockOW - PW - 1);
const int owBegin = owc > 0 ? (owc * OW + owr) / (OWC*blockOW) * blockOW + PW : 0;
const int owEnd = owc+1 < OWC ? ((owc+1) * OW + owr) / (OWC*blockOW) * blockOW + PW : OW;
const int owr = OWT * (blockOW - PW - 1);
const int owBegin = owt > 0 ? (owt * OW + owr) / (OWT*blockOW) * blockOW + PW : 0;
const int owEnd = owt+1 < OWT ? ((owt+1) * OW + owr) / (OWT*blockOW) * blockOW + PW : OW;
ispc::CPUConvKernel_run_f32(&kernel, blockOCB, ocbb * blockOCB, oh, owBegin, owEnd);
ispc::CPUConvKernel_run(&kernel, blockOCB, ocbb * blockOCB, oh, owBegin, owEnd);
});
}, ct);
}
+2 -2
View File
@@ -19,9 +19,9 @@ OIDN_NAMESPACE_BEGIN
private:
CPUEngine* engine;
int blockOCB; // block of output channel blocks
int blockOW; // output block width
int blockOW; // block of output width
int OCBB; // number of output channel block blocks
int OWC; // number of output width chunks
int OWT; // number of output width tiles
};
OIDN_NAMESPACE_END
+71 -168
View File
@@ -5,193 +5,96 @@
struct CPUConvKernel
{
uniform TensorAccessor3D_ChwBc src;
uniform TensorAccessor4D_IOhwBiBo weight;
uniform TensorAccessor3D src;
uniform TensorAccessor4D weight;
uniform TensorAccessor1D bias;
uniform TensorAccessor3D_ChwBc dst;
uniform TensorAccessor3D dst;
uniform bool relu;
};
#define _CPUConvKernel_compute(T, blockOCB) CPUConvKernel_compute_##T##_##blockOCB
#define CPUConvKernel_compute(T, blockOCB) _CPUConvKernel_compute(T, blockOCB)
#define _CPUConvKernel_computeBlock(T, blockOCB, blockOW) CPUConvKernel_computeBlock_##T##_##blockOCB##_##blockOW
#define CPUConvKernel_computeBlock(T, blockOCB, blockOW) _CPUConvKernel_computeBlock(T, blockOCB, blockOW)
#define T float
#define blockC programCount
#define KW 3 // kerned width
#define KH 3 // kerned height
#define PW 1 // padding width on each side
#define PH 1 // padding height on each side
#define blockC programCount
template<typename T, int blockOCB, int blockOW>
inline unmasked void CPUConvKernel_computeBlock(
const uniform int8* uniform srcPtr,
uniform size_t srcHByteStride,
const uniform int8* uniform weightPtr,
const uniform int8* uniform biasPtr,
uniform int8* uniform dstPtr,
uniform size_t dstCByteStride,
uniform size_t khEnd,
uniform size_t kwBegin, uniform size_t kwEnd,
uniform bool relu)
{
varying T accumRow[blockOCB][blockOW];
if (biasPtr)
{
#pragma unroll
for (uniform size_t bocb = 0; bocb < blockOCB; ++bocb)
{
#pragma unroll
for (uniform size_t bow = 0; bow < blockOW; ++bow)
accumRow[bocb][bow] = *((const varying T* uniform)biasPtr + bocb);
}
}
else
{
#pragma unroll
for (uniform size_t bocb = 0; bocb < blockOCB; ++bocb)
{
#pragma unroll
for (uniform size_t bow = 0; bow < blockOW; ++bow)
accumRow[bocb][bow] = *((const varying T* uniform)(dstPtr + bocb * dstCByteStride) + bow);
}
}
#pragma nounroll
for (uniform size_t kh = 0; kh < khEnd; ++kh)
{
#pragma nounroll
for (uniform size_t kw = kwBegin; kw < kwEnd; ++kw)
{
#pragma unroll
for (uniform size_t i = 0; i < blockC; ++i)
{
#pragma unroll
for (uniform size_t bocb = 0; bocb < blockOCB; ++bocb)
{
const varying T weightVec =
*((const varying T* uniform)weightPtr + (bocb * KW * KH + kw) * blockC + i);
#pragma unroll
for (uniform size_t bow = 0; bow < blockOW; ++bow)
{
const varying T srcVec = *((const uniform T* uniform)srcPtr + (bow + kw - PW) * blockC + i);
accumRow[bocb][bow] += srcVec * weightVec;
}
}
}
}
srcPtr += srcHByteStride;
weightPtr += KW * blockC * blockC * sizeof(uniform T);
}
if (relu)
{
#pragma unroll
for (uniform size_t bocb = 0; bocb < blockOCB; ++bocb)
{
#pragma unroll
for (uniform size_t bow = 0; bow < blockOW; ++bow)
accumRow[bocb][bow] = max(accumRow[bocb][bow], 0);
}
}
#pragma unroll
for (uniform size_t bocb = 0; bocb < blockOCB; ++bocb)
{
#pragma unroll
for (uniform size_t bow = 0; bow < blockOW; ++bow)
*((varying T* uniform)(dstPtr + bocb * dstCByteStride) + bow) = accumRow[bocb][bow];
}
}
template<typename T, int blockOCB, int blockOW>
unmasked void CPUConvKernel_compute(const uniform CPUConvKernel* uniform self,
uniform int ocb, uniform int oh,
uniform int owBegin, uniform int owEnd)
{
const uniform int oc = ocb * blockC;
#if KH == 3 && PH == 1
const uniform int khBegin = oh > 0 ? 0 : 1;
const uniform int khEnd = oh < self->dst.H-1 ? 3 : 2;
#else
const uniform int khBegin = max(PH - oh, 0);
const uniform int khEnd = KH - max(PH + oh - (self->dst.H-1), 0);
#endif
for (uniform int ic = 0; ic < self->src.C; ic += blockC)
{
const uniform int8* uniform srcPtr = Tensor_getPtr<T>(self->src, ic, oh + khBegin - PH, owBegin);
const uniform int8* uniform weightPtr = Tensor_getPtr<T>(self->weight, oc, ic, khBegin, 0);
const uniform int8* uniform biasPtr = (ic == 0) ? Tensor_getPtr<T>(self->bias, oc) : NULL;
uniform int8* uniform dstPtr = Tensor_getPtr<T>(self->dst, oc, oh, owBegin);
const uniform bool relu = self->relu && ic == (self->src.C - blockC);
uniform int ow = owBegin; // owBegin/owEnd *must* be aligned to block boundaries
while (ow < owEnd)
{
if (ow > PW - 1 && ow + blockOW + PW - 1 < self->dst.W)
{
// Fast path (no padding, width blocking)
CPUConvKernel_computeBlock<T, blockOCB, blockOW>(
srcPtr, self->src.hByteStride,
weightPtr, biasPtr,
dstPtr, self->dst.CByteStride,
khEnd - khBegin,
0, KW,
relu);
srcPtr += blockOW * blockC * sizeof(uniform T);
dstPtr += blockOW * blockC * sizeof(uniform T);
ow += blockOW;
}
else
{
// Slow path (padding, no width blocking)
CPUConvKernel_computeBlock<T, blockOCB, 1>(
srcPtr, self->src.hByteStride,
weightPtr, biasPtr,
dstPtr, self->dst.CByteStride,
khEnd - khBegin,
#if KW == 3 && PW == 1
ow > 0 ? 0 : 1,
ow < self->dst.W-1 ? 3 : 2,
#else
max(PW - ow, 0),
KW - max(PW + ow - (self->dst.W-1), 0),
#endif
relu);
srcPtr += blockC * sizeof(uniform T);
dstPtr += blockC * sizeof(uniform T);
ow++;
}
}
}
}
// Kernel variants optimized for different ISAs
#if defined(OIDN_ARCH_AVX512)
#if defined(ISPC_TARGET_AVX512SKX) || defined(ISPC_TARGET_AVX512SPR)
#define maxBlockOCB 4
#define blockOW1 10
#define blockOW2 10
#define blockOW3 7
#define blockOW4 6
#elif defined(OIDN_ARCH_AVX2)
#elif defined(ISPC_TARGET_AVX2)
#define maxBlockOCB 4
#define blockOW1 5
#define blockOW2 5
#define blockOW3 3
#define blockOW4 3
#elif defined(OIDN_ARCH_NEON)
#elif defined(ISPC_TARGET_NEON)
#define maxBlockOCB 3
#define blockOW1 9
#define blockOW2 5
#define blockOW3 3
#elif defined(OIDN_ARCH_SSE4)
#elif defined(ISPC_TARGET_SSE4) || defined(ISPC_TARGET_SSE2)
#define maxBlockOCB 1
#define blockOW1 5
#else
#error "Unsupported architecture"
#endif
#if maxBlockOCB >= 1
#define blockOCB 1
#define blockOW 1
#include "cpu_conv_compute_block.isph"
#undef blockOW
#define blockOW blockOW1
#include "cpu_conv_compute_block.isph"
#include "cpu_conv_compute.isph"
#undef blockOW
#undef blockOCB
#endif
#if maxBlockOCB >= 2
#define blockOCB 2
#define blockOW 1
#include "cpu_conv_compute_block.isph"
#undef blockOW
#define blockOW blockOW2
#include "cpu_conv_compute_block.isph"
#include "cpu_conv_compute.isph"
#undef blockOW
#undef blockOCB
#endif
#if maxBlockOCB >= 3
#define blockOCB 3
#define blockOW 1
#include "cpu_conv_compute_block.isph"
#undef blockOW
#define blockOW blockOW3
#include "cpu_conv_compute_block.isph"
#include "cpu_conv_compute.isph"
#undef blockOW
#undef blockOCB
#endif
#if maxBlockOCB >= 4
#define blockOCB 4
#define blockOW 1
#include "cpu_conv_compute_block.isph"
#undef blockOW
#define blockOW blockOW4
#include "cpu_conv_compute_block.isph"
#include "cpu_conv_compute.isph"
#undef blockOW
#undef blockOCB
#endif
export uniform int CPUConvKernel_getMaxBlockOCB()
@@ -217,21 +120,21 @@ export uniform int CPUConvKernel_getBlockOW(uniform int blockOCB)
}
}
export void CPUConvKernel_run_f32(const uniform CPUConvKernel* uniform self,
uniform int blockOCB, uniform int ocb, uniform int oh,
uniform int owBegin, uniform int owEnd)
export void CPUConvKernel_run(const uniform CPUConvKernel* uniform self,
uniform int blockOCB, uniform int ocb, uniform int oh,
uniform int owBegin, uniform int owEnd)
{
switch (blockOCB)
{
case 1: CPUConvKernel_compute<float, 1, blockOW1>(self, ocb, oh, owBegin, owEnd); break;
case 1: CPUConvKernel_compute(T, 1)(self, ocb, oh, owBegin, owEnd); break;
#if maxBlockOCB >= 2
case 2: CPUConvKernel_compute<float, 2, blockOW2>(self, ocb, oh, owBegin, owEnd); break;
case 2: CPUConvKernel_compute(T, 2)(self, ocb, oh, owBegin, owEnd); break;
#endif
#if maxBlockOCB >= 3
case 3: CPUConvKernel_compute<float, 3, blockOW3>(self, ocb, oh, owBegin, owEnd); break;
case 3: CPUConvKernel_compute(T, 3)(self, ocb, oh, owBegin, owEnd); break;
#endif
#if maxBlockOCB >= 4
case 4: CPUConvKernel_compute<float, 4, blockOW4>(self, ocb, oh, owBegin, owEnd); break;
case 4: CPUConvKernel_compute(T, 4)(self, ocb, oh, owBegin, owEnd); break;
#endif
}
}
-158
View File
@@ -1,158 +0,0 @@
// Copyright 2025 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
#include "cpu_conv_amx.h"
#include "cpu_common.h"
#include "cpu_conv_amx_ispc.h"
OIDN_NAMESPACE_BEGIN
oidn_inline void submitCPUConvAMXKernel(CPUEngine* engine, const ispc::CPUConvAMXKernel& kernel,
const Ref<CancellationToken>& ct)
{
// Block sizes of the kernel
constexpr int blockOC = 32; // output channel block size
constexpr int blockOW = 16; // output block width before post-op
constexpr int blockOH = 2; // output block height before post-op
constexpr int chunkOH = 8; // we group output height blocks into larger chunks
engine->submitFunc([=]
{
const int OC = kernel.dst.C; // output channels
const int OH = kernel.OH; // output height before post-op
const int OW = kernel.OW; // output width before post-op
const int OCB = OC / blockOC; // number of output channel blocks
const int OHC = ceil_div(OH, chunkOH); // number of output height chunks
const int OWB = ceil_div(OW, blockOW); // number of output width blocks
const size_t N = size_t(OCB) * OHC * OWB; // total number of work items
tbb::parallel_for(tbb::blocked_range<size_t>(0, N), [&](const tbb::blocked_range<size_t>& r)
{
for (size_t i = r.begin(); i != r.end(); ++i)
{
const size_t j = i / OCB;
const int ocb = int(i % OCB); // output channel block index
const int owb = int(j % OWB); // output width block index
const int ohc = int(j / OWB); // output height chunk index
const int oc = ocb * blockOC;
const int ow = owb * blockOW;
const int ohBegin = round_up(ohc * OH / OHC, blockOH);
const int ohEnd = min(round_up((ohc + 1) * OH / OHC, blockOH), OH);
const bool isFirst = (i == r.begin());
const bool isLast = (i == r.end() - 1);
ispc::CPUConvAMXKernel_run_f16(&kernel, oc, ohBegin, ohEnd, ow, isFirst, isLast);
}
});
}, ct);
}
CPUConvAMX::CPUConvAMX(CPUEngine* engine, const ConvDesc& desc)
: Conv(desc),
engine(engine)
{
if (srcDesc.layout != TensorLayout::Chw32c || srcDesc.dataType != DataType::Float16)
throw std::invalid_argument("unsupported convolution source layout/data type");
if (weightDesc.getW() != 3 || weightDesc.getH() != 3)
throw std::invalid_argument("unsupported convolution kernel size");
if (weightDesc.layout != TensorLayout::OIhw2o16i16o2i || weightDesc.dataType != DataType::Float16)
throw std::invalid_argument("unsupported convolution weight layout/data type");
if (biasDesc.layout != TensorLayout::x || biasDesc.dataType != DataType::Float16)
throw std::invalid_argument("unsupported convolution bias layout/data type");
}
void CPUConvAMX::submitKernels(const Ref<CancellationToken>& ct)
{
if (!src || !dst || !weight || !bias)
throw std::logic_error("convolution argument not set");
ispc::CPUConvAMXKernel kernel;
ispc::CPUConvAMXKernel_init(&kernel);
kernel.src[0] = *src;
kernel.src[1] = {}; // unused
kernel.weight = *weight;
kernel.bias = *bias;
kernel.dst = *dst;
kernel.numSrcs = 1;
kernel.relu = activation == Activation::ReLU;
switch (fusion)
{
case Fusion::None:
kernel.fusion = ispc::Fusion_None;
kernel.OH = dst->getH();
kernel.OW = dst->getW();
break;
case Fusion::UpsampleSrc0:
kernel.fusion = ispc::Fusion_UpsampleSrc0;
kernel.OH = dst->getH();
kernel.OW = dst->getW();
break;
case Fusion::PoolDst:
kernel.fusion = ispc::Fusion_PoolDst;
kernel.OH = src->getH();
kernel.OW = src->getW();
break;
default:
throw std::logic_error("unsupported convolution fusion");
}
submitCPUConvAMXKernel(engine, kernel, ct);
}
CPUConcatConvAMX::CPUConcatConvAMX(CPUEngine* engine, const ConcatConvDesc& desc)
: ConcatConv(desc),
engine(engine)
{
if (src0Desc.layout != TensorLayout::Chw32c || src0Desc.dataType != DataType::Float16)
throw std::invalid_argument("unsupported convolution source layout/data type");
if (src1Desc.layout != TensorLayout::Chw32c || src1Desc.dataType != DataType::Float16)
throw std::invalid_argument("unsupported convolution source layout/data type");
if (weightDesc.getW() != 3 || weightDesc.getH() != 3)
throw std::invalid_argument("unsupported convolution kernel size");
if (weightDesc.layout != TensorLayout::OIhw2o16i16o2i || weightDesc.dataType != DataType::Float16)
throw std::invalid_argument("unsupported convolution weight layout/data type");
if (biasDesc.layout != TensorLayout::x || biasDesc.dataType != DataType::Float16)
throw std::invalid_argument("unsupported convolution bias layout/data type");
}
void CPUConcatConvAMX::submitKernels(const Ref<CancellationToken>& ct)
{
if (!src0 || !src1 || !dst || !weight || !bias)
throw std::logic_error("concat+convolution argument not set");
ispc::CPUConvAMXKernel kernel;
ispc::CPUConvAMXKernel_init(&kernel);
kernel.src[0] = *src0;
kernel.src[1] = *src1;
kernel.weight = *weight;
kernel.bias = *bias;
kernel.dst = *dst;
kernel.numSrcs = 2;
kernel.OH = src1->getH();
kernel.OW = src1->getW();
kernel.relu = activation == Activation::ReLU;
switch (fusion)
{
case Fusion::None: kernel.fusion = ispc::Fusion_None; break;
case Fusion::UpsampleSrc0: kernel.fusion = ispc::Fusion_UpsampleSrc0; break;
case Fusion::PoolDst: kernel.fusion = ispc::Fusion_PoolDst; break;
default:
throw std::logic_error("unsupported concat+convolution fusion");
}
submitCPUConvAMXKernel(engine, kernel, ct);
}
OIDN_NAMESPACE_END
-36
View File
@@ -1,36 +0,0 @@
// Copyright 2025 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
#pragma once
#include "core/conv.h"
#include "core/concat_conv.h"
#include "cpu_engine.h"
OIDN_NAMESPACE_BEGIN
class CPUConvAMX final : public Conv
{
public:
CPUConvAMX(CPUEngine* engine, const ConvDesc& desc);
Engine* getEngine() const override { return engine; }
void submitKernels(const Ref<CancellationToken>& ct) override;
private:
CPUEngine* engine;
};
class CPUConcatConvAMX final : public ConcatConv
{
public:
CPUConcatConvAMX(CPUEngine* engine, const ConcatConvDesc& desc);
Engine* getEngine() const override { return engine; }
void submitKernels(const Ref<CancellationToken>& ct) override;
private:
CPUEngine* engine;
};
OIDN_NAMESPACE_END
-424
View File
@@ -1,424 +0,0 @@
// Copyright 2025 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
#include "tensor_accessor.isph"
#include <amx.isph>
struct AMXTileConfig
{
uniform uint8 palette_id;
uniform uint8 start_row;
uniform uint8 reserved_0[14];
uniform uint16 colsb[16];
uniform uint8 rows[16];
};
struct CPUConvAMXKernel
{
uniform TensorAccessor3D_ChwBc src[2]; // one or two source tensors (optional concatenation)
uniform TensorAccessor4D_OIhwPoQiRoSi weight;
uniform TensorAccessor1D bias;
uniform TensorAccessor3D_ChwBc dst;
uniform int numSrcs; // number of source tensors to concatenate
uniform int OH; // convolution output height without post-op
uniform int OW; // convolution output width without post-op
uniform bool relu;
uniform Fusion fusion;
uniform AMXTileConfig tileConfig; // set by CPUConvAMXKernel_init
};
#define KW 3 // kernel width
#define KH 3 // kernel height
#define PW 1 // padding width on each side
#define PH 1 // padding height on each side
#define blockOW 16 // output block width (= AMX tile row count)
#define blockOH 2 // output block height
#define blockIW (blockOW + KW - 1) // input block width (with borders for the kernel)
#define blockIH (blockOH + KH - 1) // input block height (with borders for the kernel)
#define wByteStride 64 // width byte stride of tensors (= AMX-512 vector size)
#if TARGET_WIDTH != blockOW
#error "Unsupported ISPC target width"
#endif
// AMX tile mapping (7 tiles)
#define tmm_accumRow0_ocb0 ((uniform uint8)0) // accumulator row 0, output channel block 0
#define tmm_accumRow0_ocb1 ((uniform uint8)1) // accumulator row 0, output channel block 1
#define tmm_accumRow1_ocb0 ((uniform uint8)2) // accumulator row 1, output channel block 0
#define tmm_accumRow1_ocb1 ((uniform uint8)3) // accumulator row 1, output channel block 1
#define tmm_inRow0 ((uniform uint8)4) // input block row 0
#define tmm_inRow1 ((uniform uint8)5) // input block row 1
#define tmm_weight ((uniform uint8)6) // weight output channel block 0/1
// Store output channel blocks
template<typename T>
inline unmasked void storeOCB(uniform int8* uniform ptr, varying T ocb0, varying T ocb1)
{
// Use streaming stores to avoid polluting the caches
streaming_store((uniform T* uniform)ptr, ocb0);
streaming_store((uniform T* uniform)(ptr + sizeof(varying T)), ocb1);
}
// Prefetches a block of input data (including borders needed for the kernel) from a source
// tensor into L1 cache
inline unmasked void CPUConvAMXKernel_prefetchInputBlock(
const uniform TensorAccessor3D_ChwBc& src,
const uniform int8* uniform srcPtr)
{
// Iterate over the input block height
#pragma unroll
for (uniform int bih = 0; bih < blockIH; ++bih)
{
// Iterate over the input block width
#pragma unroll
for (uniform int biw = 0; biw < blockIW; ++biw)
prefetch_l1(srcPtr + biw * wByteStride);
srcPtr += src.hByteStride; // next row in the source tensor
}
}
// Copy a block of input data (including borders needed for the kernel) from a source tensor
// to a temporary buffer, handling zero-padding as needed
inline unmasked void CPUConvAMXKernel_copyInputBlock(
const uniform TensorAccessor3D_ChwBc& src,
uniform int ih, uniform int iw,
const uniform int8* uniform srcPtr,
varying int* uniform tempPtr)
{
// Iterate over the input block height
#pragma unroll
for (uniform int bih = 0; bih < blockIH; ++bih)
{
// Iterate over the input block width
#pragma unroll
for (uniform int biw = 0; biw < blockIW; ++biw)
{
// Bounds check
if (ih + bih >= 0 && ih + bih < src.H &&
iw + biw >= 0 && iw + biw < src.W)
*tempPtr = *((const varying int* uniform)(srcPtr + biw * wByteStride));
else
*tempPtr = 0; // zero-padding
++tempPtr;
}
srcPtr += src.hByteStride; // next row in the source tensor
}
}
// Copy a block of input data (including borders needed for the kernel) from a virtually
// 2x upsampled source tensor to a temporary buffer, handling zero-padding as needed
// Assumes: ih%2 == PH%2 and iw%2 == PW%2
inline unmasked void CPUConvAMXKernel_upsampleInputBlock(
const uniform TensorAccessor3D_ChwBc& src,
uniform int ic, uniform int ih, uniform int iw,
const uniform int8* uniform srcPtr,
varying int* uniform tempPtr)
{
// Iterate over the input block height
#pragma unroll
for (uniform int bih = 0; bih < blockIH; ++bih)
{
// Iterate over the input block width
#pragma unroll
for (uniform int biw = 0; biw < blockIW; ++biw)
{
// Bounds check
if (ih + bih >= 0 && ih + bih < src.H * 2 &&
iw + biw >= 0 && iw + biw < src.W * 2)
*tempPtr = *((const varying int* uniform)(srcPtr + ((biw + (PW & 1)) >> 1) * wByteStride));
else
*tempPtr = 0; // zero-padding
++tempPtr;
}
if ((bih + PH) & 1)
srcPtr += src.hByteStride; // next row in the source tensor
}
}
// Copy a block of input data (including borders needed for the kernel) from a virtually
// 2x upsampled source tensor to a temporary buffer, assuming that no padding is needed
// Assumes: ih%2 == PH%2 and iw%2 == PW%2
inline unmasked void CPUConvAMXKernel_upsampleInputInnerBlock(
const uniform TensorAccessor3D_ChwBc& src,
uniform int ic, uniform int ih, uniform int iw,
const uniform int8* uniform srcPtr,
varying int* uniform tempPtr)
{
// Iterate over the input block height
#pragma unroll
for (uniform int bih = 0; bih < blockIH; ++bih)
{
// Iterate over the input block width
#pragma unroll
for (uniform int biw = 0; biw < blockIW; ++biw)
*tempPtr++ = *((const varying int* uniform)(srcPtr + ((biw + (PW & 1)) >> 1) * wByteStride));
if ((bih + PH) & 1)
srcPtr += src.hByteStride;
}
}
inline unmasked void CPUConvAMXKernel_computeBlock_f16(
const uniform CPUConvAMXKernel* uniform self,
uniform int oc, uniform int oh, uniform int ow)
{
const uniform int blockIC = 32; // input channel block size
// Temporary buffer for the input block (including borders needed for the kernel)
// Input blocks which require padding must be copied here first from the source tensor
varying int inBuffer[blockIH * blockIW]; // int is opaque type here
// AMX accumulator tiles must be stored to temporary memory before loading them into AVX registers
// Each tile holds a 16 element wide row and only a block of 16 output channels, so multiple tiles
// are needed to hold all channels of the result
varying float accumRow0_ocb0[blockOW]; // accumulator row 0, output channels 0..15
varying float accumRow0_ocb1[blockOW]; // accumulator row 0, output channels 16..31
varying float accumRow1_ocb0[blockOW]; // accumulator row 1, output channels 0..15
varying float accumRow1_ocb1[blockOW]; // accumulator row 1, output channels 16..31
// Initialize the accumulator tiles to zero
amx_tile_zero(tmm_accumRow0_ocb0);
amx_tile_zero(tmm_accumRow0_ocb1);
amx_tile_zero(tmm_accumRow1_ocb0);
amx_tile_zero(tmm_accumRow1_ocb1);
const uniform int8* uniform weightPtr = Tensor_getPtr<float16>(self->weight, oc);
// Iterate over the source tensors
for (uniform int i = 0; i < self->numSrcs; ++i)
{
const uniform TensorAccessor3D_ChwBc& src = self->src[i];
// Iterate over the input channel blocks
for (uniform int ic = 0; ic < src.C; ic += blockIC)
{
uniform int8* uniform inPtr;
uniform size_t inRowByteStride;
const uniform int ih = oh - PH;
const uniform int iw = ow - PW;
if (i == 0 && self->fusion == Fusion_UpsampleSrc0)
{
// Upsample by 2x2 nearest-neighbor
// Get a pointer to the input block in the source tensor
// Tensor_get doesn't work with negative indices, so we need to adjust the pointer
const uniform int8* uniform srcPtr = Tensor_getPtr<float16>(src, ic, oh >> 1, ow >> 1);
srcPtr -= ((PH + 1) >> 1) * src.hByteStride; // adjust for top padding
srcPtr -= ((PW + 1) >> 1) * wByteStride; // adjust for left padding
const uniform bool isInnerBlock = (ih >= 0) && (ih + blockIH <= src.H * 2) &&
(iw >= 0) && (iw + blockIW <= src.W * 2);
if (isInnerBlock)
CPUConvAMXKernel_upsampleInputInnerBlock(src, ic, ih, iw, srcPtr, inBuffer);
else
CPUConvAMXKernel_upsampleInputBlock(src, ic, ih, iw, srcPtr, inBuffer);
inPtr = (uniform int8* uniform)&inBuffer;
inRowByteStride = blockIW * wByteStride;
}
else
{
// Get a pointer to the input block in the source tensor
// Tensor_get doesn't work with negative indices, so we need to adjust the pointer
const uniform int8* uniform srcPtr = Tensor_getPtr<float16>(src, ic, oh, ow);
srcPtr -= PH * src.hByteStride; // adjust for top padding
srcPtr -= PW * wByteStride; // adjust for left padding
// Prefetch or copy+pad the input block, depending on whether it is an inner block
const uniform bool isInnerBlock = (ih >= 0) && (ih + blockIH <= src.H) &&
(iw >= 0) && (iw + blockIW <= src.W);
if (isInnerBlock)
{
// No padding needed, just prefetch from the source tensor (fast)
CPUConvAMXKernel_prefetchInputBlock(src, srcPtr);
inPtr = srcPtr;
inRowByteStride = src.hByteStride;
}
else
{
// Padding needed, copy to the temporary buffer (slow)
CPUConvAMXKernel_copyInputBlock(src, ih, iw, srcPtr, inBuffer);
inPtr = (uniform int8* uniform)&inBuffer;
inRowByteStride = blockIW * wByteStride;
}
}
// Iterate over the kernel height
#pragma unroll
for (uniform int kh = 0; kh < KH; ++kh)
{
const uniform int8* uniform inRowPtr = inPtr; // current input block row
// Iterate over the kernel width
#pragma unroll
for (uniform int kw = 0; kw < KW; ++kw)
{
const uniform bool isLast = (kh == KH - 1) && (kw == KW - 1);
// Load weights for output channel block 0
amx_tile_load(tmm_weight, weightPtr, 64);
// Load input block row 0
amx_tile_load(tmm_inRow0, inRowPtr, 64);
// Multiply-accumulate block row 0, output channel block 0
amx_dpfp16ps(tmm_accumRow0_ocb0, tmm_inRow0, tmm_weight);
if (isLast)
amx_tile_store(tmm_accumRow0_ocb0, (uniform int8* uniform)&accumRow0_ocb0, 64);
// Load input block row 1
amx_tile_load(tmm_inRow1, inRowPtr + inRowByteStride, 64);
// Multiply-accumulate block row 1, output channel block 0
amx_dpfp16ps(tmm_accumRow1_ocb0, tmm_inRow1, tmm_weight);
if (isLast)
amx_tile_store(tmm_accumRow1_ocb0, (uniform int8* uniform)&accumRow1_ocb0, 64);
// Load weights for output channel block 1
amx_tile_load(tmm_weight, weightPtr + 1024, 64);
// Multiply-accumulate block row 0, output channel block 1
amx_dpfp16ps(tmm_accumRow0_ocb1, tmm_inRow0, tmm_weight);
if (isLast)
amx_tile_store(tmm_accumRow0_ocb1, (uniform int8* uniform)&accumRow0_ocb1, 64);
// Multiply-accumulate block row 1, output channel block 1
amx_dpfp16ps(tmm_accumRow1_ocb1, tmm_inRow1, tmm_weight);
if (isLast)
amx_tile_store(tmm_accumRow1_ocb1, (uniform int8* uniform)&accumRow1_ocb1, 64);
inRowPtr += wByteStride; // shift input block by one column
weightPtr += 2048; // next pair of weight matrix tiles
}
inPtr += inRowByteStride; // shift input block by one row
}
}
}
// Load the bias
const varying float16* uniform biasPtr =
(varying float16* uniform)Tensor_getPtr<float16>(self->bias, oc);
const varying float bias_ocb0 = biasPtr[0];
const varying float bias_ocb1 = biasPtr[1];
// Compute and store the final results to the destination tensor
const uniform int validBlockOW = min(blockOW, self->OW - ow); // check width bounds
if (self->fusion == Fusion_PoolDst)
{
// Downsample by 2x2 max pooling
uniform int8* uniform dstRow = Tensor_getPtr<float16>(self->dst, oc, oh >> 1, ow >> 1);
for (uniform int bow = 0; bow < validBlockOW; bow += 2)
{
// Load the accumulated values
varying float x00_ocb0 = accumRow0_ocb0[bow];
varying float x00_ocb1 = accumRow0_ocb1[bow];
varying float x10_ocb0 = accumRow1_ocb0[bow];
varying float x10_ocb1 = accumRow1_ocb1[bow];
varying float x01_ocb0 = accumRow0_ocb0[bow+1];
varying float x01_ocb1 = accumRow0_ocb1[bow+1];
varying float x11_ocb0 = accumRow1_ocb0[bow+1];
varying float x11_ocb1 = accumRow1_ocb1[bow+1];
// Add the bias
x00_ocb0 += bias_ocb0; x00_ocb1 += bias_ocb1;
x10_ocb0 += bias_ocb0; x10_ocb1 += bias_ocb1;
x01_ocb0 += bias_ocb0; x01_ocb1 += bias_ocb1;
x11_ocb0 += bias_ocb0; x11_ocb1 += bias_ocb1;
// Apply ReLU activation
if (self->relu)
{
x00_ocb0 = max(x00_ocb0, 0); x00_ocb1 = max(x00_ocb1, 0);
x10_ocb0 = max(x10_ocb0, 0); x10_ocb1 = max(x10_ocb1, 0);
x01_ocb0 = max(x01_ocb0, 0); x01_ocb1 = max(x01_ocb1, 0);
x11_ocb0 = max(x11_ocb0, 0); x11_ocb1 = max(x11_ocb1, 0);
}
// Max pooling over 2x2 region
varying float x_ocb0 = max(max(x00_ocb0, x01_ocb0), max(x10_ocb0, x11_ocb0));
varying float x_ocb1 = max(max(x00_ocb1, x01_ocb1), max(x10_ocb1, x11_ocb1));
// Convert the result and store to the destination tensor
storeOCB(dstRow, (float16)x_ocb0, (float16)x_ocb1);
// Next destination column
dstRow += wByteStride;
}
}
else
{
uniform int8* uniform dstRow0 = Tensor_getPtr<float16>(self->dst, oc, oh, ow);
uniform int8* uniform dstRow1 = dstRow0 + self->dst.hByteStride;
for (uniform int bow = 0; bow < validBlockOW; ++bow)
{
// Load the accumulated values
varying float x0_ocb0 = accumRow0_ocb0[bow];
varying float x0_ocb1 = accumRow0_ocb1[bow];
varying float x1_ocb0 = accumRow1_ocb0[bow];
varying float x1_ocb1 = accumRow1_ocb1[bow];
// Add the bias
x0_ocb0 += bias_ocb0; x0_ocb1 += bias_ocb1;
x1_ocb0 += bias_ocb0; x1_ocb1 += bias_ocb1;
// Apply ReLU activation
if (self->relu)
{
x0_ocb0 = max(x0_ocb0, 0); x0_ocb1 = max(x0_ocb1, 0);
x1_ocb0 = max(x1_ocb0, 0); x1_ocb1 = max(x1_ocb1, 0);
}
// Convert the result and store to the destination tensor
storeOCB(dstRow0, (float16)x0_ocb0, (float16)x0_ocb1);
if (oh + 1 < self->OH) // check height bounds
storeOCB(dstRow1, (float16)x1_ocb0, (float16)x1_ocb1);
// Next destination column
dstRow0 += wByteStride;
dstRow1 += wByteStride;
}
}
}
// Initialize the AMX tile configuration data structure
export void CPUConvAMXKernel_init(uniform CPUConvAMXKernel* uniform self)
{
memset((uniform int8* uniform)&self->tileConfig, 0, sizeof(AMXTileConfig));
self->tileConfig.palette_id = 1;
self->tileConfig.start_row = 0;
// Configure 7 tiles
for (uniform int i = 0; i < 7; ++i)
{
self->tileConfig.colsb[i] = 64; // bytes per row
self->tileConfig.rows[i] = 16;
}
}
export void CPUConvAMXKernel_run_f16(
const uniform CPUConvAMXKernel* uniform self,
uniform int oc,
uniform int ohBegin, uniform int ohEnd,
uniform int ow,
uniform bool isFirst, uniform bool isLast)
{
if (isFirst)
amx_tile_loadconfig((const uniform int8* uniform)&self->tileConfig);
for (uniform int oh = ohBegin; oh < ohEnd; oh += blockOH)
CPUConvAMXKernel_computeBlock_f16(self, oc, oh, ow);
if (isLast)
amx_tile_release();
}
+67
View File
@@ -0,0 +1,67 @@
// Copyright 2024 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
unmasked void CPUConvKernel_compute(T, blockOCB)(const uniform CPUConvKernel* uniform self,
uniform int ocb, uniform int oh,
uniform int owBegin, uniform int owEnd)
{
const uniform int oc = ocb * blockC;
#if KH == 3 && PH == 1
const uniform int khBegin = oh > 0 ? 0 : 1;
const uniform int khEnd = oh < self->dst.H-1 ? 3 : 2;
#else
const uniform int khBegin = max(PH - oh, 0);
const uniform int khEnd = KH - max(PH + oh - (self->dst.H-1), 0);
#endif
for (uniform int ic = 0; ic < self->src.C; ic += blockC)
{
const uniform uint8* uniform srcPtr = Tensor_getPtr(self->src, ic, oh + khBegin - PH, owBegin);
const uniform uint8* uniform weightPtr = Tensor_getPtr(self->weight, oc, ic, khBegin, 0);
const uniform uint8* uniform biasPtr = (ic == 0) ? Tensor_getPtr(self->bias, oc) : NULL;
uniform uint8* uniform dstPtr = Tensor_getPtr(self->dst, oc, oh, owBegin);
const uniform bool relu = self->relu && ic == (self->src.C - blockC);
uniform int ow = owBegin; // owBegin/owEnd *must* be aligned to block boundaries
while (ow < owEnd)
{
if (ow > PW - 1 && ow + blockOW + PW - 1 < self->dst.W)
{
// Fast path (no padding, width blocking)
CPUConvKernel_computeBlock(T, blockOCB, blockOW)(
srcPtr, self->src.hByteStride,
weightPtr, biasPtr,
dstPtr, self->dst.CByteStride,
khEnd - khBegin,
0, KW,
relu);
srcPtr += blockOW * blockC * sizeof(uniform T);
dstPtr += blockOW * blockC * sizeof(uniform T);
ow += blockOW;
}
else
{
// Slow path (padding, no width blocking)
CPUConvKernel_computeBlock(T, blockOCB, 1)(
srcPtr, self->src.hByteStride,
weightPtr, biasPtr,
dstPtr, self->dst.CByteStride,
khEnd - khBegin,
#if KW == 3 && PW == 1
ow > 0 ? 0 : 1,
ow < self->dst.W-1 ? 3 : 2,
#else
max(PW - ow, 0),
KW - max(PW + ow - (self->dst.W-1), 0),
#endif
relu);
srcPtr += blockC * sizeof(uniform T);
dstPtr += blockC * sizeof(uniform T);
ow++;
}
}
}
}
+85
View File
@@ -0,0 +1,85 @@
// Copyright 2024 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
inline unmasked void CPUConvKernel_computeBlock(T, blockOCB, blockOW)(
const uniform uint8* uniform srcPtr,
uniform size_t srcHByteStride,
const uniform uint8* uniform weightPtr,
const uniform uint8* uniform biasPtr,
uniform uint8* uniform dstPtr,
uniform size_t dstCByteStride,
uniform size_t khEnd,
uniform size_t kwBegin, uniform size_t kwEnd,
uniform bool relu)
{
varying T accum[blockOCB][blockOW];
if (biasPtr)
{
#pragma unroll
for (uniform size_t bocb = 0; bocb < blockOCB; ++bocb)
{
#pragma unroll
for (uniform size_t bow = 0; bow < blockOW; ++bow)
accum[bocb][bow] = *((const varying T* uniform)biasPtr + bocb);
}
}
else
{
#pragma unroll
for (uniform size_t bocb = 0; bocb < blockOCB; ++bocb)
{
#pragma unroll
for (uniform size_t bow = 0; bow < blockOW; ++bow)
accum[bocb][bow] = *((const varying T* uniform)(dstPtr + bocb * dstCByteStride) + bow);
}
}
#pragma nounroll
for (uniform size_t kh = 0; kh < khEnd; ++kh)
{
#pragma nounroll
for (uniform size_t kw = kwBegin; kw < kwEnd; ++kw)
{
#pragma unroll
for (uniform size_t i = 0; i < blockC; ++i)
{
#pragma unroll
for (uniform size_t bocb = 0; bocb < blockOCB; ++bocb)
{
const varying T weightVec =
*((const varying T* uniform)weightPtr + (bocb * KW * KH + kw) * blockC + i);
#pragma unroll
for (uniform size_t bow = 0; bow < blockOW; ++bow)
{
const varying T srcVec = *((const uniform T* uniform)srcPtr + (bow + kw - PW) * blockC + i);
accum[bocb][bow] += srcVec * weightVec;
}
}
}
}
srcPtr += srcHByteStride;
weightPtr += KW * blockC * blockC * sizeof(uniform T);
}
if (relu)
{
#pragma unroll
for (uniform size_t bocb = 0; bocb < blockOCB; ++bocb)
{
#pragma unroll
for (uniform size_t bow = 0; bow < blockOW; ++bow)
accum[bocb][bow] = max(accum[bocb][bow], 0);
}
}
#pragma unroll
for (uniform size_t bocb = 0; bocb < blockOCB; ++bocb)
{
#pragma unroll
for (uniform size_t bow = 0; bow < blockOW; ++bow)
*((varying T* uniform)(dstPtr + bocb * dstCByteStride) + bow) = accum[bocb][bow];
}
}
+36 -83
View File
@@ -4,7 +4,9 @@
#include "cpu_device.h"
#include "platform_ispc.h"
#if defined(OIDN_BNNS)
#if defined(OIDN_DNNL)
#include "dnnl/dnnl_engine.h"
#elif defined(OIDN_BNNS)
#include "bnns/bnns_engine.h"
#else
#include "cpu_engine.h"
@@ -15,15 +17,7 @@
#include <intrin.h> // __cpuid
#elif !defined(__APPLE__)
#include <cpuid.h>
#include <unistd.h>
#include <sys/syscall.h>
#endif
// AMX feature control
#define ARCH_GET_XCOMP_PERM 0x1022
#define ARCH_REQ_XCOMP_PERM 0x1023
#define XFEATURE_XTILECFG 17
#define XFEATURE_XTILEDATA 18
#endif
OIDN_NAMESPACE_BEGIN
@@ -47,23 +41,12 @@ OIDN_NAMESPACE_BEGIN
std::vector<Ref<PhysicalDevice>> CPUDevice::getPhysicalDevices()
{
CPUArch arch = getNativeArch();
CPUArch arch = getArch();
if (arch == CPUArch::Unknown)
return {};
int score = 0;
switch (arch)
{
case CPUArch::AVX512_AMXFP16:
score = (15 << 16);
break;
default:
// Prefer the CPU over some low-power integrated GPUs
score = (1 << 16) + 61;
break;
}
// Prefer the CPU over some low-power integrated GPUs
int score = (1 << 16) + 61;
return {makeRef<CPUPhysicalDevice>(score)};
}
@@ -93,22 +76,19 @@ OIDN_NAMESPACE_BEGIN
return "CPU"; // fallback
}
// Returns the native CPU architecture but this may not be what we are allowed to use (e.g. AMX)
CPUArch CPUDevice::getNativeArch()
CPUArch CPUDevice::getArch()
{
switch (ispc::getCPUArch())
{
case ispc::CPUArch_SSE4: return CPUArch::SSE4;
case ispc::CPUArch_AVX2: return CPUArch::AVX2;
case ispc::CPUArch_AVX512: return CPUArch::AVX512;
case ispc::CPUArch_AVX512_AMXFP16: return CPUArch::AVX512_AMXFP16;
case ispc::CPUArch_NEON: return CPUArch::NEON;
default: return CPUArch::Unknown;
case ispc::CPUArch_SSE4: return CPUArch::SSE41;
case ispc::CPUArch_AVX2: return CPUArch::AVX2;
case ispc::CPUArch_AVX512: return CPUArch::AVX512;
case ispc::CPUArch_NEON: return CPUArch::NEON;
default: return CPUArch::Unknown;
}
}
CPUDevice::CPUDevice(const Ref<CPUPhysicalDevice>& physicalDevice)
: physicalDevice(physicalDevice)
CPUDevice::CPUDevice()
{
systemMemorySupported = true;
managedMemorySupported = true;
@@ -120,47 +100,29 @@ OIDN_NAMESPACE_BEGIN
void CPUDevice::init()
{
// Detect the architecture only once
if (physicalDevice->arch == CPUArch::Unknown)
{
physicalDevice->arch = getNativeArch();
arch = getArch();
#if defined(OIDN_ARCH_X64)
if (physicalDevice->arch == CPUArch::AVX512_AMXFP16)
{
#if defined(_WIN32)
// Check whether AMX is enabled by Windows
const DWORD64 features = GetEnabledXStateFeatures();
if ((features & (1 << XFEATURE_XTILECFG)) == 0 ||
(features & (1 << XFEATURE_XTILEDATA)) == 0)
{
printWarning("AMX not enabled by the OS");
physicalDevice->arch = CPUArch::AVX512; // fallback to plain AVX-512
}
#elif defined(__linux__)
// We must request permission to use AMX on Linux
if (syscall(SYS_arch_prctl, ARCH_REQ_XCOMP_PERM, XFEATURE_XTILEDATA))
{
printWarning("failed to get AMX enabled by the OS");
physicalDevice->arch = CPUArch::AVX512; // fallback to plain AVX-512
}
#else
// We don't support AMX on other OSes
physicalDevice->arch = CPUArch::AVX512; // fallback to plain AVX-512
#endif
}
#endif
}
arch = physicalDevice->arch;
// Set the tensor data types and layouts based on the architecture
tensorDataType = DataType::Float32;
weightDataType = DataType::Float32;
std::unique_ptr<CPUEngine> engine;
#if defined(OIDN_BNNS)
#if defined(OIDN_DNNL)
if (arch == CPUArch::AVX512)
{
tensorLayout = TensorLayout::Chw16c;
weightLayout = TensorLayout::OIhw16i16o;
tensorBlockC = 16;
}
else
{
tensorLayout = TensorLayout::Chw8c;
weightLayout = TensorLayout::OIhw8i8o;
tensorBlockC = 8;
}
engine.reset(new DNNLEngine(this, numThreads));
#elif defined(OIDN_BNNS)
tensorLayout = TensorLayout::chw;
weightLayout = TensorLayout::oihw;
tensorBlockC = 1;
@@ -173,14 +135,6 @@ OIDN_NAMESPACE_BEGIN
weightLayout = TensorLayout::IOhw16i16o;
tensorBlockC = 16;
}
else if (arch == CPUArch::AVX512_AMXFP16)
{
tensorDataType = DataType::Float16;
weightDataType = DataType::Float16;
tensorLayout = TensorLayout::Chw32c;
weightLayout = TensorLayout::OIhw2o16i16o2i;
tensorBlockC = 32;
}
else
{
tensorLayout = TensorLayout::Chw8c;
@@ -203,13 +157,12 @@ OIDN_NAMESPACE_BEGIN
std::cout << " ISA : ";
switch (arch)
{
case CPUArch::SSE2: std::cout << "SSE2"; break;
case CPUArch::SSE4: std::cout << "SSE4.1"; break;
case CPUArch::AVX2: std::cout << "AVX2"; break;
case CPUArch::AVX512: std::cout << "AVX-512"; break;
case CPUArch::AVX512_AMXFP16: std::cout << "AVX-512 AMX-FP16"; break;
case CPUArch::NEON: std::cout << "NEON"; break;
default: std::cout << "Unknown"; break;
case CPUArch::SSE2: std::cout << "SSE2"; break;
case CPUArch::SSE41: std::cout << "SSE4.1"; break;
case CPUArch::AVX2: std::cout << "AVX2"; break;
case CPUArch::AVX512: std::cout << "AVX-512"; break;
case CPUArch::NEON: std::cout << "NEON"; break;
default: std::cout << "Unknown"; break;
}
std::cout << std::endl;
+6 -11
View File
@@ -15,35 +15,35 @@ OIDN_NAMESPACE_BEGIN
{
Unknown,
SSE2,
SSE4,
SSE41,
AVX2,
AVX512,
AVX512_AMXFP16,
NEON
};
class CPUPhysicalDevice final : public PhysicalDevice
{
public:
CPUArch arch = CPUArch::Unknown;
explicit CPUPhysicalDevice(int score);
};
class CPUDevice final : public Device
{
friend class CPUEngine;
friend class DNNLEngine;
public:
static std::vector<Ref<PhysicalDevice>> getPhysicalDevices();
static std::string getName();
static CPUArch getNativeArch();
static CPUArch getArch();
explicit CPUDevice(const Ref<CPUPhysicalDevice>& physicalDevice);
CPUDevice();
DeviceType getType() const override { return DeviceType::CPU; }
#if !defined(OIDN_DNNL)
bool needWeightAndBiasOnDevice() const override { return false; } // no need to copy
#endif
Storage getPtrStorage(const void* ptr) override;
int getInt(const std::string& name) override;
@@ -51,8 +51,6 @@ OIDN_NAMESPACE_BEGIN
void wait() override;
CPUArch getArch() const { return arch; }
protected:
void init() override;
@@ -61,9 +59,6 @@ OIDN_NAMESPACE_BEGIN
int numThreads = 0; // autodetect by default
bool setAffinity = true;
// Used only for initialization
Ref<CPUPhysicalDevice> physicalDevice;
};
OIDN_NAMESPACE_END
+3 -47
View File
@@ -1,14 +1,9 @@
// Copyright 2018 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
#include "core/conv.h"
#include "core/concat_conv.h"
#include "cpu_engine.h"
#if !defined(OIDN_BNNS)
#if !defined(OIDN_DNNL) && !defined(OIDN_BNNS)
#include "cpu_conv.h"
#if defined(OIDN_ARCH_X64) && !defined(__APPLE__)
#include "cpu_conv_amx.h"
#endif
#endif
#include "cpu_pool.h"
#include "cpu_upsample.h"
@@ -64,52 +59,13 @@ OIDN_NAMESPACE_BEGIN
observer.reset();
}
bool CPUEngine::isConvSupported(Fusion fusion) const
{
if (device->getArch() == CPUArch::AVX512_AMXFP16)
{
return fusion == Fusion::None ||
fusion == Fusion::UpsampleSrc0 ||
fusion == Fusion::PoolDst;
}
else
return fusion == Fusion::None;
}
bool CPUEngine::isConcatConvSupported(Fusion fusion) const
{
if (device->getArch() == CPUArch::AVX512_AMXFP16)
{
return fusion == Fusion::None ||
fusion == Fusion::UpsampleSrc0 ||
fusion == Fusion::PoolDst;
}
else
return fusion == Fusion::None;
}
#if !defined(OIDN_BNNS)
#if !defined(OIDN_DNNL) && !defined(OIDN_BNNS)
Ref<Conv> CPUEngine::newConv(const ConvDesc& desc)
{
#if defined(OIDN_ARCH_X64) && !defined(__APPLE__)
if (device->getArch() == CPUArch::AVX512_AMXFP16)
return makeRef<CPUConvAMX>(this, desc);
else
#endif
return makeRef<CPUConv>(this, desc);
return makeRef<CPUConv>(this, desc);
}
#endif
Ref<ConcatConv> CPUEngine::newConcatConv(const ConcatConvDesc& desc)
{
#if defined(OIDN_ARCH_X64) && !defined(__APPLE__)
if (device->getArch() == CPUArch::AVX512_AMXFP16)
return makeRef<CPUConcatConvAMX>(this, desc);
else
#endif
return makeRef<PreConcatConvCHW>(this, desc);
}
Ref<Pool> CPUEngine::newPool(const PoolDesc& desc)
{
return makeRef<CPUPool>(this, desc);
+1 -4
View File
@@ -23,12 +23,9 @@ OIDN_NAMESPACE_BEGIN
int getNumThreads() const { return device->numThreads; }
// Ops
bool isConvSupported(Fusion fusion) const override;
bool isConcatConvSupported(Fusion fusion) const override;
#if !defined(OIDN_BNNS)
#if !defined(OIDN_DNNL) && !defined(OIDN_BNNS)
Ref<Conv> newConv(const ConvDesc& desc) override;
#endif
Ref<ConcatConv> newConcatConv(const ConcatConvDesc& desc) override;
Ref<Pool> newPool(const PoolDesc& desc) override;
Ref<Upsample> newUpsample(const UpsampleDesc& desc) override;
Ref<Autoexposure> newAutoexposure(const ImageDesc& srcDesc) override;
+3 -11
View File
@@ -2,8 +2,7 @@
// SPDX-License-Identifier: Apache-2.0
#include "cpu_input_process.h"
#include "cpu_input_process_f32_ispc.h"
#include "cpu_input_process_f16_ispc.h"
#include "cpu_input_process_ispc.h"
#include "cpu_common.h"
OIDN_NAMESPACE_BEGIN
@@ -11,10 +10,7 @@ OIDN_NAMESPACE_BEGIN
CPUInputProcess::CPUInputProcess(CPUEngine* engine, const InputProcessDesc& desc)
: InputProcess(engine, desc),
engine(engine)
{
if (dstDesc.dataType != DataType::Float32 && dstDesc.dataType != DataType::Float16)
throw std::invalid_argument("unsupported input process destination data type");
}
{}
void CPUInputProcess::submitKernels(const Ref<CancellationToken>& ct)
{
@@ -32,15 +28,11 @@ OIDN_NAMESPACE_BEGIN
kernel.hdr = hdr;
kernel.snorm = snorm;
auto kernelFunc = (dst->getDataType() == DataType::Float16)
? ispc::CPUInputProcessKernel_run_f16
: ispc::CPUInputProcessKernel_run_f32;
engine->submitFunc([=]
{
parallel_for(kernel.dst.H, [&](int hDst)
{
kernelFunc(&kernel, hDst);
ispc::CPUInputProcessKernel_run(&kernel, hDst);
});
}, ct);
}
@@ -1,8 +1,6 @@
// Copyright 2018 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
#pragma once
#include "tensor_accessor.isph"
#include "image_accessor.isph"
#include "color.isph"
@@ -75,8 +73,7 @@ inline vec3f getNormal(const uniform CPUInputProcessKernel* uniform self, unifor
return value;
}
template<typename DstT>
inline void CPUInputProcessKernel_run(const uniform CPUInputProcessKernel* uniform self,
export void CPUInputProcessKernel_run(const uniform CPUInputProcessKernel* uniform self,
uniform int hDst)
{
const uniform int h = hDst - self->tile.hDstBegin;
@@ -89,7 +86,7 @@ inline void CPUInputProcessKernel_run(const uniform CPUInputProcessKernel* unifo
foreach (wDst = 0 ... self->tile.wDstBegin)
{
for (uniform int c = 0; c < self->dst.C; ++c)
Tensor_set<DstT>(self->dst, c, hDst, wDst, 0);
Tensor_set(self->dst, c, hDst, wDst, 0);
}
// Reorder
@@ -98,30 +95,30 @@ inline void CPUInputProcessKernel_run(const uniform CPUInputProcessKernel* unifo
const int wSrc = w + self->tile.wSrcBegin;
const int wDst = w + self->tile.wDstBegin;
Tensor_set3<DstT>(self->dst, 0, hDst, wDst, getInput(self, hSrc, wSrc));
Tensor_set3(self->dst, 0, hDst, wDst, getInput(self, hSrc, wSrc));
uniform int c = 3;
if (self->albedo.ptr)
{
Tensor_set3<DstT>(self->dst, 3, hDst, wDst, getAlbedo(self, hSrc, wSrc));
Tensor_set3(self->dst, 3, hDst, wDst, getAlbedo(self, hSrc, wSrc));
c += 3;
if (self->normal.ptr)
{
Tensor_set3<DstT>(self->dst, 6, hDst, wDst, getNormal(self, hSrc, wSrc));
Tensor_set3(self->dst, 6, hDst, wDst, getNormal(self, hSrc, wSrc));
c += 3;
}
}
for (; c < self->dst.C; ++c)
Tensor_set<DstT>(self->dst, c, hDst, wDst, 0);
Tensor_set(self->dst, c, hDst, wDst, 0);
}
// Zero pad
foreach (wDst = self->tile.W + self->tile.wDstBegin ... self->dst.W)
{
for (uniform int c = 0; c < self->dst.C; ++c)
Tensor_set<DstT>(self->dst, c, hDst, wDst, 0);
Tensor_set(self->dst, c, hDst, wDst, 0);
}
}
else
@@ -130,7 +127,8 @@ inline void CPUInputProcessKernel_run(const uniform CPUInputProcessKernel* unifo
foreach (wDst = 0 ... self->dst.W)
{
for (uniform int c = 0; c < self->dst.C; ++c)
Tensor_set<DstT>(self->dst, c, hDst, wDst, 0);
Tensor_set(self->dst, c, hDst, wDst, 0);
}
}
}
-10
View File
@@ -1,10 +0,0 @@
// Copyright 2025 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
#include "cpu_input_process.isph"
export void CPUInputProcessKernel_run_f16(const uniform CPUInputProcessKernel* uniform self,
uniform int hDst)
{
CPUInputProcessKernel_run<float16>(self, hDst);
}
-10
View File
@@ -1,10 +0,0 @@
// Copyright 2025 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
#include "cpu_input_process.isph"
export void CPUInputProcessKernel_run_f32(const uniform CPUInputProcessKernel* uniform self,
uniform int hDst)
{
CPUInputProcessKernel_run<float>(self, hDst);
}
+1 -1
View File
@@ -12,7 +12,7 @@ OIDN_NAMESPACE_BEGIN
Ref<Device> newDevice(const Ref<PhysicalDevice>& physicalDevice) override
{
assert(physicalDevice->type == DeviceType::CPU);
return makeRef<CPUDevice>(staticRefCast<CPUPhysicalDevice>(physicalDevice));
return makeRef<CPUDevice>();
}
};
+3 -11
View File
@@ -2,8 +2,7 @@
// SPDX-License-Identifier: Apache-2.0
#include "cpu_output_process.h"
#include "cpu_output_process_f32_ispc.h"
#include "cpu_output_process_f16_ispc.h"
#include "cpu_output_process_ispc.h"
#include "cpu_common.h"
OIDN_NAMESPACE_BEGIN
@@ -11,10 +10,7 @@ OIDN_NAMESPACE_BEGIN
CPUOutputProcess::CPUOutputProcess(CPUEngine* engine, const OutputProcessDesc& desc)
: OutputProcess(desc),
engine(engine)
{
if (srcDesc.dataType != DataType::Float32 && srcDesc.dataType != DataType::Float16)
throw std::invalid_argument("unsupported output process source data type");
}
{}
void CPUOutputProcess::submitKernels(const Ref<CancellationToken>& ct)
{
@@ -29,15 +25,11 @@ OIDN_NAMESPACE_BEGIN
kernel.hdr = hdr;
kernel.snorm = snorm;
auto kernelFunc = (src->getDataType() == DataType::Float16)
? ispc::CPUOutputProcessKernel_run_f16
: ispc::CPUOutputProcessKernel_run_f32;
engine->submitFunc([=]
{
parallel_for(kernel.tile.H, [&](int h)
{
kernelFunc(&kernel, h);
ispc::CPUOutputProcessKernel_run(&kernel, h);
});
}, ct);
}
@@ -1,8 +1,6 @@
// Copyright 2018 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
#pragma once
#include "tensor_accessor.isph"
#include "image_accessor.isph"
#include "color.isph"
@@ -25,8 +23,7 @@ struct CPUOutputProcessKernel
uniform bool snorm; // signed normalized ([-1..1])
};
template<typename SrcT>
inline void CPUOutputProcessKernel_run(const uniform CPUOutputProcessKernel* uniform self,
export void CPUOutputProcessKernel_run(const uniform CPUOutputProcessKernel* uniform self,
uniform int h)
{
const uniform int hSrc = h + self->tile.hSrcBegin;
@@ -40,7 +37,7 @@ inline void CPUOutputProcessKernel_run(const uniform CPUOutputProcessKernel* uni
const int wDst = w + self->tile.wDstBegin;
// Load
vec3f value = Tensor_get3<SrcT>(self->src, 0, hSrc, wSrc);
vec3f value = Tensor_get3(self->src, 0, hSrc, wSrc);
// The CNN output may contain negative values or even NaNs, so it must be sanitized
value = clamp(nan_to_zero(value), 0.f, pos_max);
-10
View File
@@ -1,10 +0,0 @@
// Copyright 2025 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
#include "cpu_output_process.isph"
export void CPUOutputProcessKernel_run_f16(const uniform CPUOutputProcessKernel* uniform self,
uniform int h)
{
CPUOutputProcessKernel_run<float16>(self, h);
}
-10
View File
@@ -1,10 +0,0 @@
// Copyright 2025 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
#include "cpu_output_process.isph"
export void CPUOutputProcessKernel_run_f32(const uniform CPUOutputProcessKernel* uniform self,
uniform int h)
{
CPUOutputProcessKernel_run<float>(self, h);
}
+4 -12
View File
@@ -2,8 +2,7 @@
// SPDX-License-Identifier: Apache-2.0
#include "cpu_pool.h"
#include "cpu_pool_f32_ispc.h"
#include "cpu_pool_f16_ispc.h"
#include "cpu_pool_ispc.h"
#include "cpu_common.h"
OIDN_NAMESPACE_BEGIN
@@ -12,12 +11,9 @@ OIDN_NAMESPACE_BEGIN
: Pool(desc),
engine(engine)
{
if (srcDesc.layout != TensorLayout::Chw8c &&
srcDesc.layout != TensorLayout::Chw16c &&
srcDesc.layout != TensorLayout::Chw32c)
if (srcDesc.layout != TensorLayout::Chw8c &&
srcDesc.layout != TensorLayout::Chw16c)
throw std::invalid_argument("unsupported pooling source layout");
if (srcDesc.dataType != DataType::Float32 && srcDesc.dataType != DataType::Float16)
throw std::invalid_argument("unsupported pooling source data type");
}
void CPUPool::submitKernels(const Ref<CancellationToken>& ct)
@@ -31,15 +27,11 @@ OIDN_NAMESPACE_BEGIN
kernel.src = *src;
kernel.dst = *dst;
auto kernelFunc = (src->getDataType() == DataType::Float16)
? ispc::CPUPoolKernel_run_f16
: ispc::CPUPoolKernel_run_f32;
engine->submitFunc([=]
{
parallel_for(kernel.dst.C / blockC, kernel.dst.H, [&](int cb, int h)
{
kernelFunc(&kernel, cb, h);
ispc::CPUPoolKernel_run(&kernel, cb, h);
});
}, ct);
}
+33
View File
@@ -0,0 +1,33 @@
// Copyright 2022 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
#include "tensor_accessor.isph"
struct CPUPoolKernel
{
uniform TensorAccessor3D src;
uniform TensorAccessor3D dst;
};
export void CPUPoolKernel_run(const uniform CPUPoolKernel* uniform self,
uniform int cb, uniform int h)
{
const uniform size_t H = (uniform size_t)self->dst.H;
const uniform size_t W = (uniform size_t)self->dst.W;
const uniform size_t offset = (cb*H + h) * (W*B);
uniform float* const uniform dstPtr_line = (uniform float* uniform)self->dst.ptr + offset;
uniform float* const uniform srcPtr_line0 = (uniform float* uniform)self->src.ptr + offset * 4;
uniform float* const uniform srcPtr_line1 = srcPtr_line0 + W*2*B; // next line
for (uniform size_t w = 0; w < W; ++w)
{
const float value0 = *((varying float* uniform)&srcPtr_line0[w*2*B ]);
const float value1 = *((varying float* uniform)&srcPtr_line0[w*2*B+B]);
const float value2 = *((varying float* uniform)&srcPtr_line1[w*2*B ]);
const float value3 = *((varying float* uniform)&srcPtr_line1[w*2*B+B]);
const float value = max(max(value0, value1), max(value2, value3));
streaming_store(&dstPtr_line[w*B], value);
}
}
-35
View File
@@ -1,35 +0,0 @@
// Copyright 2022 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
#pragma once
#include "tensor_accessor.isph"
struct CPUPoolKernel
{
uniform TensorAccessor3D src;
uniform TensorAccessor3D dst;
};
template<typename T>
inline void CPUPoolKernel_run(const uniform CPUPoolKernel* uniform self,
uniform int cb, uniform int h)
{
const uniform int B = getBlockC<T>();
const uniform size_t W = (uniform size_t)self->dst.W;
uniform T* const uniform dstPtr_line = (uniform T* const uniform)Tensor_getPtr<T>(self->dst, cb*B, h, 0);
uniform T* const uniform srcPtr_line0 = (uniform T* const uniform)Tensor_getPtr<T>(self->src, cb*B, h*2, 0);
uniform T* const uniform srcPtr_line1 = (uniform T* const uniform)Tensor_getPtr<T>(self->src, cb*B, h*2+1, 0);
for (uniform size_t w = 0; w < W; ++w)
{
const T value0 = *((varying T* uniform)&srcPtr_line0[w*2*B ]);
const T value1 = *((varying T* uniform)&srcPtr_line0[w*2*B+B]);
const T value2 = *((varying T* uniform)&srcPtr_line1[w*2*B ]);
const T value3 = *((varying T* uniform)&srcPtr_line1[w*2*B+B]);
const T value = max(max(value0, value1), max(value2, value3));
streaming_store(&dstPtr_line[w*B], value);
}
}
-14
View File
@@ -1,14 +0,0 @@
// Copyright 2025 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
#include "cpu_pool.isph"
#if TARGET_WIDTH * 2 != BLOCKC_BYTESIZE
#error "Unsupported ISPC target width"
#endif
export void CPUPoolKernel_run_f16(const uniform CPUPoolKernel* uniform self,
uniform int cb, uniform int h)
{
CPUPoolKernel_run<float16>(self, cb, h);
}
-14
View File
@@ -1,14 +0,0 @@
// Copyright 2025 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
#include "cpu_pool.isph"
#if TARGET_WIDTH * 4 != BLOCKC_BYTESIZE
#error "Unsupported ISPC target width"
#endif
export void CPUPoolKernel_run_f32(const uniform CPUPoolKernel* uniform self,
uniform int cb, uniform int h)
{
CPUPoolKernel_run<float>(self, cb, h);
}
+5 -14
View File
@@ -2,8 +2,7 @@
// SPDX-License-Identifier: Apache-2.0
#include "cpu_upsample.h"
#include "cpu_upsample_f32_ispc.h"
#include "cpu_upsample_f16_ispc.h"
#include "cpu_upsample_ispc.h"
#include "cpu_common.h"
OIDN_NAMESPACE_BEGIN
@@ -12,14 +11,10 @@ OIDN_NAMESPACE_BEGIN
: Upsample(desc),
engine(engine)
{
if (srcDesc.layout != TensorLayout::chw &&
srcDesc.layout != TensorLayout::Chw8c &&
srcDesc.layout != TensorLayout::Chw16c &&
srcDesc.layout != TensorLayout::Chw32c)
if (srcDesc.layout != TensorLayout::chw &&
srcDesc.layout != TensorLayout::Chw8c &&
srcDesc.layout != TensorLayout::Chw16c)
throw std::invalid_argument("unsupported upsampling source layout");
if (srcDesc.dataType != DataType::Float32 &&
(srcDesc.dataType != DataType::Float16 || srcDesc.layout == TensorLayout::chw))
throw std::invalid_argument("unsupported upsampling source data type");
}
void CPUUpsample::submitKernels(const Ref<CancellationToken>& ct)
@@ -35,15 +30,11 @@ OIDN_NAMESPACE_BEGIN
kernel.src = *src;
kernel.dst = *dst;
auto kernelFunc = (src->getDataType() == DataType::Float16)
? ispc::CPUUpsampleKernel_run_f16
: ispc::CPUUpsampleKernel_run_f32;
engine->submitFunc([=]
{
parallel_for(kernel.src.C / blockC, kernel.src.H, [&](int cb, int h)
{
kernelFunc(&kernel, cb, h);
ispc::CPUUpsampleKernel_run(&kernel, cb, h);
});
}, ct);
}
@@ -1,8 +1,6 @@
// Copyright 2018 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
#pragma once
#include "tensor_accessor.isph"
struct CPUUpsampleKernel
@@ -11,20 +9,20 @@ struct CPUUpsampleKernel
uniform TensorAccessor3D dst;
};
template<typename T>
inline void CPUUpsampleKernel_run(const uniform CPUUpsampleKernel* uniform self,
export void CPUUpsampleKernel_run(const uniform CPUUpsampleKernel* uniform self,
uniform int cb, uniform int h)
{
const uniform int B = getBlockC<T>();
const uniform size_t H = (uniform size_t)self->src.H;
const uniform size_t W = (uniform size_t)self->src.W;
uniform T* const uniform srcPtr_line = (uniform T* const uniform)Tensor_getPtr<T>(self->src, cb*B, h, 0);
uniform T* const uniform dstPtr_line0 = (uniform T* const uniform)Tensor_getPtr<T>(self->dst, cb*B, h*2, 0);
uniform T* const uniform dstPtr_line1 = (uniform T* const uniform)Tensor_getPtr<T>(self->dst, cb*B, h*2+1, 0);
const uniform size_t offset = (cb*H + h) * (W*B);
uniform float* const uniform srcPtr_line = (uniform float* uniform)self->src.ptr + offset;
uniform float* const uniform dstPtr_line0 = (uniform float* uniform)self->dst.ptr + offset * 4;
uniform float* const uniform dstPtr_line1 = dstPtr_line0 + W*2*B; // next line
for (uniform size_t w = 0; w < W; ++w)
{
const T value = *((varying T* uniform)&srcPtr_line[w*B]);
const float value = *((varying float* uniform)&srcPtr_line[w*B]);
streaming_store(&dstPtr_line0[w*2*B ], value);
streaming_store(&dstPtr_line0[w*2*B+B], value);
-14
View File
@@ -1,14 +0,0 @@
// Copyright 2018 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
#include "cpu_upsample.isph"
#if TARGET_WIDTH * 2 != BLOCKC_BYTESIZE
#error "Unsupported ISPC target width"
#endif
export void CPUUpsampleKernel_run_f16(const uniform CPUUpsampleKernel* uniform self,
uniform int cb, uniform int h)
{
CPUUpsampleKernel_run<float16>(self, cb, h);
}
-14
View File
@@ -1,14 +0,0 @@
// Copyright 2018 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
#include "cpu_upsample.isph"
#if TARGET_WIDTH * 4 != BLOCKC_BYTESIZE
#error "Unsupported ISPC target width"
#endif
export void CPUUpsampleKernel_run_f32(const uniform CPUUpsampleKernel* uniform self,
uniform int cb, uniform int h)
{
CPUUpsampleKernel_run<float>(self, cb, h);
}
+94
View File
@@ -0,0 +1,94 @@
// Copyright 2018 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
#include "dnnl_common.h"
#include "dnnl_tensor.h"
OIDN_NAMESPACE_BEGIN
dnnl::memory::data_type toDNNL(DataType dataType)
{
switch (dataType)
{
case DataType::Float32:
return dnnl::memory::data_type::f32;
case DataType::Float16:
return dnnl::memory::data_type::f16;
case DataType::UInt8:
return dnnl::memory::data_type::u8;
default:
throw std::invalid_argument("unsupported data type");
}
}
dnnl::memory::desc toDNNL(const TensorDesc& td)
{
dnnl::memory::dims dnnlDims;
dnnl::memory::format_tag dnnlFormat;
switch (td.layout)
{
case TensorLayout::x:
assert(td.getRank() == 1);
dnnlDims = {td.getPaddedX()};
dnnlFormat = dnnl::memory::format_tag::x;
break;
case TensorLayout::chw:
assert(td.getRank() == 3);
dnnlDims = {1, td.getPaddedC(), td.getH(), td.getW()};
dnnlFormat = dnnl::memory::format_tag::nchw;
break;
case TensorLayout::Chw8c:
assert(td.getRank() == 3);
dnnlDims = {1, td.getPaddedC(), td.getH(), td.getW()};
dnnlFormat = dnnl::memory::format_tag::nChw8c;
break;
case TensorLayout::Chw16c:
assert(td.getRank() == 3);
dnnlDims = {1, td.getPaddedC(), td.getH(), td.getW()};
dnnlFormat = dnnl::memory::format_tag::nChw16c;
break;
case TensorLayout::oihw:
assert(td.getRank() == 4);
dnnlDims = {td.getPaddedO(), td.getPaddedI(), td.getH(), td.getW()};
dnnlFormat = dnnl::memory::format_tag::oihw;
break;
case TensorLayout::OIhw8i8o:
assert(td.getRank() == 4);
dnnlDims = {td.getPaddedO(), td.getPaddedI(), td.getH(), td.getW()};
dnnlFormat = dnnl::memory::format_tag::OIhw8i8o;
break;
case TensorLayout::OIhw16i16o:
assert(td.getRank() == 4);
dnnlDims = {td.getPaddedO(), td.getPaddedI(), td.getH(), td.getW()};
dnnlFormat = dnnl::memory::format_tag::OIhw16i16o;
break;
case TensorLayout::OIhw2o8i8o2i:
assert(td.getRank() == 4);
dnnlDims = {td.getPaddedO(), td.getPaddedI(), td.getH(), td.getW()};
dnnlFormat = dnnl::memory::format_tag::OIhw2o8i8o2i;
break;
default:
throw std::invalid_argument("unsupported tensor layout");
}
return dnnl::memory::desc(dnnlDims, toDNNL(td.dataType), dnnlFormat);
}
dnnl::memory toDNNL(const Ref<Buffer>& buffer)
{
DNNLEngine* engine = dynamic_cast<DNNLEngine*>(buffer->getEngine());
if (!engine)
throw std::logic_error("buffer does not belong to a DNNL engine");
dnnl::memory::desc desc({int64_t(buffer->getByteSize())}, dnnl::memory::data_type::u8, dnnl::memory::format_tag::x);
return {desc, engine->getDNNLEngine(), buffer->getPtr()};
}
const dnnl::memory& getDNNL(const Ref<Tensor>& tensor)
{
if (auto dnnlTensor = dynamic_cast<const DNNLTensor*>(tensor.get()))
return dnnlTensor->getDNNLMemory();
else
throw std::invalid_argument("not DNNLTensor");
}
OIDN_NAMESPACE_END
+20
View File
@@ -0,0 +1,20 @@
// Copyright 2018 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
#pragma once
#include "core/tensor.h"
#include "dnnl_engine.h"
OIDN_NAMESPACE_BEGIN
dnnl::memory::data_type toDNNL(DataType dataType);
dnnl::memory::desc toDNNL(const TensorDesc& td);
// Creates a DNNL memory structure for a buffer
dnnl::memory toDNNL(const Ref<Buffer>& buffer);
// Returns the internal DNNL memory structure of a DNNLTensor
const dnnl::memory& getDNNL(const Ref<Tensor>& tensor);
OIDN_NAMESPACE_END
+87
View File
@@ -0,0 +1,87 @@
// Copyright 2018 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
#include "dnnl_conv.h"
#include "dnnl_tensor.h"
OIDN_NAMESPACE_BEGIN
DNNLConv::DNNLConv(DNNLEngine* engine, const ConvDesc& desc)
: Conv(desc),
engine(engine)
{
const dnnl::memory::dims strides = {1, 1};
const dnnl::memory::dims padding = {1, 1};
// Incorporate activation
dnnl::primitive_attr attr;
if (activation == Activation::ReLU)
{
dnnl::post_ops ops;
ops.append_eltwise(
dnnl::algorithm::eltwise_relu,
0.f, // alpha
0.f // beta
);
attr.set_post_ops(ops);
}
attr.set_scratchpad_mode(dnnl::scratchpad_mode::user);
primDesc = dnnl::convolution_forward::primitive_desc(
engine->getDNNLEngine(),
dnnl::prop_kind::forward_inference, dnnl::algorithm::convolution_direct,
toDNNL(srcDesc),
toDNNL(weightDesc),
toDNNL(biasDesc),
toDNNL(dstDesc),
strides, padding, padding,
attr);
}
size_t DNNLConv::getScratchByteSize()
{
return primDesc.scratchpad_desc().get_size();
}
void DNNLConv::setScratch(const Ref<Buffer>& scratch)
{
this->scratch = scratch;
args[DNNL_ARG_SCRATCHPAD] = toDNNL(scratch);
}
void DNNLConv::updateSrc()
{
args[DNNL_ARG_SRC] = getDNNL(src);
}
void DNNLConv::updateWeight()
{
args[DNNL_ARG_WEIGHTS] = getDNNL(weight);
}
void DNNLConv::updateBias()
{
args[DNNL_ARG_BIAS] = getDNNL(bias);
}
void DNNLConv::updateDst()
{
args[DNNL_ARG_DST] = getDNNL(dst);
}
void DNNLConv::finalize()
{
prim = dnnl::convolution_forward(primDesc);
}
void DNNLConv::submitKernels(const Ref<CancellationToken>& ct)
{
if (!prim)
throw std::logic_error("convolution not finalized");
if (!src || !dst || !weight || !bias)
throw std::logic_error("convolution source/weight/bias/destination not set");
engine->submitFunc([=] { prim.execute(engine->getDNNLStream(), args); }, ct);
}
OIDN_NAMESPACE_END
+37
View File
@@ -0,0 +1,37 @@
// Copyright 2018 Intel Corporation
// SPDX-License-Identifier: Apache-2.0
#pragma once
#include "core/conv.h"
#include "dnnl_common.h"
OIDN_NAMESPACE_BEGIN
class DNNLConv final : public Conv
{
public:
DNNLConv(DNNLEngine* engine, const ConvDesc& desc);
Engine* getEngine() const override { return engine; }
size_t getScratchByteSize() override;
void setScratch(const Ref<Buffer>& scratch) override;
void finalize() override;
void submitKernels(const Ref<CancellationToken>& ct) override;
private:
void updateSrc() override;
void updateWeight() override;
void updateBias() override;
void updateDst() override;
DNNLEngine* engine;
dnnl::convolution_forward::primitive_desc primDesc;
dnnl::convolution_forward prim;
std::unordered_map<int, dnnl::memory> args;
Ref<Buffer> scratch;
};
OIDN_NAMESPACE_END

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