diff --git a/3rdparty/RadeonProML/RadeonProML.cpp b/3rdparty/RadeonProML/RadeonProML.cpp new file mode 100644 index 00000000..a932613b --- /dev/null +++ b/3rdparty/RadeonProML/RadeonProML.cpp @@ -0,0 +1 @@ +// Empty file to make CMake happy diff --git a/3rdparty/RadeonProML/RadeonProML.def b/3rdparty/RadeonProML/RadeonProML.def new file mode 100644 index 00000000..fb62571a --- /dev/null +++ b/3rdparty/RadeonProML/RadeonProML.def @@ -0,0 +1,15 @@ +EXPORTS + +mlCreateContext +mlReleaseContext +mlCreateImage +mlGetImageInfo +mlMapImage +mlUnmapImage +mlReleaseImage +mlCreateModel +mlGetModelInfo +mlSetModelInputInfo +mlInfer +mlReleaseModel +mlGetLastError diff --git a/3rdparty/RadeonProML/RadeonProML.h b/3rdparty/RadeonProML/RadeonProML.h new file mode 100644 index 00000000..c6ae5f48 --- /dev/null +++ b/3rdparty/RadeonProML/RadeonProML.h @@ -0,0 +1,278 @@ +#pragma once + +/** + * @file Machine learning model runner API. + * + * Typical usage is as follows: + * -# Create a context with mlCreateContext(). + * -# Set model parameters #ml_model_params. + * -# Create a model with mlCreateModel() using the parameters. + * -# Get input image information with mlGetModelInfo(). + * -# Specify required input image dimensions. + * -# Set up model input image dimensions with mlSetModelInputInfo(). + * -# Get output image information with mlGetModelInfo(). + * -# Create input image. + * -# Create output image. + * -# Fill input image with data using mlMapImage() and mlUnmapImage(). + * -# Run inference with mlInfer(). + * -# Get output image data using mlMapImage() and mlUnmapImage(). + * -# If image size is changed, repeat from the step 5. + * -# If image size is unchanged, repeat from the step 10. + * -# In a case of a failure invoke the mlGetLastError() to get details. + * -# Release the images using mlReleaseImage(). + * -# Release the model using mlReleaseModel(). + * -# Release the context using mlReleaseContext(). + */ + +#include + + +#if defined(_WIN32) + #ifdef RADEONPROML_BUILD + #define ML_API_ENTRY __declspec(dllexport) + #else + #define ML_API_ENTRY __declspec(dllimport) + #endif +#elif defined(__GNUC__) + #ifdef RADEONPROML_BUILD + #define ML_API_ENTRY __attribute__((visibility ("default"))) + #else + #define ML_API_ENTRY + #endif +#endif + + +#ifdef __cplusplus +extern "C" { +#endif + + +/** + * Model parameters. All unused values must be initialized to 0. + */ +typedef struct _ml_model_params +{ + char const* model_path; /**< Path to a model in protobuf format. */ + + char const* input_node; /**< Input graph node name, autodetect if null. */ + + char const* output_node; /**< Output graph node name, autodetect if null. */ + + float gpu_memory_fraction; /**< + * Fraction of GPU memory allowed to use + * by versions with GPU support, (0..1]. + * All memory is used by default. + */ + + char const* visible_devices; /**< + * Comma-delimited list of GPU devices + * accessible for calculations. + * All devices are visible by default. + */ +} ml_model_params; + +/** + * Context handle. + */ +typedef struct _ml_context* ml_context; + +/** + * Model handle. + */ +typedef struct _ml_model* ml_model; + +/** + * Image handle. + */ +typedef struct _ml_image* ml_image; + +/** + * Operation status. + */ +typedef enum _ml_status +{ + ML_OK, + ML_FAIL + +} ml_status; + +/** + * Image underlying data type. + */ +typedef enum _ml_data_type +{ + ML_FLOAT32, + ML_FLOAT16, + ML_INT32, + +} ml_data_type; + +/** +* Image access mode. +*/ +typedef enum _ml_access_mode +{ + ML_READ_ONLY, + ML_WRITE_ONLY, + ML_READ_WRITE + +} ml_access_mode; + +/** + * 3D image description. + */ +typedef struct _ml_image_info +{ + ml_data_type dtype; /**< Underlying data type. */ + size_t height; /**< Image height, in pixels. 0 if unspecified. */ + size_t width; /**< Image width. in pixels. 0 if unspecified. */ + size_t channels; /**< Image channel count. 0 if unspecified. */ + +} ml_image_info; + + +/** + * Creates a context. + * + * @return A valid context handle in case of success, NULL otherwise. + * The context should be released with mlReleaseContext(). + */ +ML_API_ENTRY ml_context mlCreateContext(); + +/** + * Releases a context created with mlCreateContext(), invalidates the handle. + * + * @param model A valid context handle. + */ +ML_API_ENTRY void mlReleaseContext(ml_context context); + +/** + * Creates a 3D image with a given description. + * Image dimension order is (height, width, channels). + * + * @param[in] context A valid context handle. + * @param[in] info Image description with all dimensions specified. + * @param[in] mode Image data access mode. + * + * @return A valid image handle in case of success, NULL otherwise. + * The image should be released with mlReleaseImage(). + * To get more details in case of failure, call mlGetLastError(). + */ + +ML_API_ENTRY ml_image mlCreateImage(ml_context context, ml_image_info const* info, ml_access_mode mode); + +/** + * Returns image description. + * + * @param[in] image A valid image handle. + * @param[out] info A pointer to a info structure. + * + * @return ML_OK in case of success, ML_FAIL otherwise. + */ +ML_API_ENTRY ml_status mlGetImageInfo(ml_image image, ml_image_info* info); + +/** + * Map the image data into the host address and returns a pointer + * to the mapped region. + * + * @param[in] image A valid image handle. + * @param[out] size A pointer to a size variable. If not null, the referenced + * value is set to image size, in bytes. + * + * @return A pointer to image data. + */ +ML_API_ENTRY void* mlMapImage(ml_image image, size_t* size); + +/** + * Unmaps a previously mapped image data. + * + * @param[in] image A valid image handle. + * @param[in] data A pointer to the previously mapped data. + * + * @return ML_OK in case of success, ML_FAIL otherwise. + */ +ML_API_ENTRY ml_status mlUnmapImage(ml_image image, void* data); + +/** + * Releases an image created with mlCreateImage(), invalidates the handle. + * + * @param[in] image A valid image handle. + */ +ML_API_ENTRY void mlReleaseImage(ml_image image); + + +/** + * Loads model data from a file. + * + * @param[in] context A valid context handle. + * @param[in] params Model parameters. @see #ml_model_params. + * + * @return ML_OK in case of success, ML_FAIL otherwise. + * To get more details in case of failure, call mlGetLastError(). + */ +ML_API_ENTRY ml_model mlCreateModel(ml_context context, ml_model_params const* params); + +/** + * Returns input image information. + * + * @param[in] model A valid model handle. + * @param[out] input_info A pointer to the result input info structure, may be null. + * If mlSetModelInputInfo() was not previously called, + * some dimensions may not be specified. + * @param[out] output_info A pointer to the result output info structure, may be null. + * If mlSetModelInputInfo() was not previously called, + * some dimensions may not be specified. + * + * @return ML_OK in case of success, ML_FAIL otherwise. + * To get more details in case of failure, call mlGetLastError(). + */ +ML_API_ENTRY ml_status mlGetModelInfo(ml_model model, + ml_image_info* input_info, + ml_image_info* output_info); + +/** + * Updates input image information. All image dimensions must be specified. + * @note This is a heavy operation, so avoid using it frequently. + * + * @param[in] model A valid model handle. + * @param[in] info Input image information. The specified dimensions must + * match the dimensions already known by the model. + */ +ML_API_ENTRY ml_status mlSetModelInputInfo(ml_model model, ml_image_info const* info); + +/** + * Gets an input image and fills an output image. + * + * @param[in] model A valid model handle. + * @param[in] input A valid input image descriptor. + * @param[in] output A valid output image descriptor. + * + * @return ML_OK in case of success, ML_FAIL otherwise. + * To get more details in case of failure, call mlGetLastError(). + */ +ML_API_ENTRY ml_status mlInfer(ml_model model, ml_image input, ml_image output); + +/** + * Releases a model loaded with mlCreateModel(), invalidates the handle. + * + * @param model A valid model handle. + */ +ML_API_ENTRY void mlReleaseModel(ml_model model); + +/** + * Returns a null-terminated string containing the last operation error message. + * May be called after some operation returns ML_FAIL or NULL. + * The error message is owned by the library and must NOT be freed by a client. + * The message is stored in a thread local storage, so this function + * should be called from the thread where the failure occured. + * + * @param[out] size Optional, the size of the error message (excluding the null-terminator). + * + * @return A pointer to the formatted message. + */ +ML_API_ENTRY const char* mlGetLastError(size_t* size); + + +#ifdef __cplusplus +} // extern "C" +#endif diff --git a/3rdparty/RadeonProML/RadeonProML.hpp b/3rdparty/RadeonProML/RadeonProML.hpp new file mode 100644 index 00000000..dc524798 --- /dev/null +++ b/3rdparty/RadeonProML/RadeonProML.hpp @@ -0,0 +1,195 @@ +#include + +#include +#include +#include + + +namespace RadeonProML { + +class Image +{ +public: + Image() = default; + + explicit Image(ml_image image) + : m_image(image) + { + } + + Image(Image&& other) noexcept + : m_image(other.m_image) + { + } + + ~Image() + { + if (m_image != nullptr) + { + mlReleaseImage(m_image); + } + } + + Image& operator=(Image&& other) noexcept + { + m_image = other.m_image; + other.m_image = nullptr; + return *this; + } + + void* Map(size_t* size = nullptr) + { + return mlMapImage(m_image, size); + } + + void Unmap(void* data) + { + mlUnmapImage(m_image, data); + } + + operator ml_image() const + { + return m_image; + } + +private: + ml_image m_image = nullptr; +}; + + +class Model +{ +public: + Model() = default; + + explicit Model(ml_model model) + : m_model(model) + { + } + + Model(Model&& other) noexcept + : m_model(other.m_model) + { + } + + ~Model() + { + if (m_model != nullptr) + { + mlReleaseModel(m_model); + } + } + + Model& operator =(Model&& other) noexcept + { + m_model = other.m_model; + other.m_model = nullptr; + return *this; + } + +#define CHECK_ML_STATUS(OP) CheckStatus(OP, #OP) + + void SetInputInfo(const ml_image_info& info) + { + CHECK_ML_STATUS(mlSetModelInputInfo(m_model, &info) == ML_OK); + } + + ml_image_info GetInputInfo() const + { + ml_image_info info; + CHECK_ML_STATUS(mlGetModelInfo(m_model, &info, nullptr) == ML_OK); + return info; + } + + ml_image_info GetOutputInfo() const + { + ml_image_info info; + CHECK_ML_STATUS(mlGetModelInfo(m_model, nullptr, &info) == ML_OK); + return info; + } + + void Infer(const Image& input, const Image& output) + { + CHECK_ML_STATUS(mlInfer(m_model, input, output) == ML_OK); + } + +#undef CHECK_ML_STATUS + +private: + void CheckStatus(bool status, const char* op_name) const + { + if (!status) + { + std::string func_name(op_name); + const auto end_pos = func_name.find('('); + func_name.erase(end_pos != std::string::npos ? end_pos : func_name.size()); + + throw std::runtime_error(func_name + " failed: " + mlGetLastError(nullptr)); + } + } + + ml_model m_model = nullptr; +}; + + +class Context +{ +public: + Context() = default; + + explicit Context(ml_context context) + : m_context(context) + { + } + + Context(Context&& other) noexcept + : m_context(other.m_context) + { + } + + ~Context() + { + if (m_context != nullptr) + { + mlReleaseContext(m_context); + } + } + + Context& operator =(Context&& other) noexcept + { + m_context = other.m_context; + other.m_context = nullptr; + return *this; + } + + Model CreateModel(const ml_model_params& params) + { + auto model = mlCreateModel(m_context, ¶ms); + CheckStatus(model != nullptr, "mlCreateModel"); + return Model(model); + } + + Image CreateImage(const ml_image_info& info, ml_access_mode mode) + { + auto image = mlCreateImage(m_context, &info, mode); + CheckStatus(image != nullptr, "mlCreateImage"); + return Image(image); + } + +private: + void CheckStatus(bool status, const char* op_name) const + { + if (!status) + { + std::string func_name(op_name); + const auto end_pos = func_name.find('('); + func_name.erase(end_pos != std::string::npos ? end_pos : func_name.size()); + + throw std::runtime_error(func_name + " failed: " + mlGetLastError(nullptr)); + } + } + + ml_context m_context = nullptr; +}; + +} // namespace RadeonProML diff --git a/3rdparty/RadeonProML/RadeonProML_cl.h b/3rdparty/RadeonProML/RadeonProML_cl.h new file mode 100644 index 00000000..7d87a00f --- /dev/null +++ b/3rdparty/RadeonProML/RadeonProML_cl.h @@ -0,0 +1,51 @@ +#pragma once + +/** + * @file Interop with OpenCL. + */ + +#include + +#if defined(__APPLE__) || defined(__MACOSX) +#include +#else +#include +#endif + + +#ifdef __cplusplus +extern "C" { +#endif + +/** + * Creates a context from an OpenCL context. + * + * @param[in] queue An OpenCL command queue handle. + * + * @return A valid context handle in case of success, NULL otherwise. + * The context should be released with mlReleaseContext(). + */ + +ML_API_ENTRY ml_context mlCreateContextFromClQueue(cl_command_queue queue); + +/** + * Creates an image from an OpenCL buffer. + * + * @param[in] context A valid context handle. + * @param[in] buffer A valid OpenCL memory object handle. + * @param[in] info Image description with all dimensions specified. + * @param[in] mode Image data access mode. + * + * @return A valid image handle in case of success, NULL otherwise. + * The image should be released with mlReleaseImage(). + * To get more details in case of failure, call mlGetContextError(). + */ + +ML_API_ENTRY ml_image mlCreateImageFromClBuffer(ml_context context, + cl_mem buffer, + const ml_image_info* info, + ml_access_mode mode); + +#ifdef __cplusplus +} // extern "C" +#endif diff --git a/Baikal/CMakeLists.txt b/Baikal/CMakeLists.txt index ca9b60a1..4289fbbe 100644 --- a/Baikal/CMakeLists.txt +++ b/Baikal/CMakeLists.txt @@ -13,15 +13,33 @@ set(ESTIMATORS_SOURCES set(OUTPUT_SOURCES Output/clwoutput.h Output/output.h) - + +set(POSTEFFECT_ML_SOURCES + PostEffects/ML/upsampler_preprocessor.h + PostEffects/ML/upsampler_preprocessor.cpp + PostEffects/ML/denoiser_preprocessor.h + PostEffects/ML/denoiser_preprocessor.cpp + PostEffects/ML/inference.h + PostEffects/ML/inference.cpp + PostEffects/ML/ml_post_effect.h + PostEffects/ML/ml_post_effect.cpp + PostEffects/ML/model_holder.h + PostEffects/ML/model_holder.cpp + PostEffects/ML/data_preprocessor.h + PostEffects/ML/data_preprocessor.cpp + PostEffects/ML/image.h + PostEffects/ML/image.cpp + PostEffects/ML/error_handler.h) + set(POSTEFFECT_SOURCES PostEffects/clw_post_effect.h PostEffects/post_effect.h + PostEffects/post_effect.cpp PostEffects/bilateral_denoiser.h PostEffects/wavelet_denoiser.h PostEffects/AreaMap33.h - ) - +) + set(RENDERERS_SOURCES Renderers/adaptive_renderer.cpp Renderers/adaptive_renderer.h @@ -84,11 +102,11 @@ set(SCENEGRAPH_SOURCES SceneGraph/uberv2material.h SceneGraph/inputmap.h SceneGraph/inputmaps.h) - + set(SCENEGRAPH_COLLECTOR_SOURCES SceneGraph/Collector/collector.cpp SceneGraph/Collector/collector.h) - + set(XML_SOURCES XML/tinyxml2.cpp XML/tinyxml2.h) @@ -126,6 +144,7 @@ set(SOURCES ${ESTIMATORS_SOURCES} ${OUTPUT_SOURCES} ${POSTEFFECT_SOURCES} + ${POSTEFFECT_ML_SOURCES} ${RENDERERS_SOURCES} ${RENDERFACTORY_SOURCES} ${UTILS_SOURCES} @@ -133,13 +152,13 @@ set(SOURCES ${SCENEGRAPH_COLLECTOR_SOURCES} ${SCENEGRAPH_IO_SOURCES} ${XML_SOURCES} - ${KERNELS_SOURCES} - ) + ${KERNELS_SOURCES}) source_group("Controllers" FILES ${CONTROLLERS_SOURCES}) source_group("Estimators" FILES ${ESTIMATORS_SOURCES}) source_group("Output" FILES ${OUTPUT_SOURCES}) source_group("Posteffect" FILES ${POSTEFFECT_SOURCES}) +source_group("Posteffect\\ML" FILES ${POSTEFFECT_ML_SOURCES}) source_group("Renderers" FILES ${RENDERERS_SOURCES}) source_group("RenderFactory" FILES ${RENDERFACTORY_SOURCES}) source_group("Utils" FILES ${UTILS_SOURCES}) @@ -162,10 +181,6 @@ add_library(Baikal SHARED ${SOURCES}) # actually doesn't export data symbols target_compile_definitions(Baikal PRIVATE BAIKAL_EXPORT_DATA_SYMBOLS) -if (BAIKAL_ENABLE_DENOISER) - target_compile_definitions(Baikal PUBLIC ENABLE_DENOISER) -endif(BAIKAL_ENABLE_DENOISER) - if (BAIKAL_ENABLE_RAYMASK) target_compile_definitions(Baikal PUBLIC ENABLE_RAYMASK) endif (BAIKAL_ENABLE_RAYMASK) @@ -196,6 +211,7 @@ if (WIN32) target_compile_options(Baikal PUBLIC /WX) elseif (UNIX) target_compile_options(Baikal PUBLIC -Wall -Werror) + target_link_libraries(Baikal PUBLIC dl) endif (WIN32) # Add symbolic link from binary to Baikal/Kernels directory @@ -222,5 +238,10 @@ else () install(TARGETS Baikal LIBRARY DESTINATION lib) endif () +# add ml libs +target_include_directories(Baikal PUBLIC "${Baikal_SOURCE_DIR}/3rdparty/RadeonProML") +target_link_directories(Baikal PUBLIC "${Baikal_SOURCE_DIR}/3rdparty/RadeonProML") +target_link_libraries(Baikal PUBLIC "${Baikal_SOURCE_DIR}/3rdparty/RadeonProML/RadeonProML.lib") + file(TO_CMAKE_PATH ${BAIKAL_KERNELS_SRC} BAIKAL_KERNELS_SRC_CMAKE) install(DIRECTORY ${BAIKAL_KERNELS_SRC_CMAKE} DESTINATION Baikal) diff --git a/Baikal/Controllers/clw_scene_controller.cpp b/Baikal/Controllers/clw_scene_controller.cpp index a363084d..82f7eb92 100644 --- a/Baikal/Controllers/clw_scene_controller.cpp +++ b/Baikal/Controllers/clw_scene_controller.cpp @@ -227,7 +227,6 @@ namespace Baikal auto transform = mesh->GetTransform(); shape->SetTransform(transform, inverse(transform)); shape->SetId(id++); - shape->SetMask(iter->GetVisibilityMask()); out.isect_shapes.push_back(shape); out.visible_shapes.push_back(shape); diff --git a/Baikal/Kernels/CL/denoise.cl b/Baikal/Kernels/CL/denoise.cl index 88bf8697..132c1454 100644 --- a/Baikal/Kernels/CL/denoise.cl +++ b/Baikal/Kernels/CL/denoise.cl @@ -23,6 +23,7 @@ THE SOFTWARE. #define DENOISE_CL #include <../Baikal/Kernels/CL/common.cl> +#include <../Baikal/Kernels/CL/utils.cl> // Similarity function inline float C(float3 x1, float3 x2, float sigma) @@ -112,4 +113,155 @@ void BilateralDenoise_main( } } +KERNEL +void ToneMappingExponential(GLOBAL float4* restrict dst, + GLOBAL float4 const* restrict src, + int elems_num) +{ + int id = get_global_id(0); + + if (id >= elems_num) + { + return; + } + + if (src[id].w != 0.0f) + { + dst[id].xyz = 1.f - exp(-1.2f * src[id].xyz / src[id].w); + } + else + { + dst[id] = make_float4(0.0f, 0.0f, 0.0f, 1.0f); + } +} + + +// perform division on w component +KERNEL +void DivideBySampleCount(GLOBAL float4* restrict dst, + GLOBAL float4 const* restrict src, + int elems_num) +{ + int id = get_global_id(0); + + if (id >= elems_num) + { + return; + } + + if (src[id].w != 0.0f) + { + dst[id].xyz = src[id].xyz / src[id].w; + dst[id].w = src[id].w; + } + else + { + dst[id] = make_float4(0.0f, 0.0f, 0.0f, 1.0f); + } + +} + +KERNEL +void CopyInterleaved(GLOBAL float4* restrict dst, + GLOBAL float4 const* restrict src, + int dst_width, + int dst_height, + int dst_channels_offset, // offset inside pixel in channels (not bytes) + int dst_channels_num, + int src_width, + int src_height, + int src_channels_offset, // offset inside pixel in channels (not bytes) + int src_channels_num, + int channels_to_copy) +{ + int global_id = get_global_id(0); + + int x = global_id % dst_width; + int y = global_id / dst_width; + + if ((x > dst_width) || (global_id > dst_width * dst_height)) + { + return; + } + + int src_offset = src_channels_num * (y * src_width + x) + src_channels_offset; + int dst_offset = dst_channels_num * (y * dst_width + x) + dst_channels_offset; + + GLOBAL float* dst_pixel = (GLOBAL float*)dst + dst_offset; + GLOBAL float const* src_pixel = (GLOBAL float const*)src + src_offset; + + for (int i = 0; i < channels_to_copy; i++) + { + dst_pixel[i] = src_pixel[i]; + } +} + +static float3 BicubicConvolutionCompute(float4 f0, float4 f1, float4 f2, float4 f3) +{ + const float t = 0.5f; + float4 a0 = f1; + float4 a1 = (f0 - f2) / 2.f; + float4 a2 = -f0 - 3.5f * f1 + 4 * f2 + .5f * f3; + float4 a3 = .5f * f0 + 2.5f * f1 - 2.5f * f2 - .5f * f3; + + return (a3 * t * t * t + a2 * t * t + a1 * t + a0).xyz; +} + +KERNEL +void BicubicUpscale2x_X(// size of the dst buffer should be enough + // to store 2 * sizeof(float3) * width * height + GLOBAL float4* restrict dst, + GLOBAL float4 const* restrict src, + int width, + int height) +{ + int idx = get_global_id(0); + int src_idx = idx / 2; + + dst[idx].w = src[0].w; + + if (idx % 2 == 0 || (idx + 1) % (2 * width) == 0) + { + dst[idx].xyz = src[idx / 2].xyz; + return; + } + + dst[idx].xyz = BicubicConvolutionCompute(src[src_idx - 1], + src[src_idx], + src[src_idx + 1], + src[src_idx + 2]); + + +} + + +KERNEL +void BicubicUpscale2x_Y(// size of the dst buffer should be enough to store + // 2 * sizeof(float3) * width * height + GLOBAL float4* restrict dst, + GLOBAL float4 const* restrict src, + int width, + int height) +{ + int idx = get_global_id(0); + + int x_coord = idx % width; // same for dst and src buffers + int dst_y = (idx - x_coord) / width; + int src_y = dst_y / 2; + int src_idx = src_y * width + x_coord; + + dst[idx].w = src[0].w; + + if (dst_y % 2 == 0 || dst_y == 1 || dst_y > height - 2) + { + dst[idx].xyz = src[src_idx].xyz; + return; + } + + dst[idx].xyz = BicubicConvolutionCompute(src[src_idx - width], + src[src_idx], + src[src_idx + width], + src[src_idx + 2 * width]); +} + #endif diff --git a/Baikal/Kernels/CL/fill_aovs_uberv2.cl b/Baikal/Kernels/CL/fill_aovs_uberv2.cl index de060972..3f7b7c94 100644 --- a/Baikal/Kernels/CL/fill_aovs_uberv2.cl +++ b/Baikal/Kernels/CL/fill_aovs_uberv2.cl @@ -196,10 +196,6 @@ KERNEL void FillAOVsUberV2( aov_background[idx].w += 1.0f; CORRECT_VALUE(aov_background[idx]) } - else - { - - } if (isect.shapeid > -1) @@ -370,7 +366,7 @@ KERNEL void FillAOVsUberV2( UberV2PrepareInputs(&diffgeo, input_map_values, material_attributes, TEXTURE_ARGS, &uber_shader_data); GetMaterialBxDFType(wi, &sampler, SAMPLER_ARGS, &diffgeo, &uber_shader_data); - int sampled_component = Bxdf_UberV2_GetSampledComponent(&diffgeo); + int sampled_component = Bxdf_UberV2_GetSampledComponent(&diffgeo); float gloss = 0.f; if (sampled_component == kBxdfUberV2SampleCoating) diff --git a/Baikal/Kernels/CL/monte_carlo_renderer.cl b/Baikal/Kernels/CL/monte_carlo_renderer.cl index 07c07b33..9a17be50 100644 --- a/Baikal/Kernels/CL/monte_carlo_renderer.cl +++ b/Baikal/Kernels/CL/monte_carlo_renderer.cl @@ -647,7 +647,45 @@ KERNEL void ApplyGammaAndCopyData( float4 val = clamp(native_powr(v / v.w, 1.f / gamma), 0.f, 1.f); write_imagef(img, make_int2(global_idx, global_idy), val); } -} +} + +// Copy data to interop texture if supported +KERNEL void ApplyGammaAndCopySplitData( + GLOBAL float4 const* left_data, + GLOBAL float4 const* right_data, + int img_width, + int img_height, + float gamma, + write_only image2d_t img +) +{ + int global_id = get_global_id(0); + + int global_idx = global_id % img_width; + int global_idy = global_id / img_width; + + if (global_idx < img_width && global_idy < img_height) + { + float4 v; + if (global_idx < img_width / 2) + { + v = left_data[global_id]; + } + else + { + v = right_data[global_id]; + } + +#ifdef ADAPTIVITY_DEBUG + float a = v.w < 1024 ? min(1.f, v.w / 1024.f) : 0.f; + float4 mul_color = make_float4(1.f, 1.f - a, 1.f - a, 1.f); + v *= mul_color; +#endif + + float4 val = clamp(native_powr(v / v.w, 1.f / gamma), 0.f, 1.f); + write_imagef(img, make_int2(global_idx, global_idy), val); + } +} KERNEL void AccumulateSingleSample( GLOBAL float4 const* restrict src_sample_data, diff --git a/Baikal/PostEffects/ML/data_preprocessor.cpp b/Baikal/PostEffects/ML/data_preprocessor.cpp new file mode 100644 index 00000000..b0d649ae --- /dev/null +++ b/Baikal/PostEffects/ML/data_preprocessor.cpp @@ -0,0 +1,83 @@ +/********************************************************************** + Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved. + + Permission is hereby granted, free of charge, to any person obtaining a copy + of this software and associated documentation files (the "Software"), to deal + in the Software without restriction, including without limitation the rights + to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + copies of the Software, and to permit persons to whom the Software is + furnished to do so, subject to the following conditions: + + The above copyright notice and this permission notice shall be included in + all copies or substantial portions of the Software. + + THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + THE SOFTWARE. + ********************************************************************/ + +#include "data_preprocessor.h" + +namespace Baikal +{ + namespace PostEffects + { + DataPreprocessor::DataPreprocessor(CLWContext context, + CLProgramManager const* program_manager, + std::uint32_t start_spp) +#ifdef BAIKAL_EMBED_KERNELS + : ClwClass(context, program_manager, "denoise", g_denoise_opencl, g_denoise_opencl_headers) +#else + : ClwClass(context, program_manager, "../Baikal/Kernels/CL/denoise.cl") +#endif + , m_start_spp(start_spp) + { } + + + CLWEvent DataPreprocessor::WriteToInputs(CLWBuffer const& dst_buffer, + CLWBuffer const& src_buffer, + int width, + int height, + int dst_channels_offset, + int dst_channels_num, + int src_channels_offset, + int src_channels_num, + int channels_to_copy) + { + auto copy_kernel = GetKernel("CopyInterleaved"); + + unsigned argc = 0; + copy_kernel.SetArg(argc++, dst_buffer); + copy_kernel.SetArg(argc++, src_buffer); + copy_kernel.SetArg(argc++, width); + copy_kernel.SetArg(argc++, height); + copy_kernel.SetArg(argc++, dst_channels_offset); + copy_kernel.SetArg(argc++, dst_channels_num); + // input and output buffers have the same width in pixels + copy_kernel.SetArg(argc++, width); + // input and output buffers have the same height in pixels + copy_kernel.SetArg(argc++, height); + copy_kernel.SetArg(argc++, src_channels_offset); + copy_kernel.SetArg(argc++, src_channels_num); + copy_kernel.SetArg(argc++, channels_to_copy); + + // run copy_kernel + auto thread_num = ((width * height + 63) / 64) * 64; + return GetContext().Launch1D(0, + thread_num, + 64, + copy_kernel); + } + + unsigned DataPreprocessor::ReadSpp(CLWBuffer const &buffer) + { + RadeonRays::float3 pixel; + GetContext().ReadBuffer(0, buffer, &pixel, 1).Wait(); + return static_cast(pixel.w); + } + } +} diff --git a/Baikal/PostEffects/ML/data_preprocessor.h b/Baikal/PostEffects/ML/data_preprocessor.h new file mode 100644 index 00000000..6b216d5b --- /dev/null +++ b/Baikal/PostEffects/ML/data_preprocessor.h @@ -0,0 +1,75 @@ +/********************************************************************** + Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved. + + Permission is hereby granted, free of charge, to any person obtaining a copy + of this software and associated documentation files (the "Software"), to deal + in the Software without restriction, including without limitation the rights + to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + copies of the Software, and to permit persons to whom the Software is + furnished to do so, subject to the following conditions: + + The above copyright notice and this permission notice shall be included in + all copies or substantial portions of the Software. + + THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + THE SOFTWARE. + ********************************************************************/ + +#pragma once + +#ifdef BAIKAL_EMBED_KERNELS +#include "embed_kernels.h" +#endif + +#include "CLW.h" +#include "Utils/clw_class.h" +#include "PostEffects/post_effect.h" +#include "image.h" + + +namespace Baikal +{ + namespace PostEffects + { + class DataPreprocessor : public ClwClass + { + public: + DataPreprocessor(CLWContext context, + CLProgramManager const* program_manager, + std::uint32_t start_spp = 1); + + virtual Image MakeInput(PostEffect::InputSet const& inputs) = 0; + + virtual std::set GetInputTypes() const = 0; + + // returns channels num at input and output + virtual std::tuple ChannelsNum() const = 0; + + void SetStartSpp(std::uint32_t start_spp) + { m_start_spp = start_spp; } + + protected: + template + using Handle = std::unique_ptr::type, void (*)(T)>; + + unsigned ReadSpp(CLWBuffer const& buffer); + + CLWEvent WriteToInputs(CLWBuffer const& dst_buffer, + CLWBuffer const& src_buffer, + int width, + int height, + int dst_channels_offset, + int dst_channels_num, + int src_channels_offset, + int src_channels_num, + int channels_to_copy); + + std::uint32_t m_start_spp; + }; + } +} diff --git a/Baikal/PostEffects/ML/denoiser_preprocessor.cpp b/Baikal/PostEffects/ML/denoiser_preprocessor.cpp new file mode 100644 index 00000000..94a96a06 --- /dev/null +++ b/Baikal/PostEffects/ML/denoiser_preprocessor.cpp @@ -0,0 +1,249 @@ +/********************************************************************** + Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved. + + Permission is hereby granted, free of charge, to any person obtaining a copy + of this software and associated documentation files (the "Software"), to deal + in the Software without restriction, including without limitation the rights + to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + copies of the Software, and to permit persons to whom the Software is + furnished to do so, subject to the following conditions: + + The above copyright notice and this permission notice shall be included in + all copies or substantial portions of the Software. + + THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + THE SOFTWARE. +********************************************************************/ + + +#include "PostEffects/ML/denoiser_preprocessor.h" +#include "PostEffects/ML/error_handler.h" + +#include "CLWBuffer.h" +#include "Output/clwoutput.h" +#include "math/mathutils.h" + +#include "RadeonProML.h" + +namespace Baikal +{ + namespace PostEffects + { + using float3 = RadeonRays::float3; + using OutputType = Renderer::OutputType; + + DenoiserPreprocessor::DenoiserPreprocessor(CLWContext context, + CLProgramManager const* program_manager, + std::uint32_t start_spp) + : DataPreprocessor(context, program_manager, start_spp) + , m_primitives(context) + , m_model(Model::kColorAlbedoDepthNormal9) + , m_context(mlCreateContext(), mlReleaseContext) + { + switch (m_model) + { + case Model::kColorDepthNormalGloss7: + m_layout.emplace_back(OutputType::kColor, 3); + m_layout.emplace_back(OutputType::kDepth, 1); + m_layout.emplace_back(OutputType::kViewShadingNormal, 2); + m_layout.emplace_back(OutputType::kGloss, 1); + break; + case Model::kColorAlbedoNormal8: + m_layout.emplace_back(OutputType::kColor, 3 ); + m_layout.emplace_back(OutputType::kAlbedo, 3); + m_layout.emplace_back(OutputType::kViewShadingNormal, 2); + break; + case Model::kColorAlbedoDepthNormal9: + m_layout.emplace_back(OutputType::kColor, 3 ); + m_layout.emplace_back(OutputType::kAlbedo, 3); + m_layout.emplace_back(OutputType::kDepth, 1); + m_layout.emplace_back(OutputType::kViewShadingNormal, 2); + break; + } + + for (const auto& layer: m_layout) + { + m_channels += layer.second; + } + } + + void DenoiserPreprocessor::Init(std::uint32_t width, std::uint32_t height) + { + auto context = GetContext(); + + m_cache = CLWBuffer::Create(context, + CL_MEM_READ_WRITE, + 4 * width * height); + + m_input = CLWBuffer::Create(context, + CL_MEM_READ_WRITE, + m_channels * width * height); + + ml_image_info image_info = {ML_FLOAT32, m_height, m_width, m_channels}; + m_image = mlCreateImage(m_context.get(), &image_info, ML_READ_WRITE); + + + if (!m_image) + { + ContextError(m_context.get()); + } + + m_is_initialized = true; + } + + Image DenoiserPreprocessor::MakeInput(PostEffect::InputSet const& inputs) + { + auto context = GetContext(); + unsigned channels_count = 0u; + unsigned real_sample_count = 0u; + + if (!m_is_initialized) + { + auto color = inputs.at(Renderer::OutputType::kColor); + m_width = color->width(); + m_height = color->height(); + Init(m_width, m_height); + m_is_initialized = true; + } + + for (const auto& desc : m_layout) + { + auto type = desc.first; + auto input = inputs.at(type); + + auto clw_output = dynamic_cast(input); + auto device_mem = clw_output->data(); + + switch (type) + { + case OutputType::kColor: + DivideBySampleCount(CLWBuffer::CreateFromClBuffer(m_cache), + device_mem); + + WriteToInputs(m_input, + m_cache, + m_width, + m_height, + channels_count /* dst channels offset */, + m_channels /* dst channels num */, + 0 /* src channels offset */, + 4 /* src channels num */, + 3 /* channels to copy */); + + channels_count += 3; + + real_sample_count = ReadSpp(CLWBuffer::CreateFromClBuffer(m_cache)); + + if (real_sample_count < m_start_spp) + { + return Image(real_sample_count, nullptr); + } + break; + + case OutputType::kDepth: + m_primitives.Normalize(0, + CLWBuffer::CreateFromClBuffer(device_mem), + CLWBuffer::CreateFromClBuffer(m_cache), + (int)device_mem.GetElementCount() / sizeof(cl_float3)); + + WriteToInputs(m_input, + m_cache, + m_width, + m_height, + channels_count, // dst channels offset + m_channels, // dst channels num + 0, // src channels offset + 4, // src channels num + 1).Wait(); // channels to copy + + channels_count += 1; + break; + + case OutputType::kViewShadingNormal: + case OutputType::kGloss: + case OutputType::kAlbedo: + DivideBySampleCount(CLWBuffer::CreateFromClBuffer(m_cache), + CLWBuffer::CreateFromClBuffer(device_mem)); + + WriteToInputs(m_input, + m_cache, + m_width, + m_height, + channels_count, // dst channels offset + m_channels, // dst channels num + 0, // src channels offset + 4, // src channels num + desc.second).Wait(); // channels to copy + + channels_count += desc.second; + + default: + break; + } + } + + size_t image_size = 0; + auto host_buffer = mlMapImage(m_image, &image_size); + + if (!host_buffer || image_size == 0) + { + throw std::runtime_error("map operation failed"); + } + + context.ReadBuffer(0, + m_input, + static_cast(host_buffer), + m_input.GetElementCount()).Wait(); + + if (mlUnmapImage(m_image, host_buffer) != ML_OK) + { + throw std::runtime_error("unmap operation failed"); + } + + return Image(static_cast(real_sample_count), m_image); + } + + std::tuple DenoiserPreprocessor::ChannelsNum() const + { + return {m_channels, 3}; + } + + std::set DenoiserPreprocessor::GetInputTypes() const + { + std::set out_set; + + for (const auto& type: m_layout) + { + out_set.insert(type.first); + } + + return out_set; + } + + void DenoiserPreprocessor::DivideBySampleCount(CLWBuffer const& dst, + CLWBuffer const& src) + { + assert(dst.GetElementCount() >= src.GetElementCount()); + + auto division_kernel = GetKernel("DivideBySampleCount"); + + // Set kernel parameters + unsigned argc = 0; + division_kernel.SetArg(argc++, dst); + division_kernel.SetArg(argc++, src); + division_kernel.SetArg(argc++, (int)src.GetElementCount()); + + // run DivideBySampleCount kernel + auto thread_num = ((src.GetElementCount() + 63) / 64) * 64; + GetContext().Launch1D(0, + thread_num, + 64, + division_kernel); + } + } +} diff --git a/Baikal/PostEffects/ML/denoiser_preprocessor.h b/Baikal/PostEffects/ML/denoiser_preprocessor.h new file mode 100644 index 00000000..a28bbc96 --- /dev/null +++ b/Baikal/PostEffects/ML/denoiser_preprocessor.h @@ -0,0 +1,83 @@ +/********************************************************************** +Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +********************************************************************/ + +#pragma once + +#include "data_preprocessor.h" + +#include "CLW.h" +#include "Utils/clw_class.h" + +#ifdef BAIKAL_EMBED_KERNELS +#include "embed_kernels.h" +#endif + +#include +#include +#include + + +namespace Baikal +{ + namespace PostEffects + { + enum class Model + { + kColorDepthNormalGloss7, + kColorAlbedoNormal8, + kColorAlbedoDepthNormal9 + }; + + class DenoiserPreprocessor: public DataPreprocessor + { + public: + DenoiserPreprocessor(CLWContext context, + CLProgramManager const* program_manager, + std::uint32_t start_spp = 8); + + Image MakeInput(PostEffect::InputSet const& inputs) override; + + std::set GetInputTypes() const override; + + std::tuple ChannelsNum() const override; + private: + void Init(std::uint32_t width, std::uint32_t height); + + // layout of the outputs in input tensor in terms of channels + using MemoryLayout = std::vector>; + + void DivideBySampleCount(CLWBuffer const& dst, + CLWBuffer const& src); + + bool m_is_initialized = false; + CLWParallelPrimitives m_primitives; + std::uint32_t m_width, m_height; + std::uint32_t m_channels = 0; + Model m_model; + MemoryLayout m_layout; + CLWBuffer m_cache; + CLWBuffer m_input; + Handle m_context; + ml_image m_image; + }; + } +} diff --git a/Baikal/PostEffects/ML/error_handler.h b/Baikal/PostEffects/ML/error_handler.h new file mode 100644 index 00000000..f5bf93f4 --- /dev/null +++ b/Baikal/PostEffects/ML/error_handler.h @@ -0,0 +1,47 @@ +/********************************************************************** + Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved. + + Permission is hereby granted, free of charge, to any person obtaining a copy + of this software and associated documentation files (the "Software"), to deal + in the Software without restriction, including without limitation the rights + to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + copies of the Software, and to permit persons to whom the Software is + furnished to do so, subject to the following conditions: + + The above copyright notice and this permission notice shall be included in + all copies or substantial portions of the Software. + + THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + THE SOFTWARE. + ********************************************************************/ + +#pragma once + +#include "RadeonProML.h" +#include +#include + +namespace Baikal +{ + namespace PostEffects + { + inline void CheckStatus(ml_status status) + { + if (status != ML_OK) + { + throw std::runtime_error(mlGetLastError(nullptr)); + } + } + + + inline void ContextError(ml_context context) + { + throw std::runtime_error(mlGetLastError(nullptr)); + } + } +} diff --git a/Baikal/PostEffects/ML/image.cpp b/Baikal/PostEffects/ML/image.cpp new file mode 100644 index 00000000..2d53d880 --- /dev/null +++ b/Baikal/PostEffects/ML/image.cpp @@ -0,0 +1,51 @@ +/********************************************************************** + Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved. + + Permission is hereby granted, free of charge, to any person obtaining a copy + of this software and associated documentation files (the "Software"), to deal + in the Software without restriction, including without limitation the rights + to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + copies of the Software, and to permit persons to whom the Software is + furnished to do so, subject to the following conditions: + + The above copyright notice and this permission notice shall be included in + all copies or substantial portions of the Software. + + THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + THE SOFTWARE. + ********************************************************************/ + +#include "image.h" + +namespace Baikal +{ + namespace PostEffects + { + Image::Image() + : tag(0), image(nullptr) + {} + + Image::Image(std::uint32_t tag, ml_image image) + : tag(tag), image(image) + {} + + Image::Image(Image &&img) + : tag(img.tag), image(img.image) + { + img.image = nullptr; + } + + Image &Image::operator=(Image &&img) + { + tag = img.tag; + image = img.image; + img.image = nullptr; + return *this; + } + } +} diff --git a/Baikal/PostEffects/ML/image.h b/Baikal/PostEffects/ML/image.h new file mode 100644 index 00000000..e39daf47 --- /dev/null +++ b/Baikal/PostEffects/ML/image.h @@ -0,0 +1,44 @@ +/********************************************************************** + Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved. + + Permission is hereby granted, free of charge, to any person obtaining a copy + of this software and associated documentation files (the "Software"), to deal + in the Software without restriction, including without limitation the rights + to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + copies of the Software, and to permit persons to whom the Software is + furnished to do so, subject to the following conditions: + + The above copyright notice and this permission notice shall be included in + all copies or substantial portions of the Software. + + THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + THE SOFTWARE. + ********************************************************************/ + +#pragma once + +#include +#include "RadeonProML.h" + +namespace Baikal +{ + namespace PostEffects + { + struct Image + { + Image(); + Image(std::uint32_t tag, ml_image image); + Image(Image &&); + + Image &operator=(Image &&image); + + std::uint32_t tag; + ml_image image; + }; + } +} diff --git a/Baikal/PostEffects/ML/inference.cpp b/Baikal/PostEffects/ML/inference.cpp new file mode 100644 index 00000000..2718d9f6 --- /dev/null +++ b/Baikal/PostEffects/ML/inference.cpp @@ -0,0 +1,136 @@ +/********************************************************************** +Copyright (c) 2018 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +********************************************************************/ + +#include "PostEffects/ML/error_handler.h" +#include "inference.h" +#include "RadeonProML.h" + +#include + +namespace Baikal +{ + namespace PostEffects + { + Inference::Inference(std::string const& model_path, + size_t input_height, + size_t input_width, + float gpu_memory_fraction, + std::string const& visible_devices) + : m_model(model_path, + gpu_memory_fraction, + visible_devices) + { + CheckStatus(mlGetModelInfo(m_model.GetModel(), &m_input_info, nullptr)); + // Set unspecified input tensor dimensions + m_input_info.width = input_width; + m_input_info.height = input_height; + CheckStatus(mlSetModelInputInfo(m_model.GetModel(), &m_input_info)); + + // Get output tensor shape for model + CheckStatus(mlGetModelInfo(m_model.GetModel(), nullptr, &m_output_info)); + + m_worker = std::thread(&Inference::DoInference, this); + } + + Inference::~Inference() + { + Shutdown(); + } + + ml_image_info Inference::GetInputInfo() const + { + return m_input_info; + } + + ml_image_info Inference::GetOutputInfo() const + { + return m_output_info; + } + + Image Inference::GetInputData() + { + return {0, AllocImage(m_input_info, ML_READ_WRITE)}; + } + + void Inference::PushInput(Image&& image) + { + m_input_queue.push(std::move(image)); + } + + Image Inference::TryPopOutput() + { + Image output_tensor = {0, nullptr}; + m_output_queue.try_pop(output_tensor); + return output_tensor; + } + + Image Inference::PopOutput() + { + Image output_tensor = { 0, nullptr }; + m_output_queue.wait_and_pop(output_tensor); + return output_tensor; + } + + ml_image Inference::AllocImage(ml_image_info info, ml_access_mode access_mode) + { + auto image = m_model.CreateImage(info, access_mode); + + if (image == nullptr) + { + throw std::runtime_error("Couldn't not create image"); + } + + return image; + } + + void Inference::DoInference() + { + for (;;) + { + Image input; + m_input_queue.wait_and_pop(input); + + if (input.image == nullptr) + { + break; + } + + if (m_input_queue.size() > 0) + { + continue; + } + + Image output = { input.tag, AllocImage(m_output_info, ML_READ_WRITE) }; + + CheckStatus(mlInfer(m_model.GetModel(), input.image, output.image)); + + m_output_queue.push(std::move(output)); + } + } + + void Inference::Shutdown() + { + m_input_queue.push({0, nullptr}); + m_worker.join(); + } + } +} diff --git a/Baikal/PostEffects/ML/inference.h b/Baikal/PostEffects/ML/inference.h new file mode 100644 index 00000000..69b83393 --- /dev/null +++ b/Baikal/PostEffects/ML/inference.h @@ -0,0 +1,84 @@ +/********************************************************************** +Copyright (c) 2018 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +********************************************************************/ + + +#pragma once + +#include "PostEffects/ML/model_holder.h" + +#include "../RadeonRays/RadeonRays/src/async/thread_pool.h" +#include "image.h" + +#include +#include +#include + + +namespace Baikal +{ + namespace PostEffects + { + class Inference + { + public: + using Ptr = std::unique_ptr; + + Inference(std::string const& model_path, + size_t input_height, + size_t input_width, + float gpu_memory_fraction, + std::string const& visible_devices); + + virtual ~Inference(); + + ml_image_info GetInputInfo() const; + ml_image_info GetOutputInfo() const; + + Image GetInputData(); + + void PushInput(Image&& image); + // + // Try to pop output image. + // Returns empty image, if there are no infered element in the queue + Image TryPopOutput(); + // + // Wait and pop output image. + Image PopOutput(); + + protected: + void DoInference(); + ml_image AllocImage(ml_image_info info, ml_access_mode access_mode); + + RadeonRays::thread_safe_queue m_input_queue; + RadeonRays::thread_safe_queue m_output_queue; + + ModelHolder m_model; + ml_image_info m_input_info; + ml_image_info m_output_info; + + private: + void Shutdown(); + + std::thread m_worker; + }; + } +} diff --git a/Baikal/PostEffects/ML/ml_post_effect.cpp b/Baikal/PostEffects/ML/ml_post_effect.cpp new file mode 100644 index 00000000..0f295fb6 --- /dev/null +++ b/Baikal/PostEffects/ML/ml_post_effect.cpp @@ -0,0 +1,274 @@ +/********************************************************************** + Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved. + + Permission is hereby granted, free of charge, to any person obtaining a copy + of this software and associated documentation files (the "Software"), to deal + in the Software without restriction, including without limitation the rights + to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + copies of the Software, and to permit persons to whom the Software is + furnished to do so, subject to the following conditions: + + The above copyright notice and this permission notice shall be included in + all copies or substantial portions of the Software. + + THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + THE SOFTWARE. + ********************************************************************/ + +#include "PostEffects/ML/ml_post_effect.h" +#include "denoiser_preprocessor.h" +#include "upsampler_preprocessor.h" + +#ifdef BAIKAL_EMBED_KERNELS +#include "embed_kernels.h" +#endif + +namespace Baikal +{ + namespace PostEffects { + + using float3 = RadeonRays::float3; + using OutputType = Renderer::OutputType; + + MLPostEffect::MLPostEffect(CLWContext context, const CLProgramManager* program_manager, ModelType type) +#ifdef BAIKAL_EMBED_KERNELS + : ClwPostEffect(context, program_manager, "denoise", g_denoise_opencl, g_denoise_opencl_headers), +#else + : ClwPostEffect(context, program_manager, "../Baikal/Kernels/CL/denoise.cl") +#endif + , m_inference(nullptr) + , m_type(type) + , m_start_seq(0) + , m_last_seq(0) + , m_program(program_manager) + { + RegisterParameter("gpu_memory_fraction", .7f); + RegisterParameter("visible_devices", std::string()); + RegisterParameter("start_spp", 1u); + RegisterParameter("every_frame", 0u); + + // init preprocessing + switch (m_type) + { + case ModelType::kDenoiser: + m_preproc = std::make_unique(GetContext(), m_program); + break; + case ModelType::kUpsampler: + m_preproc = std::make_unique(GetContext(), m_program); + break; + default: + throw std::logic_error("unsupported model type"); + } + } + + Inference::Ptr MLPostEffect::CreateInference() + { + auto gpu_memory_fraction = GetParameter("gpu_memory_fraction").GetFloat(); + auto visible_devices = GetParameter("visible_devices").GetString(); + m_preproc->SetStartSpp(GetParameter("start_spp").GetUint()); + m_process_every_frame = static_cast(GetParameter("every_frame").GetUint()); + + switch (m_type) + { + case ModelType::kDenoiser: + return std::make_unique( + "models/color_albedo_depth_normal_9_v3.json", + m_input_height, + m_input_width, + gpu_memory_fraction, + visible_devices); + case ModelType::kUpsampler: + return std::unique_ptr( + new Inference("models/esrgan-03x2x32-273866.json", + m_input_height, + m_input_width, + gpu_memory_fraction, + visible_devices)); + default: + throw std::logic_error("Unsupported model type"); + } + } + + void MLPostEffect::Init(InputSet const& input_set, Output& output) + { + auto aov = input_set.begin()->second; + + m_input_width = aov->width(); + m_input_height = aov->height(); + + m_inference = CreateInference(); + auto out_shape = m_inference->GetOutputInfo(); + + m_last_image = CLWBuffer::Create(GetContext(), + CL_MEM_READ_WRITE, + out_shape.width * out_shape.height); + + m_host = std::vector(out_shape.width * out_shape.height); + } + + + void MLPostEffect::Apply(InputSet const& input_set, Output& output) + { + if (m_input_width != input_set.begin()->second->width() || + m_input_height != input_set.begin()->second->height()) + { + m_is_dirty = true; + } + + if (m_is_dirty) + { + Init(input_set, output); + m_is_dirty = false; + } + + auto clw_inference_output = dynamic_cast(&output); + + if (!clw_inference_output) + { + throw std::runtime_error("MLPostEffect::Apply(...): can not cast output"); + } + + auto context = GetContext(); + auto shape = m_inference->GetInputInfo(); + auto input = m_preproc->MakeInput(input_set); + + Image res; + if (m_process_every_frame) + { + m_inference->PushInput(std::move(input)); + res = m_inference->PopOutput(); + } + else + { + if (input.tag == 1) + { + m_start_seq = m_last_seq + 1; + } + + if (input.image != nullptr) + { + input.tag = ++m_last_seq; + m_inference->PushInput(std::move(input)); + } + + res = m_inference->TryPopOutput(); + } + + if (res.image != nullptr && (res.tag >= m_start_seq || m_process_every_frame)) + { + size_t res_size; + auto res_data = static_cast(mlMapImage(res.image, &res_size)); + + if (res_data == nullptr) + { + throw std::runtime_error("map input image is failed"); + } + + auto dest = m_host.data(); + auto source = res_data; + auto output_shape = m_inference->GetOutputInfo(); + for (auto i = 0u; i < output_shape.width * output_shape.height; ++i) + { + dest->x = *source++; + dest->y = *source++; + dest->z = *source++; + dest->w = 1; + ++dest; + } + + mlUnmapImage(res.image, res_data); + + context.WriteBuffer(0, + m_last_image, + m_host.data(), + res_size / (3 * sizeof(float))); + // Copy postprocessed image + context.CopyBuffer(0, + m_last_image, + clw_inference_output->data(), + 0 /* srcOffset */, + 0 /* destOffset */, + m_last_image.GetElementCount()).Wait(); + } + else + { + // Postprocessed image is not ready yet. + // Therefore, we'll output source (not-postprocessed) image + auto color = dynamic_cast(input_set.at(OutputType::kColor))->data(); + + if (m_type == ModelType::kDenoiser) + { + context.CopyBuffer(0, + color, + clw_inference_output->data(), + 0 /* srcOffset */, + 0 /* destOffset */, + shape.width * shape.height).Wait(); + } + else + { + Resize_2x(clw_inference_output->data(), color); + } + } + } + + void MLPostEffect::SetParameter(std::string const& name, Param value) + { + auto param = GetParameter(name); + PostEffect::SetParameter(name, value); + m_is_dirty = true; + } + + PostEffect::InputTypes MLPostEffect::GetInputTypes() const + { + return m_preproc->GetInputTypes(); + } + + void MLPostEffect::Resize_2x(CLWBuffer dst, CLWBuffer src) + { + auto context = GetContext(); + + if (m_resizer_cache.GetElementCount() < 2 * src.GetElementCount()) + { + m_resizer_cache = CLWBuffer::Create(context, + CL_MEM_READ_WRITE, + 2 * src.GetElementCount()); + } + + auto scale_x = GetKernel("BicubicUpscale2x_X"); + + unsigned argc = 0; + scale_x.SetArg(argc++, m_resizer_cache); + scale_x.SetArg(argc++, src); + scale_x.SetArg(argc++, m_input_width); + scale_x.SetArg(argc++, m_input_height); + + // run BicubicUpScale2x_X kernel + auto thread_num = ((2 * m_input_width * m_input_height + 63) / 64) * 64; + context.Launch1D(0, + thread_num, + 64, + scale_x); + + auto scale_y = GetKernel("BicubicUpscale2x_Y"); + + argc = 0; + scale_y.SetArg(argc++, dst); + scale_y.SetArg(argc++, m_resizer_cache); + scale_y.SetArg(argc++, 2 * m_input_width); + scale_y.SetArg(argc++, m_input_height); + + // run BicubicUpScale2x_Y kernel + thread_num = ((4 * m_input_width * m_input_height + 63) / 64) * 64; + context.Launch1D(0, + thread_num, + 64, + scale_y).Wait(); + } + } +} diff --git a/Baikal/PostEffects/ML/ml_post_effect.h b/Baikal/PostEffects/ML/ml_post_effect.h new file mode 100644 index 00000000..e0038170 --- /dev/null +++ b/Baikal/PostEffects/ML/ml_post_effect.h @@ -0,0 +1,75 @@ +/********************************************************************** +Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +********************************************************************/ + +#pragma once + +#include "data_preprocessor.h" +#include "PostEffects/ML/inference.h" +#include "PostEffects/clw_post_effect.h" + + +namespace Baikal +{ + namespace PostEffects + { + enum class ModelType + { + kDenoiser, + kUpsampler, + }; + + class MLPostEffect : public ClwPostEffect + { + public: + MLPostEffect(CLWContext context, const CLProgramManager* program_manager, ModelType type); + + void Apply(InputSet const& input_set, Output& output) override; + + void SetParameter(std::string const& name, Param value) override; + + InputTypes GetInputTypes() const override; + + void Resize_2x(CLWBuffer dst, CLWBuffer src); + private: + Inference::Ptr CreateInference(); + void Init(InputSet const& input_set, Output& output); + + Inference::Ptr m_inference; + ModelType m_type; + bool m_is_dirty = true; + bool m_process_every_frame = false; + bool m_has_postprocessed_image = false; + + std::vector m_host; + CLWBuffer m_last_image; + CLWBuffer m_resizer_cache; + std::unique_ptr m_preproc; + + std::uint32_t m_input_width = 0; + std::uint32_t m_input_height = 0; + + std::uint32_t m_start_seq = 0; + std::uint32_t m_last_seq = 0; + const CLProgramManager *m_program; + }; + } +} diff --git a/Baikal/PostEffects/ML/model_holder.cpp b/Baikal/PostEffects/ML/model_holder.cpp new file mode 100644 index 00000000..a28ac9c0 --- /dev/null +++ b/Baikal/PostEffects/ML/model_holder.cpp @@ -0,0 +1,72 @@ +/********************************************************************** +Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +********************************************************************/ + +#include "model_holder.h" + +#include +#include + +namespace Baikal +{ + namespace PostEffects + { + ModelHolder::ModelHolder(std::string const& model_path, + float gpu_memory_fraction, + std::string const& visible_devices) + : m_context(mlCreateContext(), mlReleaseContext) + , m_model(nullptr, nullptr) + { + if (m_context == nullptr) + { + throw std::runtime_error("can't create ml context"); + } + + ml_model_params params = {}; + params.model_path = model_path.c_str(); + params.gpu_memory_fraction = gpu_memory_fraction; + params.visible_devices = !visible_devices.empty() ? + visible_devices.c_str() : nullptr; + + m_model = Handle(mlCreateModel(m_context.get(), ¶ms), mlReleaseModel); + + if (m_model == nullptr) + { + throw std::runtime_error(mlGetLastError(nullptr)); + } + } + + ml_image ModelHolder::CreateImage(ml_image_info const& info, ml_access_mode access_mode) + { + auto tensor = mlCreateImage(m_context.get(), &info, access_mode); + + if (tensor == nullptr) + { + throw std::runtime_error("can not create model image"); + } + + return tensor; + } + + ModelHolder::~ModelHolder() + {} + } +} diff --git a/Baikal/PostEffects/ML/model_holder.h b/Baikal/PostEffects/ML/model_holder.h new file mode 100644 index 00000000..fd2583d2 --- /dev/null +++ b/Baikal/PostEffects/ML/model_holder.h @@ -0,0 +1,66 @@ +/********************************************************************** +Copyright (c) 2018 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +********************************************************************/ + +#pragma once + +#include "RadeonProML.h" + +#include +#include + + +namespace Baikal +{ + namespace PostEffects + { + // non copyable/movable + class ModelHolder + { + public: + ModelHolder(std::string const& model_path, + float gpu_memory_fraction, + std::string const& visible_devices); + + + const ml_model GetModel() + { + return m_model.get(); + } + + ml_image CreateImage(ml_image_info const& info, ml_access_mode access_mode); + + ~ModelHolder(); + + ModelHolder(const ModelHolder&) = delete; + ModelHolder(ModelHolder&&) = delete; + ModelHolder& operator = (const ModelHolder&) = delete; + ModelHolder& operator = (ModelHolder&&) = delete; + + private: + template + using Handle = std::unique_ptr::type, void (*)(T)>; + + Handle m_context; + Handle m_model; + }; + } +} diff --git a/Baikal/PostEffects/ML/upsampler_preprocessor.cpp b/Baikal/PostEffects/ML/upsampler_preprocessor.cpp new file mode 100644 index 00000000..c75ad344 --- /dev/null +++ b/Baikal/PostEffects/ML/upsampler_preprocessor.cpp @@ -0,0 +1,154 @@ +/********************************************************************** +Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +********************************************************************/ + +#include "upsampler_preprocessor.h" +#include "Output/clwoutput.h" +#include "CLWBuffer.h" + +namespace Baikal +{ + namespace PostEffects + { + using uint32_t = std::uint32_t; + using float3 = RadeonRays::float3; + + UpsamplerPreprocessor::UpsamplerPreprocessor(CLWContext context, + Baikal::CLProgramManager const *program_manager, + std::uint32_t start_spp) + : DataPreprocessor(context, program_manager, start_spp) + , m_context(mlCreateContext(), mlReleaseContext) + {} + + void UpsamplerPreprocessor::Init(std::uint32_t width, std::uint32_t height) + { + m_cache = CLWBuffer::Create(GetContext(), + CL_MEM_READ_WRITE, + width * height); + + m_input = CLWBuffer::Create(GetContext(), + CL_MEM_READ_WRITE, + 3 * width * height); + + ml_image_info image_info = {ML_FLOAT32, width, height, 3}; + m_image = mlCreateImage(m_context.get(), &image_info, ML_READ_WRITE); + + if (!m_image) + { + throw std::runtime_error("can not create ml_image"); + } + } + + std::tuple UpsamplerPreprocessor::ChannelsNum() const + { + return std::tuple(3, 3); + } + + Image UpsamplerPreprocessor::MakeInput(PostEffect::InputSet const& inputs) + { + auto color_aov = inputs.begin()->second; + + if (!m_is_init) + { + m_width = color_aov->width(); + m_height = color_aov->height(); + Init(m_width, m_height); + m_is_init = true; + } + + auto clw_input = dynamic_cast(color_aov); + + if (clw_input == nullptr) + { + throw std::runtime_error("UpsamplerPreprocessor::MakeInput(..): incorrect input"); + } + + auto context = GetContext(); + + // read spp from first pixel as 4th channel + auto sample_count = ReadSpp(clw_input->data()); + + if (m_start_spp > sample_count) + { + return Image(sample_count, nullptr); + } + + ApplyToneMapping(m_cache, clw_input->data()); + + // delete 4th channel + WriteToInputs(m_input, + CLWBuffer::CreateFromClBuffer(m_cache), + m_width, + m_height, + 0, // dst channels offset + 3, // dst channels num + 0, // src channels offset + 4, // src channels num + 3); // channels to copy + + size_t image_size = 0; + auto host_buffer = mlMapImage(m_image, &image_size); + + if (!host_buffer || image_size == 0) + { + throw std::runtime_error("map operation failed"); + } + + context.ReadBuffer(0, + m_input, + static_cast(host_buffer), + m_input.GetElementCount()).Wait(); + + if (mlUnmapImage(m_image, host_buffer) != ML_OK) + { + throw std::runtime_error("unmap operation failed"); + } + + return Image(sample_count, m_image); + } + + void UpsamplerPreprocessor::ApplyToneMapping(CLWBuffer const& dst, + CLWBuffer const& src) + { + assert (dst.GetElementCount() >= src.GetElementCount()); + + auto tonemapping = GetKernel("ToneMappingExponential"); + + // Set kernel parameters + unsigned argc = 0; + tonemapping.SetArg(argc++, dst); + tonemapping.SetArg(argc++, src); + tonemapping.SetArg(argc++, (int)src.GetElementCount()); + + // run DivideBySampleCount kernel + auto thread_num = ((src.GetElementCount() + 63) / 64) * 64; + GetContext().Launch1D(0, + thread_num, + 64, + tonemapping); + } + + std::set UpsamplerPreprocessor::GetInputTypes() const + { + return std::set({Renderer::OutputType::kColor}); + } + } +} diff --git a/Baikal/PostEffects/ML/upsampler_preprocessor.h b/Baikal/PostEffects/ML/upsampler_preprocessor.h new file mode 100644 index 00000000..d6861606 --- /dev/null +++ b/Baikal/PostEffects/ML/upsampler_preprocessor.h @@ -0,0 +1,65 @@ +/********************************************************************** + Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved. + + Permission is hereby granted, free of charge, to any person obtaining a copy + of this software and associated documentation files (the "Software"), to deal + in the Software without restriction, including without limitation the rights + to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + copies of the Software, and to permit persons to whom the Software is + furnished to do so, subject to the following conditions: + + The above copyright notice and this permission notice shall be included in + all copies or substantial portions of the Software. + + THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + THE SOFTWARE. + ********************************************************************/ + +#pragma once + +#ifdef BAIKAL_EMBED_KERNELS +#include "embed_kernels.h" +#endif + +#include "data_preprocessor.h" + + +namespace Baikal +{ + namespace PostEffects + { + class UpsamplerPreprocessor : public DataPreprocessor + { + public: + UpsamplerPreprocessor(CLWContext context, + Baikal::CLProgramManager const *program_manager, + std::uint32_t spp = 1); + + + Image MakeInput(PostEffect::InputSet const& inputs) override; + + std::set GetInputTypes() const override; + + std::tuple ChannelsNum() const override; + + private: + void Init(std::uint32_t width, std::uint32_t height); + + void ApplyToneMapping(CLWBuffer const& dst, + CLWBuffer const& src); + + bool m_is_init = false; + std::uint32_t m_width, m_height; + CLWBuffer m_input; + CLWBuffer m_resizer_cache; + CLWBuffer m_cache; + Handle m_context; + ml_image m_image; + }; + } +} diff --git a/Baikal/PostEffects/bilateral_denoiser.h b/Baikal/PostEffects/bilateral_denoiser.h index 595f6eee..e9e2db0d 100644 --- a/Baikal/PostEffects/bilateral_denoiser.h +++ b/Baikal/PostEffects/bilateral_denoiser.h @@ -20,7 +20,8 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. ********************************************************************/ #pragma once -#include "clw_post_effect.h" + +#include "PostEffects/clw_post_effect.h" #ifdef BAIKAL_EMBED_KERNELS #include "embed_kernels.h" @@ -43,18 +44,30 @@ namespace Baikal * kColor * kWorldShadingNormal * kWorldPosition + * kAlbedo */ class BilateralDenoiser : public ClwPostEffect { public: // Constructor - BilateralDenoiser(CLWContext context, const CLProgramManager *program_manager); + BilateralDenoiser(CLWContext context, const CLProgramManager* program_manager); // Apply filter void Apply(InputSet const& input_set, Output& output) override; + InputTypes GetInputTypes() const override + { + return std::set( + { + Renderer::OutputType::kColor, + Renderer::OutputType::kWorldShadingNormal, + Renderer::OutputType::kWorldPosition, + Renderer::OutputType::kAlbedo + }); + } + private: // Find required output - ClwOutput* FindOutput(InputSet const& input_set, Renderer::OutputType type); + static ClwOutput* FindOutput(InputSet const& input_set, Renderer::OutputType type); CLWProgram m_program; }; @@ -67,11 +80,11 @@ namespace Baikal #endif { // Add necessary params - RegisterParameter("radius", RadeonRays::float4(5.f, 0.f, 0.f, 0.f)); - RegisterParameter("color_sensitivity", RadeonRays::float4(5.f, 0.f, 0.f, 0.f)); - RegisterParameter("position_sensitivity", RadeonRays::float4(5.f, 0.f, 0.f, 0.f)); - RegisterParameter("normal_sensitivity", RadeonRays::float4(0.1f, 0.f, 0.f, 0.f)); - RegisterParameter("albedo_sensitivity", RadeonRays::float4(0.1f, 0.f, 0.f, 0.f)); + RegisterParameter("radius", 5.f); + RegisterParameter("color_sensitivity", 5.f); + RegisterParameter("position_sensitivity", 5.f); + RegisterParameter("normal_sensitivity", 0.1f); + RegisterParameter("albedo_sensitivity", 0.1f); } inline ClwOutput* BilateralDenoiser::FindOutput(InputSet const& input_set, Renderer::OutputType type) @@ -88,11 +101,11 @@ namespace Baikal inline void BilateralDenoiser::Apply(InputSet const& input_set, Output& output) { - auto radius = static_cast(GetParameter("radius").x); - auto sigma_color = GetParameter("color_sensitivity").x; - auto sigma_position = GetParameter("position_sensitivity").x; - auto sigma_normal = GetParameter("normal_sensitivity").x; - auto sigma_albedo = GetParameter("albedo_sensitivity").x; + auto radius = static_cast(GetParameter("radius").GetFloat()); + auto sigma_color = GetParameter("color_sensitivity").GetFloat(); + auto sigma_position = GetParameter("position_sensitivity").GetFloat(); + auto sigma_normal = GetParameter("normal_sensitivity").GetFloat(); + auto sigma_albedo = GetParameter("albedo_sensitivity").GetFloat(); auto color = FindOutput(input_set, Renderer::OutputType::kColor); auto normal = FindOutput(input_set, Renderer::OutputType::kWorldShadingNormal); @@ -118,12 +131,9 @@ namespace Baikal denoise_kernel.SetArg(argc++, out_color->data()); // Run shading kernel - { - size_t gs[] = { static_cast((output.width() + 7) / 8 * 8), static_cast((output.height() + 7) / 8 * 8) }; - size_t ls[] = { 8, 8 }; + size_t gs[] = { static_cast((output.width() + 7) / 8 * 8), static_cast((output.height() + 7) / 8 * 8) }; + size_t ls[] = { 8, 8 }; - GetContext().Launch2D(0, gs, ls, denoise_kernel); - } + GetContext().Launch2D(0, gs, ls, denoise_kernel); } - } diff --git a/Baikal/PostEffects/clw_post_effect.h b/Baikal/PostEffects/clw_post_effect.h index 01d7be77..afab538f 100644 --- a/Baikal/PostEffects/clw_post_effect.h +++ b/Baikal/PostEffects/clw_post_effect.h @@ -21,15 +21,16 @@ THE SOFTWARE. ********************************************************************/ #pragma once -#include "post_effect.h" - -#include "CLW.h" +#include "PostEffects/post_effect.h" #include "Output/clwoutput.h" #include "Utils/clw_class.h" +#include "CLW.h" + #include #include + namespace Baikal { class CLProgramManager; diff --git a/Baikal/PostEffects/post_effect.cpp b/Baikal/PostEffects/post_effect.cpp new file mode 100644 index 00000000..44a85b8a --- /dev/null +++ b/Baikal/PostEffects/post_effect.cpp @@ -0,0 +1,134 @@ +/********************************************************************** +Copyright (c) 2016 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +********************************************************************/ + +#include "post_effect.h" + +namespace Baikal +{ + + PostEffect::Param::Param(std::uint32_t value) + : m_type(ParamType::kUint), + m_uint_value(value) + { } + + PostEffect::Param::Param(float value) + : m_type(ParamType::kFloat), + m_float_value(value) + { } + + PostEffect::Param::Param(RadeonRays::float2 const& value) + : m_type(ParamType::kFloat2), + m_float2_value(value) + { } + + PostEffect::Param::Param(RadeonRays::float4 const& value) + : m_type(ParamType::kFloat4), + m_float4_value(value) + { } + + PostEffect::Param::Param(std::string const& value) + : m_type(ParamType::kString), + m_str_value(value) + { } + + PostEffect::ParamType PostEffect::Param::GetType() const + { + return m_type; + } + + void PostEffect::Param::AssertType(ParamType type) const + { + if (m_type != type) + { + throw std::runtime_error("Attempt to get incorrect param type value"); + } + } + + float PostEffect::Param::GetFloat() const + { + AssertType(ParamType::kFloat); + return m_float_value; + } + + std::uint32_t PostEffect::Param::GetUint() const + { + AssertType(ParamType::kUint); + return m_uint_value; + } + + const RadeonRays::float4& PostEffect::Param::GetFloat4() const + { + AssertType(ParamType::kFloat4); + return m_float4_value; + } + + const RadeonRays::float2& PostEffect::Param::GetFloat2() const + { + AssertType(ParamType::kFloat2); + return m_float2_value; + } + + const std::string& PostEffect::Param::GetString() const + { + AssertType(ParamType::kString); + return m_str_value; + } + + const PostEffect::Param& PostEffect::GetParameter(std::string const& name) const + { + auto iter = m_parameters.find(name); + + if (iter == m_parameters.cend()) + { + throw std::runtime_error("PostEffect: no such parameter " + name); + } + + return iter->second; + } + + void PostEffect::SetParameter(std::string const& name, Param value) + { + auto iter = m_parameters.find(name); + + if (iter == m_parameters.cend()) + { + throw std::runtime_error("PostEffect: no such parameter " + name); + } + + if (value.GetType() != iter->second.GetType()) + { + throw std::runtime_error("PostEffect: attemp to change type of registred parameter " + name); + } + + iter->second = std::move(value); + } + + void PostEffect::RegisterParameter(std::string const &name, Param initial_value) + { + if (m_parameters.find(name) != m_parameters.cend()) + { + throw std::runtime_error("Attempt to register already existing name"); + } + + m_parameters.emplace(name, initial_value); + } +} diff --git a/Baikal/PostEffects/post_effect.h b/Baikal/PostEffects/post_effect.h index 62becbd1..bf48aba3 100644 --- a/Baikal/PostEffects/post_effect.h +++ b/Baikal/PostEffects/post_effect.h @@ -25,12 +25,16 @@ THE SOFTWARE. #include "Output/output.h" #include +#include #include #include #include + namespace Baikal { + class Camera; + /** \brief Interface for post-processing effects. @@ -45,59 +49,78 @@ namespace Baikal class PostEffect { public: - // Data type to pass all necessary content into the post effect. + + enum class ParamType + { + kFloat = 0, + kUint, + kFloat2, + kFloat4, + kString, + }; + + class Param + { + public: + ParamType GetType() const; + + float GetFloat() const; + std::uint32_t GetUint() const; + const RadeonRays::float2& GetFloat2() const; + const RadeonRays::float4& GetFloat4() const; + const std::string& GetString() const; + + Param(float value); + Param(std::uint32_t value); + Param(RadeonRays::float2 const& value); + Param(RadeonRays::float4 const& value); + Param(std::string const& value); + + operator float() const { return GetFloat(); } + operator std::uint32_t() const { return GetUint(); } + operator const RadeonRays::float2&() const { return GetFloat2(); } + operator const RadeonRays::float4&() const { return GetFloat4(); } + operator const std::string&() const { return GetString(); } + + private: + void AssertType(ParamType type) const; + + ParamType m_type; + union { + std::uint32_t m_uint_value; + float m_float_value; + RadeonRays::float2 m_float2_value; + RadeonRays::float4 m_float4_value; + }; + std::string m_str_value; + }; + + // Data type to pass all necessary content into the post effect. using InputSet = std::map; + // Specification of the input set types + using InputTypes = std::set; + // Default constructor & destructor PostEffect() = default; virtual ~PostEffect() = default; + virtual InputTypes GetInputTypes() const = 0; + // Apply post effect and use output for the result virtual void Apply(InputSet const& input_set, Output& output) = 0; - // Set scalar parameter - void SetParameter(std::string const& name, RadeonRays::float4 const& value); + virtual void SetParameter(std::string const& name, Param value); - // Get scalar parameter - RadeonRays::float4 GetParameter(std::string const& name) const; + const Param& GetParameter(std::string const& name) const; protected: - // Adds scalar parameter into the parameter map - void RegisterParameter(std::string const& name, RadeonRays::float4 const& initial_value); + + void RegisterParameter(std::string const& name, Param init_value); private: + // Parameter map - std::map m_parameters; + std::map m_parameters; }; - - inline void PostEffect::SetParameter(std::string const& name, RadeonRays::float4 const& value) - { - auto iter = m_parameters.find(name); - - if (iter == m_parameters.cend()) - { - throw std::runtime_error("PostEffect: no such parameter " + name); - } - - iter->second = value; - } - - inline RadeonRays::float4 PostEffect::GetParameter(std::string const& name) const - { - auto iter = m_parameters.find(name); - - if (iter == m_parameters.cend()) - { - throw std::runtime_error("PostEffect: no such parameter " + name); - } - - return iter->second; - } - - inline void PostEffect::RegisterParameter(std::string const& name, RadeonRays::float4 const& initial_value) - { - assert(m_parameters.find(name) == m_parameters.cend()); - - m_parameters.emplace(name, initial_value); - } } diff --git a/Baikal/PostEffects/wavelet_denoiser.h b/Baikal/PostEffects/wavelet_denoiser.h index 4c801271..f3d02d01 100644 --- a/Baikal/PostEffects/wavelet_denoiser.h +++ b/Baikal/PostEffects/wavelet_denoiser.h @@ -20,19 +20,22 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. ********************************************************************/ #pragma once -#include "clw_post_effect.h" -#include -#include -#include -#include "AreaMap33.h" +#include "PostEffects/clw_post_effect.h" +#include "PostEffects/AreaMap33.h" +#include "SceneGraph/camera.h" + +#include "math/matrix.h" +#include "math/mathutils.h" #include +#include #ifdef BAIKAL_EMBED_KERNELS #include "embed_kernels.h" #endif + namespace Baikal { /** @@ -58,9 +61,21 @@ namespace Baikal // Constructor WaveletDenoiser(CLWContext context, const CLProgramManager *program_manager); virtual ~WaveletDenoiser(); + + InputTypes GetInputTypes() const override + { + return std::set( + { + Renderer::OutputType::kColor, + Renderer::OutputType::kWorldShadingNormal, + Renderer::OutputType::kWorldPosition, + Renderer::OutputType::kAlbedo, + Renderer::OutputType::kMeshID + }); + } + // Apply filter void Apply(InputSet const& input_set, Output& output) override; - void Update(PerspectiveCamera* camera); private: // Find required output @@ -103,7 +118,7 @@ namespace Baikal bool m_buffers_initialized; }; - inline WaveletDenoiser::WaveletDenoiser(CLWContext context, const CLProgramManager *program_manager) + inline WaveletDenoiser::WaveletDenoiser(CLWContext context, const CLProgramManager* program_manager) #ifdef BAIKAL_EMBED_KERNELS : ClwPostEffect(context, program_manager, "wavelet_denoise", g_wavelet_denoise_opencl, g_wavelet_denoise_opencl_headers) #else @@ -116,9 +131,18 @@ namespace Baikal , m_buffers_initialized(false) { // Add necessary params - RegisterParameter("color_sensitivity", RadeonRays::float4(0.07f, 0.f, 0.f, 0.f)); - RegisterParameter("position_sensitivity", RadeonRays::float4(0.03f, 0.f, 0.f, 0.f)); - RegisterParameter("normal_sensitivity", RadeonRays::float4(0.01f, 0.f, 0.f, 0.f)); + RegisterParameter("color_sensitivity", 0.07f); + RegisterParameter("position_sensitivity", 0.03f); + RegisterParameter("normal_sensitivity", 0.01f); + + RegisterParameter("camera_focal_length", 0.f); + RegisterParameter("camera_sensor_size", RadeonRays::float2()); + RegisterParameter("camera_depth_range", RadeonRays::float2()); + RegisterParameter("camera_up_vector", RadeonRays::float3()); + RegisterParameter("camera_forward_vector", RadeonRays::float3()); + RegisterParameter("camera_right_vector", RadeonRays::float3()); + RegisterParameter("camera_position", RadeonRays::float3()); + RegisterParameter("camera_aspect_ratio", 1.f); for (uint32_t buffer_index = 0; buffer_index < m_num_tmp_buffers; buffer_index++) { @@ -199,11 +223,41 @@ namespace Baikal inline void WaveletDenoiser::Apply(InputSet const& input_set, Output& output) { + m_prev_view_proj = m_view_proj; + + const float focal_length = GetParameter("camera_focal_length"); + const float2 sensor_size = GetParameter("camera_sensor_size"); + float2 z_range = GetParameter("camera_depth_range"); + + // Nan-avoidance in perspective matrix + z_range.x = std::max(z_range.x, std::numeric_limits::epsilon()); + + const float fovy = std::atan(sensor_size.y / (2.0f * focal_length)); + + const float3 up = GetParameter("camera_up_vector"); + const float3 right = GetParameter("camera_right_vector"); + const float3 forward = GetParameter("camera_forward_vector"); + const float3 pos = GetParameter("camera_position"); + const float aspect_ratio = GetParameter("camera_aspect_ratio"); + + const matrix proj = perspective_proj_fovy_rh_gl(fovy, aspect_ratio, z_range.x, z_range.y); + const float3 ip = float3(-dot(right, pos), -dot(up, pos), -dot(forward, pos)); + + const matrix view = matrix(right.x, right.y, right.z, ip.x, + up.x, up.y, up.z, ip.y, + forward.x, forward.y, forward.z, ip.z, + 0.0f, 0.0f, 0.0f, 1.0f); + + m_view_proj = proj * view; + + GetContext().WriteBuffer(0, m_view_proj_buffer, &m_view_proj.m[0][0], 16).Wait(); + GetContext().WriteBuffer(0, m_prev_view_proj_buffer, &m_prev_view_proj.m[0][0], 16).Wait(); + uint32_t prev_buffer_index = m_current_buffer_index; m_current_buffer_index = (m_current_buffer_index + 1) % m_num_tmp_buffers; - auto sigma_color = GetParameter("color_sensitivity").x; - auto sigma_position = GetParameter("position_sensitivity").x; + auto sigma_color = GetParameter("color_sensitivity").GetFloat(); + auto sigma_position = GetParameter("position_sensitivity").GetFloat(); auto color = FindOutput(input_set, Renderer::OutputType::kColor); auto normal = FindOutput(input_set, Renderer::OutputType::kWorldShadingNormal); @@ -523,37 +577,4 @@ namespace Baikal } } } - - inline void WaveletDenoiser::Update(PerspectiveCamera* camera) - { - m_prev_view_proj = m_view_proj; - - const float focal_length = camera->GetFocalLength(); - const RadeonRays::float2 sensor_size = camera->GetSensorSize(); - - RadeonRays::float2 z_range = camera->GetDepthRange(); - - // Nan-avoidance in perspective matrix - z_range.x = std::max(z_range.x, std::numeric_limits::epsilon()); - - const float fovy = atan(sensor_size.y / (2.0f * focal_length)); - - const RadeonRays::float3 up = camera->GetUpVector(); - const RadeonRays::float3 right = -camera->GetRightVector(); - const RadeonRays::float3 forward = camera->GetForwardVector(); - const RadeonRays::float3 pos = camera->GetPosition(); - - const RadeonRays::matrix proj = RadeonRays::perspective_proj_fovy_rh_gl(fovy, camera->GetAspectRatio(), z_range.x, z_range.y); - const RadeonRays::float3 ip = RadeonRays::float3(-dot(right, pos), -dot(up, pos), -dot(forward, pos)); - - const RadeonRays::matrix view = RadeonRays::matrix(right.x, right.y, right.z, ip.x, - up.x, up.y, up.z, ip.y, - forward.x, forward.y, forward.z, ip.z, - 0.0f, 0.0f, 0.0f, 1.0f); - - m_view_proj = proj * view; - - GetContext().WriteBuffer(0, m_view_proj_buffer, &m_view_proj.m[0][0], 16).Wait(); - GetContext().WriteBuffer(0, m_prev_view_proj_buffer, &m_prev_view_proj.m[0][0], 16).Wait(); - } } diff --git a/Baikal/RenderFactory/clw_render_factory.cpp b/Baikal/RenderFactory/clw_render_factory.cpp index 5c8b4d8c..24d4be01 100644 --- a/Baikal/RenderFactory/clw_render_factory.cpp +++ b/Baikal/RenderFactory/clw_render_factory.cpp @@ -4,14 +4,15 @@ #include "Renderers/monte_carlo_renderer.h" #include "Renderers/adaptive_renderer.h" #include "Estimators/path_tracing_estimator.h" +#include "Controllers/scene_controller.h" -#ifdef ENABLE_DENOISER #include "PostEffects/bilateral_denoiser.h" #include "PostEffects/wavelet_denoiser.h" -#endif +#include "PostEffects/ML/ml_post_effect.h" #include + namespace Baikal { ClwRenderFactory::ClwRenderFactory(CLWContext context, std::string const& cache_path) @@ -31,8 +32,7 @@ namespace Baikal } // Create a renderer of specified type - std::unique_ptr ClwRenderFactory::CreateRenderer( - RendererType type) const + std::unique_ptr ClwRenderFactory::CreateRenderer(RendererType type) const { switch (type) { @@ -55,24 +55,25 @@ namespace Baikal return std::unique_ptr(new ClwOutput(m_context, w, h)); } - std::unique_ptr ClwRenderFactory::CreatePostEffect( - PostEffectType type) const + std::unique_ptr ClwRenderFactory::CreatePostEffect(PostEffectType type) const { -#ifdef ENABLE_DENOISER switch (type) { case PostEffectType::kBilateralDenoiser: - return std::unique_ptr( - new BilateralDenoiser(m_context, &m_program_manager)); + return std::make_unique(m_context, &m_program_manager); case PostEffectType::kWaveletDenoiser: - return std::unique_ptr( - new WaveletDenoiser(m_context, &m_program_manager)); + return std::make_unique(m_context, &m_program_manager); + case PostEffectType::kMLDenoiser: + return std::make_unique(m_context, + &m_program_manager, + PostEffects::ModelType::kDenoiser); + case PostEffectType::kMLUpsampler: + return std::make_unique(m_context, + &m_program_manager, + PostEffects::ModelType::kUpsampler); default: throw std::runtime_error("PostEffect is not supported"); } -#else - throw std::runtime_error("PostEffect is not supported"); -#endif } std::unique_ptr> ClwRenderFactory::CreateSceneController() const diff --git a/Baikal/RenderFactory/clw_render_factory.h b/Baikal/RenderFactory/clw_render_factory.h index ef995b95..ae5ca47f 100644 --- a/Baikal/RenderFactory/clw_render_factory.h +++ b/Baikal/RenderFactory/clw_render_factory.h @@ -22,13 +22,13 @@ ********************************************************************/ #pragma once -#include "render_factory.h" +#include "RenderFactory/render_factory.h" +#include "Utils/cl_program_manager.h" +#include "SceneGraph/clwscene.h" + #include "radeon_rays_cl.h" #include "CLW.h" -#include "SceneGraph/clwscene.h" -#include "Utils/cl_program_manager.h" - #include #include @@ -54,8 +54,7 @@ namespace Baikal std::unique_ptr CreateOutput(std::uint32_t w, std::uint32_t h) const override; // Create post effect of specified type - std::unique_ptr - CreatePostEffect(PostEffectType type) const override; + std::unique_ptr CreatePostEffect(PostEffectType type) const override; std::unique_ptr> CreateSceneController() const override; diff --git a/Baikal/RenderFactory/render_factory.h b/Baikal/RenderFactory/render_factory.h index 29ebdbdc..f0f7986f 100644 --- a/Baikal/RenderFactory/render_factory.h +++ b/Baikal/RenderFactory/render_factory.h @@ -31,7 +31,15 @@ namespace Baikal class Renderer; class Output; class PostEffect; - + + enum class PostEffectType + { + kBilateralDenoiser, + kWaveletDenoiser, + kMLDenoiser, + kMLUpsampler, + }; + /** \brief RenderFactory class is in charge of render entities creation. @@ -47,12 +55,6 @@ namespace Baikal { kUnidirectionalPathTracer }; - - enum class PostEffectType - { - kBilateralDenoiser, - kWaveletDenoiser - }; RenderFactory() = default; virtual ~RenderFactory() = default; @@ -72,6 +74,4 @@ namespace Baikal RenderFactory(RenderFactory const&) = delete; RenderFactory const& operator = (RenderFactory const&) = delete; }; - - } diff --git a/Baikal/Renderers/monte_carlo_renderer.cpp b/Baikal/Renderers/monte_carlo_renderer.cpp index 3f99b0ba..424cb79c 100644 --- a/Baikal/Renderers/monte_carlo_renderer.cpp +++ b/Baikal/Renderers/monte_carlo_renderer.cpp @@ -369,6 +369,11 @@ namespace Baikal return GetKernel("ApplyGammaAndCopyData"); } + CLWKernel MonteCarloRenderer::GetCopySplitKernel() + { + return GetKernel("ApplyGammaAndCopySplitData"); + } + CLWKernel MonteCarloRenderer::GetAccumulateKernel() { return GetKernel("AccumulateData"); @@ -422,5 +427,5 @@ namespace Baikal GetContext().Launch1D(0, ((size + 63) / 64) * 64, 64, misskernel); } } - + } diff --git a/Baikal/Renderers/monte_carlo_renderer.h b/Baikal/Renderers/monte_carlo_renderer.h index 76792d83..6d8a0c78 100644 --- a/Baikal/Renderers/monte_carlo_renderer.h +++ b/Baikal/Renderers/monte_carlo_renderer.h @@ -70,8 +70,10 @@ namespace Baikal void SetRandomSeed(std::uint32_t seed) override; - // Interop function + // Interop functions CLWKernel GetCopyKernel(); + CLWKernel GetCopySplitKernel(); + // Add function CLWKernel GetAccumulateKernel(); // Run render benchmark diff --git a/BaikalDataGenerator/CMakeLists.txt b/BaikalDataGenerator/CMakeLists.txt index da25701e..dc2fae58 100644 --- a/BaikalDataGenerator/CMakeLists.txt +++ b/BaikalDataGenerator/CMakeLists.txt @@ -25,9 +25,13 @@ target_compile_features(BaikalDataGenerator PRIVATE cxx_std_17) target_include_directories(BaikalDataGenerator PRIVATE ${Baikal_SOURCE_DIR} PRIVATE .) - + target_link_libraries(BaikalDataGenerator PRIVATE Baikal BaikalIO) +if (NOT MSVC) + target_link_libraries(BaikalDataGenerator PRIVATE stdc++fs) +endif () + set_target_properties(BaikalDataGenerator PROPERTIES VS_DEBUGGER_WORKING_DIRECTORY ${Baikal_SOURCE_DIR}/BaikalDataGenerator) diff --git a/BaikalDataGenerator/Source/render.cpp b/BaikalDataGenerator/Source/render.cpp index 64a4867d..5da2c94d 100644 --- a/BaikalDataGenerator/Source/render.cpp +++ b/BaikalDataGenerator/Source/render.cpp @@ -20,6 +20,12 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. ********************************************************************/ +#include "CLW.h" +#include "Controllers/scene_controller.h" +#include "Renderers/renderer.h" +#include "RenderFactory/clw_render_factory.h" +#include "SceneGraph/camera.h" +#include "scene_io.h" #include "input_info.h" #include "logging.h" #include "material_io.h" diff --git a/BaikalIO/scene_io.cpp b/BaikalIO/scene_io.cpp index 846f63c9..54dc4392 100644 --- a/BaikalIO/scene_io.cpp +++ b/BaikalIO/scene_io.cpp @@ -81,7 +81,7 @@ namespace Baikal m_texture_cache[name] = texture; return texture; } - catch (std::runtime_error &) + catch (std::runtime_error&) { LogInfo("Missing texture: ", name, "\n"); return nullptr; diff --git a/BaikalIO/scene_obj_io.cpp b/BaikalIO/scene_obj_io.cpp index 48d32baf..66c30637 100644 --- a/BaikalIO/scene_obj_io.cpp +++ b/BaikalIO/scene_obj_io.cpp @@ -224,8 +224,8 @@ namespace Baikal // TODO: temporary code to add directional light auto light = DirectionalLight::Create(); - light->SetDirection(RadeonRays::float3(.1f, -1.f, -.1f)); - light->SetEmittedRadiance(RadeonRays::float3(1.f, 1.f, 1.f)); + light->SetDirection(RadeonRays::float3(-1.f, -1.f, -1.f)); + light->SetEmittedRadiance(RadeonRays::float3(10.f, 8.f, 6.f)); scene->AttachLight(light); diff --git a/BaikalStandalone/Application/app_utils.cpp b/BaikalStandalone/Application/app_utils.cpp index ac293590..b775838f 100644 --- a/BaikalStandalone/Application/app_utils.cpp +++ b/BaikalStandalone/Application/app_utils.cpp @@ -39,6 +39,8 @@ namespace Baikal { AppSettings s; + s.help = m_cmd_parser.OptionExists("-help"); + s.path = m_cmd_parser.GetOption("-p", s.path); s.modelname = m_cmd_parser.GetOption("-f", s.modelname); @@ -58,6 +60,8 @@ namespace Baikal s.num_bounces = m_cmd_parser.GetOption("-nb", s.num_bounces); + s.split_output = m_cmd_parser.OptionExists("-split"); + s.camera_pos.x = m_cmd_parser.GetOption("-cpx", s.camera_pos.x); s.camera_pos.y = m_cmd_parser.GetOption("-cpy", s.camera_pos.y); @@ -90,6 +94,51 @@ namespace Baikal s.image_file_format = m_cmd_parser.GetOption("-iff", s.image_file_format); + s.light_file = m_cmd_parser.GetOption("-lights", s.light_file); + + s.gpu_mem_fraction = m_cmd_parser.GetOption("-gmf", s.gpu_mem_fraction); + + s.visible_devices = m_cmd_parser.GetOption("-vds", s.visible_devices); + + auto has_primary_device = [](std::string const& str) + { + if (str.empty()) + { + return true; + } + + std::stringstream ss(str); + + while (ss.good()) + { + std::string substr; + std::getline( ss, substr, ',' ); + + if (substr == "0") + { + return true; + } + } + return false; + }; + + float max_gmf = 1.f; + if (has_primary_device(s.visible_devices)) + { + max_gmf = .5f; + } + + if (s.gpu_mem_fraction < 0.0f) + { + std::cout << "WARNING: '-gmf' option value clamped to zero" << std::endl; + s.gpu_mem_fraction = 0.0f; + } + else if (s.gpu_mem_fraction > max_gmf) + { + std::cout << "WARNING: '-gmf' option value clamped to one or 0.5 in case primary device" << std::endl; + s.gpu_mem_fraction = 1.f; + } + if (m_cmd_parser.OptionExists("-ct")) { auto camera_type = m_cmd_parser.GetOption("-ct"); @@ -111,15 +160,15 @@ namespace Baikal auto cfg = m_cmd_parser.GetOption("-config"); if (cfg == "cpu") - s.mode = ConfigManager::Mode::kUseSingleCpu; + s.mode = Mode::kUseSingleCpu; else if (cfg == "gpu") - s.mode = ConfigManager::Mode::kUseSingleGpu; + s.mode = Mode::kUseSingleGpu; else if (cfg == "mcpu") - s.mode = ConfigManager::Mode::kUseCpus; + s.mode = Mode::kUseCpus; else if (cfg == "mgpu") - s.mode = ConfigManager::Mode::kUseGpus; + s.mode = Mode::kUseGpus; else if (cfg == "all") - s.mode = ConfigManager::Mode::kUseAll; + s.mode = Mode::kUseAll; } s.platform_index = m_cmd_parser.GetOption("-platform", s.platform_index); @@ -142,6 +191,41 @@ namespace Baikal s.cmd_line_mode = true; } + if (m_cmd_parser.OptionExists("-postprocess")) + { + auto post_processing_type = m_cmd_parser.GetOption("-postprocess"); + + if (post_processing_type == "denoise-bilateral") + { + s.postprocess_type = PostProcessingType::kBilateralDenoiser; + } + else if (post_processing_type == "denoise-wavelet") + { + s.postprocess_type = PostProcessingType::kWaveletDenoser; + } + else if (post_processing_type == "denoise-ml") + { + s.postprocess_type = PostProcessingType::kMLDenoiser; + } + else if (post_processing_type == "upsample-ml-2x") + { + s.postprocess_type = PostProcessingType::kMLUpsample; + } + else + { + std::cerr << "WARNING: unknown postprocessing type" << std::endl; + } + } + + s.postprocess_every_frame = m_cmd_parser.GetOption("-every_frame", s.postprocess_every_frame); + if (s.postprocess_every_frame) + { + s.postprocess_start_spp = 1; + } + else + { + s.postprocess_start_spp = m_cmd_parser.GetOption("-start_spp", s.postprocess_start_spp); + } return s; } @@ -151,7 +235,8 @@ namespace Baikal } AppSettings::AppSettings() - : path("../Resources/CornellBox") + : help(false) + , path("../Resources/CornellBox") , modelname("orig.objm") , envmapname("../Resources/Textures/studio015.hdr") //render @@ -161,7 +246,7 @@ namespace Baikal , num_samples(-1) , interop(true) , cspeed(10.25f) - , mode(ConfigManager::Mode::kUseSingleGpu) + , mode(Mode::kUseSingleGpu) //ao , ao_radius(1.f) , num_ao_rays(1) @@ -202,4 +287,4 @@ namespace Baikal , device_index(-1) { } -} \ No newline at end of file +} diff --git a/BaikalStandalone/Application/app_utils.h b/BaikalStandalone/Application/app_utils.h index 0b82a241..57c376fe 100644 --- a/BaikalStandalone/Application/app_utils.h +++ b/BaikalStandalone/Application/app_utils.h @@ -36,26 +36,29 @@ namespace Baikal { AppSettings(); - //model load settings + bool help; + + // model load settings std::string path; std::string modelname; std::string envmapname; - //render - int width; - int height; + // render + std::uint32_t width; + std::uint32_t height; int num_bounces; int num_samples; bool interop; float cspeed; - ConfigManager::Mode mode; + Mode mode; + bool split_output = false; - //ao + // ao float ao_radius; int num_ao_rays; bool ao_enabled; - //camera + // camera RadeonRays::float3 camera_pos; RadeonRays::float3 camera_at; RadeonRays::float3 camera_up; @@ -66,48 +69,56 @@ namespace Baikal float camera_focal_length; CameraType camera_type; - //folder to store camera position output + // folder to store camera position output std::string camera_out_folder; - //app + // app bool progressive; bool cmd_line_mode; bool recording_enabled; bool benchmark; bool gui_visible; - //bencmark + // bencmark Estimator::RayTracingStats stats; bool time_benchmarked; bool rt_benchmarked; bool time_benchmark; float time_benchmark_time; - //image file + // image file std::string base_image_file_name; std::string image_file_format; - //unused + std::string light_file; + + // unused int num_shadow_rays; int samplecount; float envmapmul; + // OpenCL platform & device settings int platform_index; int device_index; + // device settings + float gpu_mem_fraction = 0; // float number from 0 to 1, percentage of max used device memory, 0 for default behavior + std::string visible_devices; + + // postprocessing settings + PostProcessingType postprocess_type = PostProcessingType::kNone; + std::uint32_t postprocess_start_spp = 1; + std::uint32_t postprocess_every_frame = 0; // TODO: make it bool }; class AppCliParser { public: - AppCliParser(int argc, char * argv[]); AppSettings Parse(); + static void ShowHelp(); private: - - void ShowHelp(); - CmdParser m_cmd_parser; }; } diff --git a/BaikalStandalone/Application/application.cpp b/BaikalStandalone/Application/application.cpp index c0f7ef91..f75dd095 100644 --- a/BaikalStandalone/Application/application.cpp +++ b/BaikalStandalone/Application/application.cpp @@ -94,6 +94,26 @@ namespace Baikal static float2 g_mouse_delta = float2(0, 0); static float2 g_scroll_delta = float2(0, 0); + static const std::vector> kBaikalOutputs = + { + { Renderer::OutputType::kColor, "Color" }, + { Renderer::OutputType::kOpacity, "Opacity" }, + { Renderer::OutputType::kVisibility, "Visibility" }, + { Renderer::OutputType::kWorldPosition, "World Position" }, + { Renderer::OutputType::kWorldShadingNormal, "World Shading Normal" }, + { Renderer::OutputType::kWorldGeometricNormal, "World Geometric Normal" }, + { Renderer::OutputType::kUv, "Texture Coordinates" }, + { Renderer::OutputType::kWireframe, "Wireframe" }, + { Renderer::OutputType::kAlbedo, "Albedo" }, + { Renderer::OutputType::kWorldTangent, "Tangent" }, + { Renderer::OutputType::kWorldBitangent, "Bitangent" }, + { Renderer::OutputType::kGloss, "Glossiness" }, + { Renderer::OutputType::kMeshID, "Object ID" }, + { Renderer::OutputType::kGroupID, "Object Group ID" }, + { Renderer::OutputType::kBackground, "Background" }, + { Renderer::OutputType::kDepth, "Depth" } + }; + const std::string kCameraLogFile("camera.xml"); //ls - light set const std::string kLightLogFile("light.xml"); @@ -333,7 +353,7 @@ namespace Baikal { ImGuiIO& io = ImGui::GetIO(); Application* app = static_cast(glfwGetWindowUserPointer(window)); - + const bool press_or_repeat = action == GLFW_PRESS || action == GLFW_REPEAT; if (action == GLFW_PRESS) @@ -414,24 +434,23 @@ namespace Baikal camroty = -delta.y; + const float kMovementSpeed = m_settings.cspeed; + if (!g_is_middle_pressed) { - if (std::abs(camroty) > 0.001f) { - camera->Tilt(camroty); + camera->Tilt(camroty * kMovementSpeed); update = true; } if (std::abs(camrotx) > 0.001f) { - - camera->Rotate(camrotx); + camera->Rotate(camrotx * kMovementSpeed); update = true; } } - const float kMovementSpeed = m_settings.cspeed; if (std::abs(scroll_delta.y) > 0.001f) { camera->Zoom(scroll_delta.y * kMovementSpeed); @@ -555,9 +574,9 @@ namespace Baikal { std::swap_ranges(data + channels * w * i, data + channels * w * (i + 1) - 1, data + channels * w * (h - (i + 1))); } - + const auto filename = m_settings.path + "/" + m_settings.base_image_file_name + "-" + std::to_string(time.time_since_epoch().count()) + "." + m_settings.image_file_format; - + auto out = ImageOutput::create(filename); if (out) { @@ -571,7 +590,7 @@ namespace Baikal { std::cout << "Wrong file format\n"; } - + delete[] data; } @@ -612,7 +631,7 @@ namespace Baikal return false; } - Application::Application(int argc, char * argv[]) + Application::Application(int argc, char* argv[]) : m_window(nullptr, glfwDestroyWindow) // Add custom deleter to shared_ptr , m_num_triangles(0) , m_num_instances(0) @@ -621,6 +640,12 @@ namespace Baikal AppCliParser cli(argc, argv); m_settings = cli.Parse(); + if (m_settings.help) + { + AppCliParser::ShowHelp(); + std::exit(0); + } + if (!m_settings.cmd_line_mode) { // Initialize GLFW @@ -709,6 +734,12 @@ namespace Baikal if (!m_settings.cmd_line_mode) { + // Add outputs selectable in UI + for (auto& output : kBaikalOutputs) + { + m_cl->AddOutput(output.first); + } + try { m_cl->StartRenderThreads(); @@ -726,9 +757,9 @@ namespace Baikal m_cl->StopRenderThreads(); } - catch (std::runtime_error&) + catch (std::runtime_error& e) { - std::cout << "Caught exception in Application::Run()\n"; + std::cout << "Caught exception in Application::Run(): " << e.what() << "\n"; throw; } @@ -767,34 +798,13 @@ namespace Baikal } } - - bool Application::UpdateGui() { static const ImVec2 win_size(380, 580); - static float aperture = 0.0f; + static float aperture = m_settings.camera_aperture; static float focal_length = 35.f; static float focus_distance = 1.f; - static int num_bounces = 5; - static const std::vector> kBaikalOutputs = - { - { Renderer::OutputType::kColor, "Color" }, - { Renderer::OutputType::kOpacity, "Opacity" }, - { Renderer::OutputType::kVisibility, "Visibility" }, - { Renderer::OutputType::kWorldPosition, "World Position" }, - { Renderer::OutputType::kWorldShadingNormal, "Shading Normal" }, - { Renderer::OutputType::kWorldGeometricNormal, "Geometric Normal" }, - { Renderer::OutputType::kUv, "Texture Coordinates" }, - { Renderer::OutputType::kWireframe, "Wireframe" }, - { Renderer::OutputType::kAlbedo, "Albedo" }, - { Renderer::OutputType::kWorldTangent, "Tangent" }, - { Renderer::OutputType::kWorldBitangent, "Bitangent" }, - { Renderer::OutputType::kGloss, "Glossiness" }, - { Renderer::OutputType::kMeshID, "Object ID" }, - { Renderer::OutputType::kGroupID, "Object Group ID" }, - { Renderer::OutputType::kBackground, "Background" }, - { Renderer::OutputType::kDepth, "Depth" } - }; + static int num_bounces = m_settings.num_bounces; static int output = 0; bool update = false; @@ -956,27 +966,36 @@ namespace Baikal ImGui::Text("Shadow rays: %f Mrays/s", stats.shadow_throughput * 1e-6f); } -#ifdef ENABLE_DENOISER - ImGui::Separator(); + if (m_cl->GetDenoiserType() == PostProcessingType::kBilateralDenoiser || + m_cl->GetDenoiserType() == PostProcessingType::kWaveletDenoser) + { + ImGui::Separator(); - static float sigmaPosition = m_cl->GetDenoiserFloatParam("position_sensitivity").x; - static float sigmaNormal = m_cl->GetDenoiserFloatParam("normal_sensitivity").x; - static float sigmaColor = m_cl->GetDenoiserFloatParam("color_sensitivity").x; + static float sigmaPosition = m_cl->GetDenoiserFloatParam("position_sensitivity"); + static float sigmaNormal = m_cl->GetDenoiserFloatParam("normal_sensitivity"); + static float sigmaColor = m_cl->GetDenoiserFloatParam("color_sensitivity"); - ImGui::Text("Denoiser settings"); - ImGui::SliderFloat("Position sigma", &sigmaPosition, 0.f, 0.3f); - ImGui::SliderFloat("Normal sigma", &sigmaNormal, 0.f, 5.f); - ImGui::SliderFloat("Color sigma", &sigmaColor, 0.f, 5.f); + ImGui::Text("Denoiser settings"); + ImGui::SliderFloat("Position sigma", &sigmaPosition, 0.f, 0.3f); + ImGui::SliderFloat("Normal sigma", &sigmaNormal, 0.f, 5.f); + ImGui::SliderFloat("Color sigma", &sigmaColor, 0.f, 5.f); - if (m_cl->GetDenoiserFloatParam("position_sensitivity").x != sigmaPosition || - m_cl->GetDenoiserFloatParam("normal_sensitivity").x != sigmaNormal || - m_cl->GetDenoiserFloatParam("color_sensitivity").x != sigmaColor) + if (m_cl->GetDenoiserFloatParam("position_sensitivity") != sigmaPosition || + m_cl->GetDenoiserFloatParam("normal_sensitivity") != sigmaNormal || + m_cl->GetDenoiserFloatParam("color_sensitivity") != sigmaColor) + { + m_cl->SetDenoiserFloatParam("position_sensitivity", sigmaPosition); + m_cl->SetDenoiserFloatParam("normal_sensitivity", sigmaNormal); + m_cl->SetDenoiserFloatParam("color_sensitivity", sigmaColor); + } + } + + if (m_cl->GetDenoiserType() != PostProcessingType::kNone) { - m_cl->SetDenoiserFloatParam("position_sensitivity", sigmaPosition); - m_cl->SetDenoiserFloatParam("normal_sensitivity", sigmaNormal); - m_cl->SetDenoiserFloatParam("color_sensitivity", sigmaColor); + ImGui::Separator(); + ImGui::Checkbox("Split output", &m_settings.split_output); } -#endif + ImGui::End(); // Get shape/material info from renderer diff --git a/BaikalStandalone/Application/cl_render.cpp b/BaikalStandalone/Application/cl_render.cpp index da0c84aa..25a876cf 100644 --- a/BaikalStandalone/Application/cl_render.cpp +++ b/BaikalStandalone/Application/cl_render.cpp @@ -23,40 +23,61 @@ THE SOFTWARE. #include "image_io.h" #include "Application/cl_render.h" +#include "Application/scene_load_utils.h" #include "Application/gl_render.h" +#include "Utils/output_accessor.h" #include "SceneGraph/scene1.h" #include "SceneGraph/camera.h" #include "SceneGraph/material.h" -#include "scene_io.h" -#include "material_io.h" -#include "SceneGraph/material.h" - +#include "BaikalIO/scene_io.h" +#include "BaikalIO/image_io.h" +#include "BaikalIO/material_io.h" #include "Renderers/monte_carlo_renderer.h" #include "Renderers/adaptive_renderer.h" +#include "Utils/clw_class.h" + +#include "OpenImageIO/imageio.h" #include #include #include #include +#include -#ifdef ENABLE_DENOISER -#include "PostEffects/wavelet_denoiser.h" -#endif -#include "Utils/clw_class.h" +#include "PostEffects/ML/ml_post_effect.h" namespace Baikal { - AppClRender::AppClRender(AppSettings& settings, GLuint tex) : m_tex(tex), m_output_type(Renderer::OutputType::kColor) + + namespace + { + constexpr float kGamma = 2.2f; + + // Define the IMAGE_DUMP_PATH value to dump render data +#ifdef IMAGE_DUMP_PATH + std::unique_ptr s_output_accessor; +#endif + } + + AppClRender::AppClRender(AppSettings& settings, GLuint tex) + : m_tex(tex) + , m_post_processing_type(settings.postprocess_type) { InitCl(settings, m_tex); - LoadScene(settings); + InitPostEffect(settings); + InitScene(settings); + +#ifdef IMAGE_DUMP_PATH + s_output_accessor = std::make_unique( + IMAGE_DUMP_PATH, m_width, m_height); +#endif } void AppClRender::InitCl(AppSettings& settings, GLuint tex) { // Create cl context - ConfigManager::CreateConfigs( + CreateConfigs( settings.mode, settings.interop, m_cfgs, @@ -64,8 +85,11 @@ namespace Baikal settings.platform_index, settings.device_index); - m_width = (std::uint32_t)settings.width; - m_height = (std::uint32_t)settings.height; + m_width = (m_post_processing_type == PostProcessingType::kMLUpsample) ? + settings.width / 2 : settings.width; + + m_height = (m_post_processing_type == PostProcessingType::kMLUpsample) ? + settings.height / 2 : settings.height; std::cout << "Running on devices: \n"; @@ -76,14 +100,13 @@ namespace Baikal settings.interop = false; - m_outputs.resize(m_cfgs.size()); - m_ctrl.reset(new ControlData[m_cfgs.size()]); + m_ctrl = std::make_unique(m_cfgs.size()); for (std::size_t i = 0; i < m_cfgs.size(); ++i) { - if (m_cfgs[i].type == ConfigManager::kPrimary) + if (m_cfgs[i].type == DeviceType::kPrimary) { - m_primary = static_cast(i); + m_primary = i; if (m_cfgs[i].caninterop) { @@ -108,99 +131,83 @@ namespace Baikal std::cout << "OpenGL interop mode disabled\n"; } + m_outputs.resize(m_cfgs.size()); + //create renderer for (std::size_t i = 0; i < m_cfgs.size(); ++i) { - m_outputs[i].output = m_cfgs[i].factory->CreateOutput(m_width, m_height); - m_outputs[i].dummy_output = m_cfgs[i].factory->CreateOutput(m_width, m_height); - -#ifdef ENABLE_DENOISER - if (m_cfgs[i].type == ConfigManager::kPrimary) - { - CreateDenoiserOutputs(i, settings.width, settings.height); - SetDenoiserOutputs(i); - m_outputs[i].denoiser = m_cfgs[i].factory->CreatePostEffect(Baikal::RenderFactory::PostEffectType::kWaveletDenoiser); - } -#endif - m_cfgs[i].renderer->SetOutput(Baikal::Renderer::OutputType::kColor, m_outputs[i].output.get()); - - m_outputs[i].fdata.resize(settings.width * settings.height); - m_outputs[i].udata.resize(settings.width * settings.height * 4); + AddRendererOutput(i, Renderer::OutputType::kColor); + m_outputs[i].dummy_output = m_cfgs[i].factory->CreateOutput(m_width, m_height); // TODO: mldenoiser, clear? - if (m_cfgs[i].type == ConfigManager::kPrimary) - { - m_outputs[i].copybuffer = m_cfgs[i].context.CreateBuffer(m_width * m_height, CL_MEM_READ_WRITE); - } + m_outputs[i].fdata.resize(m_width * m_height); + m_outputs[i].udata.resize(4 * m_width * m_height); } m_shape_id_data.output = m_cfgs[m_primary].factory->CreateOutput(m_width, m_height); - m_cfgs[m_primary].renderer->Clear(RadeonRays::float3(0, 0, 0), *m_outputs[m_primary].output); m_cfgs[m_primary].renderer->Clear(RadeonRays::float3(0, 0, 0), *m_shape_id_data.output); + m_copybuffer = m_cfgs[m_primary].context.CreateBuffer(m_width * m_height, CL_MEM_READ_WRITE); } - - void AppClRender::LoadScene(AppSettings& settings) + void AppClRender::AddPostEffect(size_t device_idx, PostEffectType type) { - rand_init(); + m_post_effect = m_cfgs[device_idx].factory->CreatePostEffect(type); - // Load obj file - std::string basepath = settings.path; - basepath += "/"; - std::string filename = basepath + settings.modelname; + // create or get inputs for post-effect + for (auto required_input : m_post_effect->GetInputTypes()) + { + AddRendererOutput(device_idx, required_input); + m_post_effect_inputs[required_input] = GetRendererOutput(device_idx, required_input); + } + // create buffer for post-effect output + if (m_post_processing_type != PostProcessingType::kMLUpsample) { - m_scene = Baikal::SceneIo::LoadScene(filename, basepath); + m_post_effect_output = m_cfgs[device_idx].factory->CreateOutput(m_width, m_height); + } + else + { + m_post_effect_output = m_cfgs[device_idx].factory->CreateOutput(2 * m_width, 2 * m_height); + m_upscaled_img = m_cfgs[device_idx].factory->CreateOutput(2 * m_width, 2 * m_height); + } - { - #ifdef WIN32 - #undef LoadImage - #endif - auto image_io(ImageIo::CreateImageIo()); - auto ibl_texture = image_io->LoadImage(settings.envmapname); - - auto ibl = ImageBasedLight::Create(); - ibl->SetTexture(ibl_texture); - m_scene->AttachLight(ibl); - } + m_shape_id_data.output = m_cfgs[m_primary].factory->CreateOutput(m_width, m_height); + m_cfgs[m_primary].renderer->Clear(RadeonRays::float3(0, 0, 0), *m_shape_id_data.output); + } - // Enable this to generate new materal mapping for a model -#if 0 - auto material_io{Baikal::MaterialIo::CreateMaterialIoXML()}; - material_io->SaveMaterialsFromScene(basepath + "materials.xml", *m_scene); - material_io->SaveIdentityMapping(basepath + "mapping.xml", *m_scene); -#endif + PostProcessingType AppClRender::GetDenoiserType() const + { + return m_post_processing_type; + } - // Check it we have material remapping - std::ifstream in_materials(basepath + "materials.xml"); - std::ifstream in_mapping(basepath + "mapping.xml"); + void AppClRender::SetDenoiserFloatParam(std::string const& name, float value) + { + m_post_effect->SetParameter(name, value); + } - if (in_materials && in_mapping) - { - in_materials.close(); - in_mapping.close(); + float AppClRender::GetDenoiserFloatParam(std::string const& name) const + { + return m_post_effect->GetParameter(name).GetFloat(); + } - auto material_io = Baikal::MaterialIo::CreateMaterialIoXML(); - auto mats = material_io->LoadMaterials(basepath + "materials.xml"); - auto mapping = material_io->LoadMaterialMapping(basepath + "mapping.xml"); + void AppClRender::InitScene(AppSettings& settings) + { + rand_init(); - material_io->ReplaceSceneMaterials(*m_scene, *mats, mapping); - } - } + m_scene = LoadScene(settings); switch (settings.camera_type) { case CameraType::kPerspective: m_camera = Baikal::PerspectiveCamera::Create( - settings.camera_pos - , settings.camera_at - , settings.camera_up); - + settings.camera_pos + , settings.camera_at + , settings.camera_up); break; case CameraType::kOrthographic: m_camera = Baikal::OrthographicCamera::Create( - settings.camera_pos - , settings.camera_at - , settings.camera_up); + settings.camera_pos + , settings.camera_at + , settings.camera_up); break; default: throw std::runtime_error("AppClRender::InitCl(...): unsupported camera type"); @@ -236,30 +243,41 @@ namespace Baikal { for (std::size_t i = 0; i < m_cfgs.size(); ++i) { - if (i == static_cast(m_primary)) + if (i == m_primary) { - m_cfgs[i].renderer->Clear(float3(0, 0, 0), *m_outputs[i].output); + m_cfgs[i].renderer->Clear(float3(), *GetRendererOutput(i, Renderer::OutputType::kColor)); m_cfgs[i].controller->CompileScene(m_scene); ++m_ctrl[i].scene_state; -#ifdef ENABLE_DENOISER - ClearDenoiserOutputs(i); -#endif - + if (m_post_effect_output) + { + m_post_effect_output->Clear(float3()); + } } else m_ctrl[i].clear.store(true); } + + for (auto& output : m_outputs[m_primary].render_outputs) + { + output.second->Clear(float3()); + } } void AppClRender::Update(AppSettings& settings) { ++settings.samplecount; +#ifdef IMAGE_DUMP_PATH + s_output_accessor->SaveAllOutputs(m_outputs); +#endif + for (std::size_t i = 0; i < m_cfgs.size(); ++i) { - if (m_cfgs[i].type == ConfigManager::kPrimary) + if (m_cfgs[i].type == DeviceType::kPrimary) + { continue; + } int desired = 1; if (std::atomic_compare_exchange_strong(&m_ctrl[i].newdata, &desired, 0)) @@ -270,18 +288,21 @@ namespace Baikal continue; } - { - m_cfgs[m_primary].context.WriteBuffer(0, m_outputs[m_primary].copybuffer, &m_outputs[i].fdata[0], settings.width * settings.height); - } + m_cfgs[m_primary].context.WriteBuffer( + 0, m_copybuffer, + &m_outputs[i].fdata[0], + m_width * m_height); - auto acckernel = static_cast(m_cfgs[m_primary].renderer.get())->GetAccumulateKernel(); + auto acckernel = static_cast(m_cfgs[m_primary].renderer.get())->GetAccumulateKernel(); int argc = 0; - acckernel.SetArg(argc++, m_outputs[m_primary].copybuffer); - acckernel.SetArg(argc++, settings.width * settings.width); - acckernel.SetArg(argc++, static_cast(m_outputs[m_primary].output.get())->data()); + acckernel.SetArg(argc++, m_copybuffer); + acckernel.SetArg(argc++, m_width * m_height); + acckernel.SetArg(argc++, + static_cast(GetRendererOutput( + m_primary, Renderer::OutputType::kColor))->data()); - int globalsize = settings.width * settings.height; + int globalsize = m_width * m_height; m_cfgs[m_primary].context.Launch1D(0, ((globalsize + 63) / 64) * 64, 64, acckernel); settings.samplecount += m_ctrl[i].new_samples_count; } @@ -289,60 +310,67 @@ namespace Baikal if (!settings.interop) { -#ifdef ENABLE_DENOISER - m_outputs[m_primary].output_denoised->GetData(&m_outputs[m_primary].fdata[0]); -#else - m_outputs[m_primary].output->GetData(&m_outputs[m_primary].fdata[0]); -#endif - - float gamma = 2.2f; - for (int i = 0; i < (int)m_outputs[m_primary].fdata.size(); ++i) + if (m_post_effect_output) + { + m_post_effect_output->GetData(&m_outputs[m_primary].fdata[0]); + ApplyGammaCorrection(m_primary, kGamma, false); + } + else { - m_outputs[m_primary].udata[4 * i] = (unsigned char)clamp(clamp(pow(m_outputs[m_primary].fdata[i].x / m_outputs[m_primary].fdata[i].w, 1.f / gamma), 0.f, 1.f) * 255, 0, 255); - m_outputs[m_primary].udata[4 * i + 1] = (unsigned char)clamp(clamp(pow(m_outputs[m_primary].fdata[i].y / m_outputs[m_primary].fdata[i].w, 1.f / gamma), 0.f, 1.f) * 255, 0, 255); - m_outputs[m_primary].udata[4 * i + 2] = (unsigned char)clamp(clamp(pow(m_outputs[m_primary].fdata[i].z / m_outputs[m_primary].fdata[i].w, 1.f / gamma), 0.f, 1.f) * 255, 0, 255); - m_outputs[m_primary].udata[4 * i + 3] = 1; + GetOutputData(m_primary, Renderer::OutputType::kColor, &m_outputs[m_primary].fdata[0]); + ApplyGammaCorrection(m_primary, kGamma, true); } glActiveTexture(GL_TEXTURE0); glBindTexture(GL_TEXTURE_2D, m_tex); - glTexSubImage2D(GL_TEXTURE_2D, 0, 0, 0, m_outputs[m_primary].output->width(), m_outputs[m_primary].output->height(), GL_RGBA, GL_UNSIGNED_BYTE, &m_outputs[m_primary].udata[0]); + glTexSubImage2D(GL_TEXTURE_2D, 0, 0, 0, m_width, m_height, GL_RGBA, GL_UNSIGNED_BYTE, &m_outputs[m_primary].udata[0]); glBindTexture(GL_TEXTURE_2D, 0); } else { - std::vector objects; - objects.push_back(m_cl_interop_image); - m_cfgs[m_primary].context.AcquireGLObjects(0, objects); - - auto copykernel = static_cast(m_cfgs[m_primary].renderer.get())->GetCopyKernel(); - -#ifdef ENABLE_DENOISER - auto output = m_outputs[m_primary].output_denoised.get(); -#else - auto output = m_outputs[m_primary].output.get(); -#endif - - int argc = 0; - - copykernel.SetArg(argc++, static_cast(output)->data()); - copykernel.SetArg(argc++, output->width()); - copykernel.SetArg(argc++, output->height()); - copykernel.SetArg(argc++, 2.2f); - copykernel.SetArg(argc++, m_cl_interop_image); - - int globalsize = output->width() * output->height(); - m_cfgs[m_primary].context.Launch1D(0, ((globalsize + 63) / 64) * 64, 64, copykernel); - - m_cfgs[m_primary].context.ReleaseGLObjects(0, objects); - m_cfgs[m_primary].context.Finish(0); + if (!m_post_effect_output) + { + CopyToGL(GetRendererOutput(m_primary, m_output_type)); + } + else + { + if (m_output_type == Renderer::OutputType::kColor) + { + if (settings.split_output) + { + if (m_post_processing_type != PostProcessingType::kMLUpsample) + { + CopyToGL(GetRendererOutput(m_primary, Renderer::OutputType::kColor), + m_post_effect_output.get()); + } + else + { + auto ml_post_effect = dynamic_cast(m_post_effect.get()); + + ml_post_effect->Resize_2x(dynamic_cast(m_upscaled_img.get())->data(), + dynamic_cast(GetRendererOutput( + m_primary, + Renderer::OutputType::kColor))->data()); + + CopyToGL(m_upscaled_img.get(), m_post_effect_output.get()); + } + } + else + { + CopyToGL(m_post_effect_output.get()); + } + } + else + { + CopyToGL(GetRendererOutput(m_primary, m_output_type)); + } + } } - if (settings.benchmark) { auto& scene = m_cfgs[m_primary].controller->CompileScene(m_scene); - static_cast(m_cfgs[m_primary].renderer.get())->Benchmark(scene, settings.stats); + dynamic_cast(m_cfgs[m_primary].renderer.get())->Benchmark(scene, settings.stats); settings.benchmark = false; settings.rt_benchmarked = true; @@ -352,14 +380,7 @@ namespace Baikal void AppClRender::Render(int sample_cnt) { -#ifdef ENABLE_DENOISER - WaveletDenoiser* wavelet_denoiser = dynamic_cast(m_outputs[m_primary].denoiser.get()); - - if (wavelet_denoiser != nullptr) - { - wavelet_denoiser->Update(static_cast(m_camera.get())); - } -#endif + SetPostEffectParams(sample_cnt); auto& scene = m_cfgs[m_primary].controller->GetCachedScene(m_scene); m_cfgs[m_primary].renderer->Render(scene); @@ -377,34 +398,10 @@ namespace Baikal m_shape_id_requested = false; } -#ifdef ENABLE_DENOISER - Baikal::PostEffect::InputSet input_set; - input_set[Baikal::Renderer::OutputType::kColor] = m_outputs[m_primary].output.get(); - input_set[Baikal::Renderer::OutputType::kWorldShadingNormal] = m_outputs[m_primary].output_normal.get(); - input_set[Baikal::Renderer::OutputType::kWorldPosition] = m_outputs[m_primary].output_position.get(); - input_set[Baikal::Renderer::OutputType::kAlbedo] = m_outputs[m_primary].output_albedo.get(); - input_set[Baikal::Renderer::OutputType::kMeshID] = m_outputs[m_primary].output_mesh_id.get(); - - auto radius = 10U - RadeonRays::clamp((sample_cnt / 16), 1U, 9U); - auto position_sensitivity = 5.f + 10.f * (radius / 10.f); - - const bool is_bilateral_denoiser = dynamic_cast(m_outputs[m_primary].denoiser.get()) != nullptr; - - if (is_bilateral_denoiser) + if (m_post_effect && m_output_type == Renderer::OutputType::kColor) { - - auto normal_sensitivity = 0.1f + (radius / 10.f) * 0.15f; - auto color_sensitivity = (radius / 10.f) * 2.f; - auto albedo_sensitivity = 0.5f + (radius / 10.f) * 0.5f; - m_outputs[m_primary].denoiser->SetParameter("radius", static_cast(radius)); - m_outputs[m_primary].denoiser->SetParameter("color_sensitivity", color_sensitivity); - m_outputs[m_primary].denoiser->SetParameter("normal_sensitivity", normal_sensitivity); - m_outputs[m_primary].denoiser->SetParameter("position_sensitivity", position_sensitivity); - m_outputs[m_primary].denoiser->SetParameter("albedo_sensitivity", albedo_sensitivity); + m_post_effect->Apply(m_post_effect_inputs, *m_post_effect_output); } - - m_outputs[m_primary].denoiser->Apply(input_set, *m_outputs[m_primary].output_denoised); -#endif } void AppClRender::SaveFrameBuffer(AppSettings& settings) @@ -415,10 +412,14 @@ namespace Baikal std::vector output_data; if (settings.interop) { - auto output = m_outputs[m_primary].output.get(); - auto buffer = static_cast(output)->data(); + auto output = GetRendererOutput(m_primary, Renderer::OutputType::kColor); + auto buffer = dynamic_cast(output)->data(); output_data.resize(buffer.GetElementCount()); - m_cfgs[m_primary].context.ReadBuffer(0, static_cast(output)->data(), &output_data[0], output_data.size()).Wait(); + m_cfgs[m_primary].context.ReadBuffer( + 0, + dynamic_cast(output)->data(), + &output_data[0], + output_data.size()).Wait(); } //use already copied to CPU cl data in case of no interop @@ -456,9 +457,9 @@ namespace Baikal float3 val = data[(height - 1 - y) * width + x]; tempbuf[y * width + x] = (1.f / val.w) * val; - tempbuf[y * width + x].x = std::pow(tempbuf[y * width + x].x, 1.f / 2.2f); - tempbuf[y * width + x].y = std::pow(tempbuf[y * width + x].y, 1.f / 2.2f); - tempbuf[y * width + x].z = std::pow(tempbuf[y * width + x].z, 1.f / 2.2f); + tempbuf[y * width + x].x = std::pow(tempbuf[y * width + x].x, 1.f / kGamma); + tempbuf[y * width + x].y = std::pow(tempbuf[y * width + x].y, 1.f / kGamma); + tempbuf[y * width + x].z = std::pow(tempbuf[y * width + x].z, 1.f / kGamma); } ImageOutput* out = ImageOutput::create(name); @@ -479,7 +480,6 @@ namespace Baikal { auto renderer = m_cfgs[cd.idx].renderer.get(); auto controller = m_cfgs[cd.idx].controller.get(); - auto output = m_outputs[cd.idx].output.get(); auto updatetime = std::chrono::high_resolution_clock::now(); @@ -492,7 +492,10 @@ namespace Baikal bool update = false; if (std::atomic_compare_exchange_strong(&cd.clear, &result, 0)) { - renderer->Clear(float3(0, 0, 0), *output); + for (auto& output : m_outputs[cd.idx].render_outputs) + { + output.second->Clear(float3()); + } controller->CompileScene(m_scene); scene_state = m_ctrl[m_primary].scene_state; update = true; @@ -508,7 +511,7 @@ namespace Baikal if (update) { - m_outputs[cd.idx].output->GetData(&m_outputs[cd.idx].fdata[0]); + GetOutputData(cd.idx, Renderer::OutputType::kColor, &m_outputs[cd.idx].fdata[0]); updatetime = now; m_ctrl[cd.idx].scene_state = scene_state; m_ctrl[cd.idx].new_samples_count = new_samples_count; @@ -524,32 +527,29 @@ namespace Baikal { for (std::size_t i = 0; i < m_cfgs.size(); ++i) { - if (i != static_cast(m_primary)) + if (i != m_primary) { - m_renderthreads.push_back(std::thread(&AppClRender::RenderThread, this, std::ref(m_ctrl[i]))); + m_renderthreads.emplace_back(&AppClRender::RenderThread, this, std::ref(m_ctrl[i])); } } - std::cout << m_cfgs.size() << " OpenCL submission threads started\n"; + std::cout << m_renderthreads.size() << " OpenCL submission threads started\n"; } void AppClRender::StopRenderThreads() { for (std::size_t i = 0; i < m_cfgs.size(); ++i) { - if (i == static_cast(m_primary)) + if (i != m_primary) { - continue; + m_ctrl[i].stop.store(true); } - - m_ctrl[i].stop.store(true); } for (std::size_t i = 0; i < m_renderthreads.size(); ++i) { m_renderthreads[i].join(); } - } void AppClRender::RunBenchmark(AppSettings& settings) @@ -569,15 +569,8 @@ namespace Baikal settings.time_benchmark_time = delta / 1000.f; - m_outputs[m_primary].output->GetData(&m_outputs[m_primary].fdata[0]); - float gamma = 2.2f; - for (int i = 0; i < (int)m_outputs[m_primary].fdata.size(); ++i) - { - m_outputs[m_primary].udata[4 * i] = (unsigned char)clamp(clamp(pow(m_outputs[m_primary].fdata[i].x / m_outputs[m_primary].fdata[i].w, 1.f / gamma), 0.f, 1.f) * 255, 0, 255); - m_outputs[m_primary].udata[4 * i + 1] = (unsigned char)clamp(clamp(pow(m_outputs[m_primary].fdata[i].y / m_outputs[m_primary].fdata[i].w, 1.f / gamma), 0.f, 1.f) * 255, 0, 255); - m_outputs[m_primary].udata[4 * i + 2] = (unsigned char)clamp(clamp(pow(m_outputs[m_primary].fdata[i].z / m_outputs[m_primary].fdata[i].w, 1.f / gamma), 0.f, 1.f) * 255, 0, 255); - m_outputs[m_primary].udata[4 * i + 3] = 1; - } + GetOutputData(m_primary, Renderer::OutputType::kColor, &m_outputs[m_primary].fdata[0]); + ApplyGammaCorrection(m_primary, kGamma, true); auto& fdata = m_outputs[m_primary].fdata; std::vector data(fdata.size()); @@ -596,39 +589,56 @@ namespace Baikal std::cout << "Running RT benchmark...\n"; auto& scene = m_cfgs[m_primary].controller->GetCachedScene(m_scene); - static_cast(m_cfgs[m_primary].renderer.get())->Benchmark(scene, settings.stats); + dynamic_cast(m_cfgs[m_primary].renderer.get())->Benchmark(scene, settings.stats); } - void AppClRender::SetNumBounces(int num_bounces) + void AppClRender::ApplyGammaCorrection(size_t device_idx, float gamma, bool divide_by_spp = true) { - for (std::size_t i = 0; i < m_cfgs.size(); ++i) + auto adjust_gamma = [gamma](float val) { - static_cast(m_cfgs[i].renderer.get())->SetMaxBounces(num_bounces); - } - } + return (unsigned char)clamp(std::pow(val , 1.f / gamma) * 255, 0, 255); + }; - void AppClRender::SetOutputType(Renderer::OutputType type) - { - for (std::size_t i = 0; i < m_cfgs.size(); ++i) + if (divide_by_spp) { -#ifdef ENABLE_DENOISER - RestoreDenoiserOutput(i, m_output_type); -#else - m_cfgs[i].renderer->SetOutput(m_output_type, nullptr); -#endif - if (type == Renderer::OutputType::kOpacity || type == Renderer::OutputType::kVisibility) + for (int i = 0; i < (int)m_outputs[device_idx].fdata.size(); ++i) { - m_cfgs[i].renderer->SetOutput(Renderer::OutputType::kColor, m_outputs[i].dummy_output.get()); + auto &fdata = m_outputs[device_idx].fdata[i]; + auto &udata = m_outputs[device_idx].udata; + + udata[4 * i] = adjust_gamma(fdata.x / fdata.w); + udata[4 * i + 1] = adjust_gamma(fdata.y / fdata.w); + udata[4 * i + 2] = adjust_gamma(fdata.z / fdata.w); + udata[4 * i + 3] = 1; } - else + } + else + { + for (int i = 0; i < (int)m_outputs[device_idx].fdata.size(); ++i) { - m_cfgs[i].renderer->SetOutput(Renderer::OutputType::kColor, nullptr); + auto &fdata = m_outputs[device_idx].fdata[i]; + auto &udata = m_outputs[device_idx].udata; + + udata[4 * i] = adjust_gamma(fdata.x); + udata[4 * i + 1] = adjust_gamma(fdata.y); + udata[4 * i + 2] = adjust_gamma(fdata.z); + udata[4 * i + 3] = 1; } - m_cfgs[i].renderer->SetOutput(type, m_outputs[i].output.get()); } - m_output_type = type; } + void AppClRender::SetNumBounces(int num_bounces) + { + for (std::size_t i = 0; i < m_cfgs.size(); ++i) + { + dynamic_cast(m_cfgs[i].renderer.get())->SetMaxBounces(num_bounces); + } + } + + void AppClRender::SetOutputType(Renderer::OutputType type) + { + m_output_type = type; + } std::future AppClRender::GetShapeId(std::uint32_t x, std::uint32_t y) { @@ -664,63 +674,152 @@ namespace Baikal return nullptr; } -#ifdef ENABLE_DENOISER - void AppClRender::SetDenoiserFloatParam(const std::string& name, const float4& value) + void AppClRender::AddRendererOutput(size_t device_idx, Renderer::OutputType type) { - m_outputs[m_primary].denoiser->SetParameter(name, value); + auto it = m_outputs[device_idx].render_outputs.find(type); + if (it == m_outputs[device_idx].render_outputs.end()) + { + const auto& config = m_cfgs.at(device_idx); + auto output = config.factory->CreateOutput(m_width, m_height); + config.renderer->SetOutput(type, output.get()); + config.renderer->Clear(RadeonRays::float3(0, 0, 0), *output); + + m_outputs[device_idx].render_outputs.emplace(type, std::move(output)); + } } - float4 AppClRender::GetDenoiserFloatParam(const std::string& name) + Output* AppClRender::GetRendererOutput(size_t device_idx, Renderer::OutputType type) { - return m_outputs[m_primary].denoiser->GetParameter(name); + return m_outputs.at(device_idx).render_outputs.at(type).get(); } - void AppClRender::CreateDenoiserOutputs(std::size_t cfg_index, int width, int height) + void AppClRender::GetOutputData(size_t device_idx, Renderer::OutputType type, RadeonRays::float3* data) const { - m_outputs[cfg_index].output_denoised = m_cfgs[cfg_index].factory->CreateOutput(width, height); - m_outputs[cfg_index].output_normal = m_cfgs[cfg_index].factory->CreateOutput(width, height); - m_outputs[cfg_index].output_position = m_cfgs[cfg_index].factory->CreateOutput(width, height); - m_outputs[cfg_index].output_albedo = m_cfgs[cfg_index].factory->CreateOutput(width, height); - m_outputs[cfg_index].output_mesh_id = m_cfgs[cfg_index].factory->CreateOutput(width, height); + m_outputs.at(device_idx).render_outputs.at(type)->GetData(data); } - void AppClRender::SetDenoiserOutputs(std::size_t cfg_index) const + void AppClRender::CopyToGL(Output* output) { - m_cfgs[cfg_index].renderer->SetOutput(Renderer::OutputType::kWorldShadingNormal, m_outputs[cfg_index].output_normal.get()); - m_cfgs[cfg_index].renderer->SetOutput(Renderer::OutputType::kWorldPosition, m_outputs[cfg_index].output_position.get()); - m_cfgs[cfg_index].renderer->SetOutput(Renderer::OutputType::kAlbedo, m_outputs[cfg_index].output_albedo.get()); - m_cfgs[cfg_index].renderer->SetOutput(Renderer::OutputType::kMeshID, m_outputs[cfg_index].output_mesh_id.get()); + std::vector objects = {m_cl_interop_image}; + m_cfgs[m_primary].context.AcquireGLObjects(0, objects); + + auto copykernel = dynamic_cast( + m_cfgs[m_primary].renderer.get())->GetCopyKernel(); + + copykernel.SetArg(0, dynamic_cast(output)->data()); + copykernel.SetArg(1, output->width()); + copykernel.SetArg(2, output->height()); + copykernel.SetArg(3, kGamma); + copykernel.SetArg(4, m_cl_interop_image); + + int globalsize = output->width() * output->height(); + m_cfgs[m_primary].context.Launch1D(0, (globalsize + 63) / 64 * 64, 64, copykernel); + + m_cfgs[m_primary].context.ReleaseGLObjects(0, objects); + m_cfgs[m_primary].context.Finish(0); } - void AppClRender::ClearDenoiserOutputs(std::size_t cfg_index) const + void AppClRender::CopyToGL(Output* left_output, Output* right_output) { - m_cfgs[cfg_index].renderer->Clear(float3(0, 0, 0), *m_outputs[cfg_index].output_normal); - m_cfgs[cfg_index].renderer->Clear(float3(0, 0, 0), *m_outputs[cfg_index].output_position); - m_cfgs[cfg_index].renderer->Clear(float3(0, 0, 0), *m_outputs[cfg_index].output_albedo); - m_cfgs[cfg_index].renderer->Clear(float3(0, 0, 0), *m_outputs[cfg_index].output_mesh_id); + assert(left_output->width() == right_output->width()); + assert(left_output->height() == right_output->height()); + + std::vector objects = {m_cl_interop_image}; + m_cfgs[m_primary].context.AcquireGLObjects(0, objects); + + auto copykernel = dynamic_cast( + m_cfgs[m_primary].renderer.get())->GetCopySplitKernel(); + + copykernel.SetArg(0, dynamic_cast(left_output)->data()); + copykernel.SetArg(1, dynamic_cast(right_output)->data()); + copykernel.SetArg(2, left_output->width()); + copykernel.SetArg(3, left_output->height()); + copykernel.SetArg(4, kGamma); + copykernel.SetArg(5, m_cl_interop_image); + + int globalsize = left_output->width() * left_output->height(); + m_cfgs[m_primary].context.Launch1D(0, (globalsize + 63) / 64 * 64, 64, copykernel); + + m_cfgs[m_primary].context.ReleaseGLObjects(0, objects); + m_cfgs[m_primary].context.Finish(0); } - void AppClRender::RestoreDenoiserOutput(std::size_t cfg_index, Renderer::OutputType type) const + void AppClRender::AddOutput(Renderer::OutputType type) { - switch (type) + for (std::size_t idx = 0; idx < m_cfgs.size(); ++idx) { - case Renderer::OutputType::kWorldShadingNormal: - m_cfgs[cfg_index].renderer->SetOutput(type, m_outputs[cfg_index].output_normal.get()); + AddRendererOutput(idx, type); + } + } + + void AppClRender::InitPostEffect(AppSettings const& settings) + { + switch (m_post_processing_type) + { + case PostProcessingType::kNone: + break; + case PostProcessingType::kBilateralDenoiser: + if (settings.camera_type != CameraType::kPerspective) + { + throw std::logic_error("Bilateral denoiser requires perspective camera"); + } + AddPostEffect(m_primary, PostEffectType::kBilateralDenoiser); + break; + case PostProcessingType::kWaveletDenoser: + AddPostEffect(m_primary, PostEffectType::kWaveletDenoiser); + break; + case PostProcessingType::kMLDenoiser: + AddPostEffect(m_primary, PostEffectType::kMLDenoiser); + m_post_effect->SetParameter("gpu_memory_fraction", settings.gpu_mem_fraction); + m_post_effect->SetParameter("visible_devices", settings.visible_devices); + m_post_effect->SetParameter("start_spp", settings.postprocess_start_spp); + m_post_effect->SetParameter("every_frame", settings.postprocess_every_frame); break; - case Renderer::OutputType::kWorldPosition: - m_cfgs[cfg_index].renderer->SetOutput(type, m_outputs[cfg_index].output_position.get()); + case PostProcessingType::kMLUpsample: + AddPostEffect(m_primary, PostEffectType::kMLUpsampler); + m_post_effect->SetParameter("gpu_memory_fraction", settings.gpu_mem_fraction); + m_post_effect->SetParameter("visible_devices", settings.visible_devices); + m_post_effect->SetParameter("start_spp", settings.postprocess_start_spp); + m_post_effect->SetParameter("every_frame", settings.postprocess_every_frame); break; - case Renderer::OutputType::kAlbedo: - m_cfgs[cfg_index].renderer->SetOutput(type, m_outputs[cfg_index].output_albedo.get()); + default: + throw std::runtime_error("AppClRender(...): Unsupported denoiser type"); + } + } + + void AppClRender::SetPostEffectParams(int sample_cnt) + { + switch (m_post_processing_type) + { + case PostProcessingType::kBilateralDenoiser: + { + const auto radius = 10U - RadeonRays::clamp((sample_cnt / 16), 1U, 9U); + m_post_effect->SetParameter("radius", static_cast(radius)); + m_post_effect->SetParameter("color_sensitivity", (radius / 10.f) * 2.f); + m_post_effect->SetParameter("normal_sensitivity", 0.1f + (radius / 10.f) * 0.15f); + m_post_effect->SetParameter("position_sensitivity", 5.f + 10.f * (radius / 10.f)); + m_post_effect->SetParameter("albedo_sensitivity", 0.5f + (radius / 10.f) * 0.5f); break; - case Renderer::OutputType::kMeshID: - m_cfgs[cfg_index].renderer->SetOutput(type, m_outputs[cfg_index].output_mesh_id.get()); + } + case PostProcessingType::kWaveletDenoser: + { + auto pCamera = dynamic_cast(m_camera.get()); + m_post_effect->SetParameter("camera_focal_length", pCamera->GetFocalLength()); + m_post_effect->SetParameter("camera_sensor_size", pCamera->GetSensorSize()); + m_post_effect->SetParameter("camera_depth_range", pCamera->GetDepthRange()); + m_post_effect->SetParameter("camera_up_vector", pCamera->GetUpVector()); + m_post_effect->SetParameter("camera_forward_vector", pCamera->GetForwardVector()); + m_post_effect->SetParameter("camera_right_vector", -pCamera->GetRightVector()); + m_post_effect->SetParameter("camera_position", pCamera->GetPosition()); + m_post_effect->SetParameter("camera_aspect_ratio", pCamera->GetAspectRatio()); break; + } + case PostProcessingType::kMLDenoiser: + case PostProcessingType::kMLUpsample: + case PostProcessingType::kNone: default: - // Nothing to restore - m_cfgs[cfg_index].renderer->SetOutput(type, nullptr); break; } } -#endif + } // Baikal diff --git a/BaikalStandalone/Application/cl_render.h b/BaikalStandalone/Application/cl_render.h index 0082e2c6..5c12958e 100644 --- a/BaikalStandalone/Application/cl_render.h +++ b/BaikalStandalone/Application/cl_render.h @@ -22,11 +22,6 @@ ********************************************************************/ #pragma once -#include -#include -#include -#include - #include "RenderFactory/render_factory.h" #include "Renderers/monte_carlo_renderer.h" #include "Output/clwoutput.h" @@ -34,10 +29,16 @@ #include "Utils/config_manager.h" #include "Application/gl_render.h" #include "SceneGraph/camera.h" +#include "PostEffects/post_effect.h" +#include "Output/output.h" -#ifdef ENABLE_DENOISER -#include "PostEffects/bilateral_denoiser.h" -#endif +#include +#include +#include +#include +#include +#include +#include namespace Baikal @@ -46,21 +47,11 @@ namespace Baikal { struct OutputData { + std::map> render_outputs; std::unique_ptr output; std::unique_ptr dummy_output; - -#ifdef ENABLE_DENOISER - std::unique_ptr output_position; - std::unique_ptr output_normal; - std::unique_ptr output_albedo; - std::unique_ptr output_mesh_id; - std::unique_ptr output_denoised; - std::unique_ptr denoiser; -#endif - std::vector fdata; std::vector udata; - CLWBuffer copybuffer; }; struct ControlData @@ -94,7 +85,7 @@ namespace Baikal inline Baikal::Camera::Ptr GetCamera() { return m_camera; }; inline Baikal::Scene1::Ptr GetScene() { return m_scene; }; - inline const std::vector& GetConfigs() const { return m_cfgs; }; + inline const std::vector& GetConfigs() const { return m_cfgs; }; inline Renderer::OutputType GetOutputType() { return m_output_type; }; void SetNumBounces(int num_bounces); @@ -103,20 +94,29 @@ namespace Baikal std::future GetShapeId(std::uint32_t x, std::uint32_t y); Baikal::Shape::Ptr GetShapeById(int shape_id); -#ifdef ENABLE_DENOISER - // Denoiser - void SetDenoiserFloatParam(const std::string& name, const float4& value); - float4 GetDenoiserFloatParam(const std::string& name); - void CreateDenoiserOutputs(std::size_t cfg_index, int width, int height); - void SetDenoiserOutputs(std::size_t cfg_index) const; - void ClearDenoiserOutputs(std::size_t cfg_index) const; - void RestoreDenoiserOutput(std::size_t cfg_index, Renderer::OutputType type) const; -#endif + PostProcessingType GetDenoiserType() const; + void SetDenoiserFloatParam(std::string const& name, float value); + float GetDenoiserFloatParam(std::string const& name) const; + + void AddOutput(Renderer::OutputType type); + private: + using RendererOutputs = std::map>; + void InitCl(AppSettings& settings, GLuint tex); - void LoadScene(AppSettings& settings); + void InitScene(AppSettings& settings); void RenderThread(ControlData& cd); + Output* GetRendererOutput(size_t device_idx, Renderer::OutputType type); + void AddRendererOutput(size_t device_idx, Renderer::OutputType type); + void GetOutputData(size_t device_idx, Renderer::OutputType type, RadeonRays::float3* data) const; + void AddPostEffect(size_t device_idx, PostEffectType type); + void CopyToGL(Output* output); + void CopyToGL(Output* left_output, Output* right_output); + void ApplyGammaCorrection(size_t device_idx, float gamma, bool divide_by_spp); + void InitPostEffect(AppSettings const& settings); + void SetPostEffectParams(int sample_cnt); + Baikal::Scene1::Ptr m_scene; Baikal::Camera::Ptr m_camera; @@ -124,17 +124,25 @@ namespace Baikal bool m_shape_id_requested = false; OutputData m_shape_id_data; RadeonRays::float2 m_shape_id_pos; - std::vector m_cfgs; + std::vector m_cfgs; + std::vector m_outputs; std::unique_ptr m_ctrl; std::vector m_renderthreads; - int m_primary = -1; + size_t m_primary; std::uint32_t m_width, m_height; + CLWBuffer m_copybuffer; //if interop CLWImage2D m_cl_interop_image; //save GL tex for no interop case GLuint m_tex; - Renderer::OutputType m_output_type; + Renderer::OutputType m_output_type = Renderer::OutputType::kColor; + + PostProcessingType m_post_processing_type = PostProcessingType::kNone; + std::unique_ptr m_post_effect; + PostEffect::InputSet m_post_effect_inputs; + std::unique_ptr m_post_effect_output; + std::unique_ptr m_upscaled_img; }; } diff --git a/BaikalStandalone/Application/gl_render.h b/BaikalStandalone/Application/gl_render.h index 10e5e6d9..906b8421 100644 --- a/BaikalStandalone/Application/gl_render.h +++ b/BaikalStandalone/Application/gl_render.h @@ -31,7 +31,6 @@ #include "GLFW/glfw3.h" #elif WIN32 #define NOMINMAX -#include #include "GL/glew.h" #include "GLFW/glfw3.h" #else diff --git a/BaikalStandalone/Application/scene_load_utils.cpp b/BaikalStandalone/Application/scene_load_utils.cpp new file mode 100644 index 00000000..8dea9766 --- /dev/null +++ b/BaikalStandalone/Application/scene_load_utils.cpp @@ -0,0 +1,148 @@ +/********************************************************************** + Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved. + + Permission is hereby granted, free of charge, to any person obtaining a copy + of this software and associated documentation files (the "Software"), to deal + in the Software without restriction, including without limitation the rights + to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + copies of the Software, and to permit persons to whom the Software is + furnished to do so, subject to the following conditions: + + The above copyright notice and this permission notice shall be included in + all copies or substantial portions of the Software. + + THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + THE SOFTWARE. + ********************************************************************/ + +#include "scene_load_utils.h" +#include +#include +#include +#include +#include "XML/tinyxml2.h" + +using namespace Baikal; + +namespace +{ + void LoadMaterials(std::string const& basepath, Scene1::Ptr scene) + { + // Check it we have material remapping + std::ifstream in_materials(basepath + "materials.xml"); + std::ifstream in_mapping(basepath + "mapping.xml"); + + if (!in_materials || !in_mapping) + return; + + auto material_io = Baikal::MaterialIo::CreateMaterialIoXML(); + auto mats = material_io->LoadMaterials(basepath + "materials.xml"); + auto mapping = material_io->LoadMaterialMapping(basepath + "mapping.xml"); + material_io->ReplaceSceneMaterials(*scene, *mats, mapping); + } + + void LoadLights(std::string const& light_file, Scene1::Ptr scene) + { + if (light_file.empty()) + { + return; + } + + tinyxml2::XMLDocument doc; + doc.LoadFile(light_file.c_str()); + + auto root = doc.FirstChildElement("light_list"); + + if (!root) + { + throw std::runtime_error("Failed to open lights set file."); + } + + auto elem = root->FirstChildElement("light"); + + while (elem) + { + Light::Ptr light; + std::string type = elem->Attribute("type"); + + if (type == "point") + { + light = PointLight::Create(); + } + else if (type == "spot") + { + auto spot = SpotLight::Create(); + RadeonRays::float2 cone_shape(elem->FloatAttribute("csx"), + elem->FloatAttribute("csy")); + spot->SetConeShape(cone_shape); + light = spot; + } + else if (type == "direct") + { + light = DirectionalLight::Create(); + } + else if (type == "ibl") + { + auto image_io = ImageIo::CreateImageIo(); + auto tex = image_io->LoadImage(elem->Attribute("tex")); + auto ibl = ImageBasedLight::Create(); + ibl->SetTexture(tex); + ibl->SetMultiplier(elem->FloatAttribute("mul")); + light = ibl; + } + else + { + throw std::runtime_error(std::string("Unknown light type: ") + type); + } + + RadeonRays::float3 rad; + + light->SetPosition({ + elem->FloatAttribute("posx"), + elem->FloatAttribute("posy"), + elem->FloatAttribute("posz")}); + + light->SetDirection({ + elem->FloatAttribute("dirx"), + elem->FloatAttribute("diry"), + elem->FloatAttribute("dirz")}); + + light->SetEmittedRadiance({ + elem->FloatAttribute("radx"), + elem->FloatAttribute("rady"), + elem->FloatAttribute("radz")}); + + scene->AttachLight(light); + + elem = elem->NextSiblingElement("light"); + } + } +} + + +Scene1::Ptr LoadScene(Baikal::AppSettings const& settings) +{ +#ifdef WIN32 + std::string basepath = settings.path + "\\"; +#else + std::string basepath = settings.path + "/"; +#endif + + std::string filename = basepath + settings.modelname; + + auto scene = Baikal::SceneIo::LoadScene(filename, basepath); + + if (scene == nullptr) + { + throw std::runtime_error("LoadScene(...): cannot create scene"); + } + + LoadMaterials(basepath, scene); + LoadLights(settings.light_file, scene); + return scene; +} diff --git a/BaikalStandalone/Application/scene_load_utils.h b/BaikalStandalone/Application/scene_load_utils.h new file mode 100644 index 00000000..e536201c --- /dev/null +++ b/BaikalStandalone/Application/scene_load_utils.h @@ -0,0 +1,33 @@ +/********************************************************************** + Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved. + + Permission is hereby granted, free of charge, to any person obtaining a copy + of this software and associated documentation files (the "Software"), to deal + in the Software without restriction, including without limitation the rights + to use, copy, modify, merge, publish, distribute, sublicense, and/or sell + copies of the Software, and to permit persons to whom the Software is + furnished to do so, subject to the following conditions: + + The above copyright notice and this permission notice shall be included in + all copies or substantial portions of the Software. + + THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR + IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, + FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE + AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER + LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, + OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN + THE SOFTWARE. + ********************************************************************/ + +#pragma once + +#include +#include "app_utils.h" + +namespace Baikal +{ + class Scene1; +} + +std::shared_ptr LoadScene(Baikal::AppSettings const& settings); diff --git a/BaikalStandalone/CMakeLists.txt b/BaikalStandalone/CMakeLists.txt index 854d0dfc..717c92c7 100644 --- a/BaikalStandalone/CMakeLists.txt +++ b/BaikalStandalone/CMakeLists.txt @@ -14,7 +14,9 @@ set(APPLICATION_SOURCES Application/graph_scheme.h Application/graph_scheme.cpp Application/material_explorer.h - Application/material_explorer.cpp) + Application/material_explorer.cpp + Application/scene_load_utils.h + Application/scene_load_utils.cpp) set(IMGUI_SORUCES ImGUI/imconfig.h @@ -36,7 +38,8 @@ set(UTILS_SOURCES Utils/config_manager.cpp Utils/config_manager.h Utils/shader_manager.cpp - Utils/shader_manager.h) + Utils/shader_manager.h + Utils/output_accessor.h) set(KERNEL_SOURCES Kernels/GLSL/simple.fsh @@ -53,7 +56,7 @@ source_group("Application" FILES ${APPLICATION_SOURCES}) source_group("ImGUI" FILES ${IMGUI_SORUCES}) source_group("Utils" FILES ${UTILS_SOURCES}) source_group("Kernels\\GLSL" FILES ${KERNEL_SOURCES}) - + add_executable(BaikalStandalone ${SOURCES}) target_compile_features(BaikalStandalone PRIVATE cxx_std_14) @@ -64,11 +67,6 @@ target_include_directories(BaikalStandalone target_link_libraries(BaikalStandalone PRIVATE Baikal BaikalIO glfw3::glfw3 OpenGL::GL GLEW::GLEW) -# Compile definitions -if (BAIKAL_ENABLE_DENOISER) - target_compile_definitions(BaikalStandalone PUBLIC ENABLE_DENOISER) -endif(BAIKAL_ENABLE_DENOISER) - if (BAIKAL_ENABLE_GLTF) target_link_libraries(BaikalStandalone PRIVATE RadeonProRender RprSupport) endif (BAIKAL_ENABLE_GLTF) diff --git a/BaikalStandalone/Utils/config_manager.cpp b/BaikalStandalone/Utils/config_manager.cpp index 20f9d70f..8a7de46e 100644 --- a/BaikalStandalone/Utils/config_manager.cpp +++ b/BaikalStandalone/Utils/config_manager.cpp @@ -23,6 +23,7 @@ THE SOFTWARE. #include "CLW.h" #include "RenderFactory/render_factory.h" +#include "Controllers/scene_controller.h" #ifndef APP_BENCHMARK @@ -39,7 +40,10 @@ THE SOFTWARE. #include #endif -void ConfigManager::CreateConfigs( +// Enable forward declarations of T for std::unique_ptr +Config::~Config() = default; + +void CreateConfigs( Mode mode, bool interop, std::vector& configs, @@ -85,10 +89,10 @@ void ConfigManager::CreateConfigs( { if (req_platform_index < 0) { - if ((mode == kUseGpus || mode == kUseSingleGpu) && platforms[i].GetDevice(d).GetType() != CL_DEVICE_TYPE_GPU) + if ((mode == Mode::kUseGpus || mode == Mode::kUseSingleGpu) && platforms[i].GetDevice(d).GetType() != CL_DEVICE_TYPE_GPU) continue; - if ((mode == kUseCpus || mode == kUseSingleCpu) && platforms[i].GetDevice(d).GetType() != CL_DEVICE_TYPE_CPU) + if ((mode == Mode::kUseCpus || mode == Mode::kUseSingleCpu) && platforms[i].GetDevice(d).GetType() != CL_DEVICE_TYPE_CPU) continue; } @@ -132,7 +136,7 @@ void ConfigManager::CreateConfigs( try { cfg.context = CLWContext::Create(platforms[i].GetDevice(d), props); - cfg.type = kPrimary; + cfg.type = DeviceType::kPrimary; cfg.caninterop = true; hasprimary = true; create_without_interop = false; @@ -154,16 +158,16 @@ void ConfigManager::CreateConfigs( if (create_without_interop) { cfg.context = CLWContext::Create(platforms[i].GetDevice(d)); - cfg.type = kSecondary; + cfg.type = DeviceType::kSecondary; } configs.push_back(std::move(cfg)); - if (mode == kUseSingleGpu || mode == kUseSingleCpu) + if (mode == Mode::kUseSingleGpu || mode == Mode::kUseSingleCpu) break; } - if (configs.size() == 1 && (mode == kUseSingleGpu || mode == kUseSingleCpu)) + if (configs.size() == 1 && (mode == Mode::kUseSingleGpu || mode == Mode::kUseSingleCpu)) break; } @@ -175,7 +179,7 @@ void ConfigManager::CreateConfigs( if (!hasprimary) { - configs[0].type = kPrimary; + configs[0].type = DeviceType::kPrimary; } for (std::size_t i = 0; i < configs.size(); ++i) @@ -187,7 +191,7 @@ void ConfigManager::CreateConfigs( } #else -void ConfigManager::CreateConfigs( +void CreateConfigs( Mode mode, bool interop, std::vector& configs, diff --git a/BaikalStandalone/Utils/config_manager.h b/BaikalStandalone/Utils/config_manager.h index e91c619c..985b1b5d 100644 --- a/BaikalStandalone/Utils/config_manager.h +++ b/BaikalStandalone/Utils/config_manager.h @@ -19,12 +19,14 @@ LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. ********************************************************************/ -#ifndef CONFIG_MANAGER_H -#define CONFIG_MANAGER_H + +#pragma once #include "CLW.h" #include "RenderFactory/clw_render_factory.h" #include "Renderers/renderer.h" +#include "SceneGraph/clwscene.h" + #include #include @@ -33,53 +35,49 @@ namespace Baikal class Renderer; } -class ConfigManager +enum class DeviceType { -public: - - enum DeviceType - { - kPrimary, - kSecondary - }; - - enum Mode - { - kUseAll, - kUseGpus, - kUseSingleGpu, - kUseSingleCpu, - kUseCpus - }; - - struct Config - { - DeviceType type; - std::unique_ptr renderer; - std::unique_ptr> controller; - std::unique_ptr> factory; - CLWContext context; - bool caninterop; - - Config() = default; + kPrimary = 0, + kSecondary +}; - Config(Config&& cfg) = default; +enum class PostProcessingType +{ + kNone = 0, + kBilateralDenoiser, + kWaveletDenoser, + kMLDenoiser, + kMLUpsample +}; - ~Config() - { - } - }; +enum class Mode +{ + kUseAll = 0, + kUseGpus, + kUseSingleGpu, + kUseSingleCpu, + kUseCpus +}; - static void CreateConfigs( - Mode mode, - bool interop, - std::vector& renderers, - int initial_num_bounces, - int req_platform_index = -1, - int req_device_index = -1); +struct Config +{ + DeviceType type; + std::unique_ptr renderer; + std::unique_ptr> controller; + std::unique_ptr> factory; + CLWContext context; + bool caninterop; -private: + Config() = default; + Config(Config&& cfg) = default; + ~Config(); }; -#endif // CONFIG_MANAGER_H +void CreateConfigs( + Mode mode, + bool interop, + std::vector& configs, + int initial_num_bounces, + int req_platform_index = -1, + int req_device_index = -1); diff --git a/BaikalStandalone/Utils/output_accessor.h b/BaikalStandalone/Utils/output_accessor.h new file mode 100644 index 00000000..a4d10265 --- /dev/null +++ b/BaikalStandalone/Utils/output_accessor.h @@ -0,0 +1,171 @@ +#pragma once + +#include "Renderers/renderer.h" +#include "Output/clwoutput.h" +#include "RadeonRays/CLW/CLWContext.h" + +#include +#include +#include +#include + + +namespace Baikal +{ + class RendererOutputAccessor + { + public: + RendererOutputAccessor(std::string output_dir, size_t width, size_t height) + : m_output_dir(std::move(output_dir)), m_width(width), m_height(height) + { + m_image_data.resize(m_width * m_height); + m_image_rgb.resize(m_width * m_height); + }; + + template + void SaveAllOutputs(const std::vector& outputs) + { + if (m_frame_count % kDumpPeriod == 0) + { + for (std::size_t idx = 0; idx < outputs.size(); ++idx) + { + for (auto& output : outputs[idx].render_outputs) + { + SaveImageFromRendererOutput(idx, + output.first, + output.second.get(), + m_frame_count); + } + } + } + ++m_frame_count; + } + + void SaveImageFromRendererOutput(size_t device_idx, + Renderer::OutputType output_type, + Output *output, size_t index) + { + GetOutputData(output); + NormalizeImage(); + SaveImage(device_idx, GetOutputName(output_type), index); + } + + void LoadImageToRendererOutput(const CLWContext& context, + Output* output, + const std::string& image_path) + { + using float3 = RadeonRays::float3; + std::ifstream ifs(image_path, std::ios::binary | std::ios::ate); + std::ifstream::pos_type pos = ifs.tellg(); + + std::vector result(pos); + ifs.seekg(0, std::ios::beg); + ifs.read(result.data(), pos); + + auto clw_output = dynamic_cast(output); + + context.WriteBuffer(0, + clw_output->data(), + reinterpret_cast(result.data()), + sizeof(RadeonRays::float3)).Wait(); + } + + private: + struct Color + { + float r; + float g; + float b; + }; + + void GetOutputData(Output* output) + { + output->GetData(&m_image_data[0]); + } + + static std::string GetOutputName(Renderer::OutputType output_type) + { + switch (output_type) + { + case Renderer::OutputType::kColor: + return "color"; + case Renderer::OutputType::kOpacity: + return "opacity"; + case Renderer::OutputType::kVisibility: + return "visibility"; + case Baikal::Renderer::OutputType::kMaxMultiPassOutput: + return "max_multi_pass_output"; + case Baikal::Renderer::OutputType::kMeshID: + return "mesh_id"; + case Renderer::OutputType::kAlbedo: + return "albedo"; + case Renderer::OutputType::kWorldPosition: + return "world_position"; + case Renderer::OutputType::kWorldTangent: + return "world_tangent"; + case Renderer::OutputType::kWorldBitangent: + return "world_bitangent"; + case Renderer::OutputType::kWorldShadingNormal: + return "world_shading_normal"; + case Renderer::OutputType::kWorldGeometricNormal: + return "world_shading_normal"; + case Renderer::OutputType::kViewShadingNormal: + return "view_shading_normal"; + case Renderer::OutputType::kBackground: + return "background"; + case Renderer::OutputType::kDepth: + return "depth"; + case Renderer::OutputType::kGloss: + return "gloss"; + case Renderer::OutputType::kGroupID: + return "group_id"; + case Renderer::OutputType::kMax: + return "max"; + case Renderer::OutputType::kShapeId: + return "shape_id"; + case Renderer::OutputType::kWireframe: + return "wireframe"; + default: + throw std::runtime_error("Output is not supported"); + } + } + + void NormalizeImage() + { + std::transform(m_image_data.cbegin(), m_image_data.cend(), m_image_rgb.begin(), + [](RadeonRays::float3 const &v) + { + if (v.w == 0) + { + return Color({0, 0, 0}); + } + else + { + float invw = 1.f / v.w; + return Color({v.x * invw, v.y * invw, v.z * invw}); + } + }); + } + + void SaveImage(size_t device_idx, std::string const& output_name, std::size_t index) + { + std::ostringstream path; + path << m_output_dir + << "/device_" << device_idx + << "_frame_" << index + << "_" << output_name << ".bin"; + std::ofstream out(path.str(), std::ios_base::binary); + out.write(reinterpret_cast(m_image_rgb.data()), + m_width * m_height * sizeof(Color)); + std::cerr << "Written: " << path.str() << "\n"; + } + + std::vector m_image_data; + std::vector m_image_rgb; + std::string m_output_dir; + std::size_t m_width; + std::size_t m_height; + std::size_t m_frame_count = 0; + static constexpr std::size_t kDumpPeriod = 20; + }; +} diff --git a/BaikalStandalone/Utils/shader_manager.h b/BaikalStandalone/Utils/shader_manager.h index f762b248..a0ac9984 100644 --- a/BaikalStandalone/Utils/shader_manager.h +++ b/BaikalStandalone/Utils/shader_manager.h @@ -29,7 +29,6 @@ THE SOFTWARE. #include "GLFW/glfw3.h" #elif WIN32 #define NOMINMAX -#include #include "GL/glew.h" #else #include diff --git a/BaikalStandalone/cache/placeholder.txt b/BaikalStandalone/cache/placeholder.txt deleted file mode 100644 index cb0bb35c..00000000 --- a/BaikalStandalone/cache/placeholder.txt +++ /dev/null @@ -1 +0,0 @@ -placeholder for output images \ No newline at end of file diff --git a/BaikalStandalone/main.cpp b/BaikalStandalone/main.cpp index c7a82ed6..efcdc2b5 100644 --- a/BaikalStandalone/main.cpp +++ b/BaikalStandalone/main.cpp @@ -20,6 +20,7 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE SOFTWARE. ********************************************************************/ #include "Application/application.h" +#include "except.h" int main(int argc, char * argv[]) { @@ -42,6 +43,11 @@ int main(int argc, char * argv[]) return -1; } + catch(Calc::Exception& ex) + { + std::cout << ex.what() << std::endl; + return -1; + } return 0; } diff --git a/BaikalTest/basic.h b/BaikalTest/basic.h index b69e8418..5781b6f0 100644 --- a/BaikalTest/basic.h +++ b/BaikalTest/basic.h @@ -24,6 +24,7 @@ THE SOFTWARE. #include "gtest/gtest.h" #include "CLW.h" +#include "Controllers/clw_scene_controller.h" #include "Renderers/renderer.h" #include "RenderFactory/clw_render_factory.h" #include "Output/output.h" diff --git a/BaikalTest/test_scenes.h b/BaikalTest/test_scenes.h index d0dbd11d..773bbf6d 100644 --- a/BaikalTest/test_scenes.h +++ b/BaikalTest/test_scenes.h @@ -32,7 +32,7 @@ class TestScenesTest : public BasicTest { m_scene = Baikal::SceneIo::LoadScene(file_name, file_path); } - catch (std::exception &) + catch (std::exception&) { return false; } diff --git a/CMakeLists.txt b/CMakeLists.txt index 5697e0c9..33251bea 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -3,7 +3,6 @@ cmake_minimum_required(VERSION 3.8) project(Baikal CXX) option(BAIKAL_ENABLE_RAYMASK "Enable visibility flags for shapes (slows down an intersector)" OFF) -option(BAIKAL_ENABLE_DENOISER "Use denoising on output" OFF) option(BAIKAL_ENABLE_RPR "Enable RadeonProRender API lib" OFF) option(BAIKAL_ENABLE_TESTS "Enable tests" ON) option(BAIKAL_ENABLE_STANDALONE "Enable standalone application build" ON) @@ -147,3 +146,9 @@ if (WIN32) else () install(FILES $ DESTINATION lib) endif () + +# Ignore warning caused by old tiny_obj_loader.h code +if ((CMAKE_CXX_COMPILER_ID STREQUAL GNU) AND + (CMAKE_CXX_COMPILER_VERSION VERSION_GREATER_EQUAL 8.0.0)) + target_compile_options(Baikal PUBLIC -Wno-class-memaccess) +endif() diff --git a/README.md b/README.md index cef78063..141c46bb 100644 --- a/README.md +++ b/README.md @@ -87,8 +87,6 @@ git clone --recursive https://github.com/GPUOpen-LibrariesAndSDKs/RadeonProRende Available premake options: -- `BAIKAL_ENABLE_DENOISER` enables EAW denoiser in interactive output: - - `BAIKAL_ENABLE_RPR` generates RadeonProRender API implemenatiton C-library and couple of RPR tutorials. ## Run diff --git a/RadeonRays b/RadeonRays index 58aad072..ff90589a 160000 --- a/RadeonRays +++ b/RadeonRays @@ -1 +1 @@ -Subproject commit 58aad072734ac79ebe591280bc788fce03cc8a3c +Subproject commit ff90589afbfe27d72d3447dc08bdb4fce92e6d77 diff --git a/Resources/Textures/pano_port_001.jpg b/Resources/Textures/pano_port_001.jpg new file mode 100644 index 00000000..887e8a57 Binary files /dev/null and b/Resources/Textures/pano_port_001.jpg differ