Skip to content

Commit

Permalink
use half3
Browse files Browse the repository at this point in the history
  • Loading branch information
@zhouhang95
zhouhang95 committed Jun 19, 2024
1 parent 353b674 commit f9c0f99
Show file tree
Hide file tree
Showing 3 changed files with 69 additions and 32 deletions.
40 changes: 33 additions & 7 deletions zenovis/xinxinoptix/PTKernel.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -92,6 +92,32 @@ vec3 PhysicalCamera(vec3 in,
mapped = in * exposure;
return enableExposure? (enableACES? ACESFilm(mapped):mapped ) : (enableACES? ACESFilm(in) : in);
}

static __inline__ __device__
ushort3 float3_to_half3(float3 in)
{
half x = __float2half(in.x);
half y = __float2half(in.y);
half z = __float2half(in.z);
ushort3 v;
v.x = reinterpret_cast<unsigned short&>(x);
v.y = reinterpret_cast<unsigned short&>(y);
v.z = reinterpret_cast<unsigned short&>(z);
return v;
}

static __inline__ __device__
float3 half3_to_float3(ushort3 in)
{
half x = reinterpret_cast<half&>(in.x);
half y = reinterpret_cast<half&>(in.y);
half z = reinterpret_cast<half&>(in.z);
float3 v;
v.x = __half2float(x);
v.y = __half2float(y);
v.z = __half2float(z);
return v;
}
extern "C" __global__ void __raygen__rg()
{

Expand Down Expand Up @@ -362,7 +388,7 @@ extern "C" __global__ void __raygen__rg()
const float3 accum_color_prev_s = make_float3( params.accum_buffer_S[ image_index ]);
const float3 accum_color_prev_t = make_float3( params.accum_buffer_T[ image_index ]);
const float3 accum_color_prev_b = make_float3( params.accum_buffer_B[ image_index ]);
const float3 accum_mask_prev = params.frame_buffer_M[ image_index ];
const float3 accum_mask_prev = half3_to_float3(params.frame_buffer_M[ image_index ]);
accum_color = mix( vec3(accum_color_prev), accum_color, a );
accum_color_d = mix( vec3(accum_color_prev_d), accum_color_d, a );
accum_color_s = mix( vec3(accum_color_prev_s), accum_color_s, a );
Expand Down Expand Up @@ -399,12 +425,12 @@ extern "C" __global__ void __raygen__rg()
float3 out_color_t = t_mapped;
float3 out_color_b = accum_color_b;
params.frame_buffer[ image_index ] = make_color ( out_color );
params.frame_buffer_C[ image_index ] = out_color;
params.frame_buffer_D[ image_index ] = out_color_d;
params.frame_buffer_S[ image_index ] = out_color_s;
params.frame_buffer_T[ image_index ] = out_color_t;
params.frame_buffer_B[ image_index ] = accum_color_b;
params.frame_buffer_M[ image_index ] = accum_mask;
params.frame_buffer_C[ image_index ] = float3_to_half3(out_color);
params.frame_buffer_D[ image_index ] = float3_to_half3(out_color_d);
params.frame_buffer_S[ image_index ] = float3_to_half3(out_color_s);
params.frame_buffer_T[ image_index ] = float3_to_half3(out_color_t);
params.frame_buffer_B[ image_index ] = float3_to_half3(accum_color_b);
params.frame_buffer_M[ image_index ] = float3_to_half3(accum_mask);

if (params.denoise) {
params.albedo_buffer[ image_index ] = tmp_albedo;
Expand Down
49 changes: 30 additions & 19 deletions zenovis/xinxinoptix/optixPathTracer.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -253,12 +253,12 @@ ushort2 halfNormal(float4 in)
#endif

std::optional<sutil::CUDAOutputBuffer<uchar4>> output_buffer_o;
std::optional<sutil::CUDAOutputBuffer<float3>> output_buffer_color;
std::optional<sutil::CUDAOutputBuffer<float3>> output_buffer_diffuse;
std::optional<sutil::CUDAOutputBuffer<float3>> output_buffer_specular;
std::optional<sutil::CUDAOutputBuffer<float3>> output_buffer_transmit;
std::optional<sutil::CUDAOutputBuffer<float3>> output_buffer_background;
std::optional<sutil::CUDAOutputBuffer<float3>> output_buffer_mask;
std::optional<sutil::CUDAOutputBuffer<ushort3>> output_buffer_color;
std::optional<sutil::CUDAOutputBuffer<ushort3>> output_buffer_diffuse;
std::optional<sutil::CUDAOutputBuffer<ushort3>> output_buffer_specular;
std::optional<sutil::CUDAOutputBuffer<ushort3>> output_buffer_transmit;
std::optional<sutil::CUDAOutputBuffer<ushort3>> output_buffer_background;
std::optional<sutil::CUDAOutputBuffer<ushort3>> output_buffer_mask;
using Vertex = float4;

struct PathTracerState
Expand Down Expand Up @@ -3805,6 +3805,15 @@ void *optixgetimg_extra(std::string name) {
}
throw std::runtime_error("invalid optixgetimg_extra name: " + name);
}

std::vector<float> optixgetimg_extra(std::string name, int w, int h) {
std::vector<float> data(w * h * 3);
short *ptr = (short*)optixgetimg_extra(name);
for (auto i = 0; i < data.size(); i++) {
data[i] = glm::detail::toFloat32(ptr[i]);
}
return data;
}
static void save_exr(float3* ptr, int w, int h, std::string path) {
std::vector<float3> data(w * h);
std::copy_n(ptr, w * h, data.data());
Expand Down Expand Up @@ -3879,20 +3888,20 @@ void optixrender(int fbo, int samples, bool denoise, bool simpleRender) {
auto exr_path = path.substr(0, path.size() - 4) + ".exr";
if (enable_output_mask) {
path = path.substr(0, path.size() - 4);
save_png_data(path + "_mask.png", w, h, (float*)optixgetimg_extra("mask"));
save_png_data(path + "_mask.png", w, h, optixgetimg_extra("mask", w, h).data());
}
// AOV
if (enable_output_aov) {
if (enable_output_exr) {
zeno::create_directories_when_write_file(exr_path);
SaveMultiLayerEXR(
{
(float*)optixgetimg_extra("color"),
(float*)optixgetimg_extra("diffuse"),
(float*)optixgetimg_extra("specular"),
(float*)optixgetimg_extra("transmit"),
(float*)optixgetimg_extra("background"),
(float*)optixgetimg_extra("mask"),
optixgetimg_extra("color", w, h).data(),
optixgetimg_extra("diffuse", w, h).data(),
optixgetimg_extra("specular", w, h).data(),
optixgetimg_extra("transmit", w, h).data(),
optixgetimg_extra("background", w, h).data(),
optixgetimg_extra("mask", w, h).data(),
},
w,
h,
Expand All @@ -3909,18 +3918,20 @@ void optixrender(int fbo, int samples, bool denoise, bool simpleRender) {

}
else {
zeno::create_directories_when_write_file(exr_path);
save_exr((float3 *)optixgetimg_extra("color", w, h).data(), w, h, exr_path);
path = path.substr(0, path.size() - 4);
save_png_color(path + ".aov.diffuse.png", w, h, (float*)optixgetimg_extra("diffuse"));
save_png_color(path + ".aov.specular.png", w, h, (float*)optixgetimg_extra("specular"));
save_png_color(path + ".aov.transmit.png", w, h, (float*)optixgetimg_extra("transmit"));
save_png_data(path + ".aov.background.png", w, h, (float*)optixgetimg_extra("background"));
save_png_data(path + ".aov.mask.png", w, h, (float*)optixgetimg_extra("mask"));
save_png_color(path + ".aov.diffuse.png", w, h, optixgetimg_extra("diffuse", w, h).data());
save_png_color(path + ".aov.specular.png", w, h, optixgetimg_extra("specular", w, h).data());
save_png_color(path + ".aov.transmit.png", w, h, optixgetimg_extra("transmit", w, h).data());
save_png_data(path + ".aov.background.png", w, h, optixgetimg_extra("background", w, h).data());
save_png_data(path + ".aov.mask.png", w, h, optixgetimg_extra("mask", w, h).data());
}
}
else {
if (enable_output_exr) {
zeno::create_directories_when_write_file(exr_path);
save_exr((float3 *)optixgetimg_extra("color"), w, h, exr_path);
save_exr((float3 *)optixgetimg_extra("color", w, h).data(), w, h, exr_path);
}
else {
std::string jpg_native_path = zeno::create_directories_when_write_file(path);
Expand Down
12 changes: 6 additions & 6 deletions zenovis/xinxinoptix/optixPathTracer.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -160,12 +160,12 @@ struct Params
float4* accum_buffer_T;
float4* accum_buffer_B;
uchar4* frame_buffer;
float3* frame_buffer_C;
float3* frame_buffer_D;
float3* frame_buffer_S;
float3* frame_buffer_T;
float3* frame_buffer_B;
float3* frame_buffer_M;
ushort3* frame_buffer_C;
ushort3* frame_buffer_D;
ushort3* frame_buffer_S;
ushort3* frame_buffer_T;
ushort3* frame_buffer_B;
ushort3* frame_buffer_M;

float3* debug_buffer;
float3* albedo_buffer;
Expand Down

0 comments on commit f9c0f99

Please sign in to comment.