Merge pull request #1946 from zenustech/optix-shrink-screen-buffer-v2

Optix shrink screen buffer v2
zenustech · Jun 24, 2024 · 97902a9 · 97902a9
2 parents 2e4aa77 + 4c446cb
commit 97902a9
Show file tree

Hide file tree

Showing 6 changed files with 226 additions and 155 deletions.
diff --git a/zenovis/xinxinoptix/ChiefDesignerEXR.h b/zenovis/xinxinoptix/ChiefDesignerEXR.h
@@ -161,4 +161,40 @@ inline void SaveMultiLayerEXR(
     file.writePixels (height);
 }
 
+inline void SaveMultiLayerEXR_half(
+    std::vector<half*> pixels
+    , int width
+    , int height
+    , std::vector<std::string> channels
+    , const char* exrFilePath
+) {
+    using namespace Imath;
+    using namespace Imf;
+
+    Header header(width, height);
+    ChannelList channelList;
+
+    const char *std_suffix = "RGB";
+    for (auto channel: channels) {
+        for (int i = 0; i < 3; i++) {
+            std::string name = zeno::format("{}{}", channel, std_suffix[i]);
+            channelList.insert(name, Channel(HALF));
+        }
+    }
+
+    header.channels() = channelList;
+
+    OutputFile file (exrFilePath, header);
+    FrameBuffer frameBuffer;
+
+    for (auto i = 0; i < channels.size(); i++) {
+        frameBuffer.insert (zeno::format("{}R", channels[i]), Slice ( HALF, (char*) &pixels[i][0], sizeof (half) * 3, sizeof (half) * width * 3));
+        frameBuffer.insert (zeno::format("{}G", channels[i]), Slice ( HALF, (char*) &pixels[i][1], sizeof (half) * 3, sizeof (half) * width * 3));
+        frameBuffer.insert (zeno::format("{}B", channels[i]), Slice ( HALF, (char*) &pixels[i][2], sizeof (half) * 3, sizeof (half) * width * 3));
+    }
+
+    file.setFrameBuffer (frameBuffer);
+    file.writePixels (height);
+}
+
 }
diff --git a/zenovis/xinxinoptix/PTKernel.cu b/zenovis/xinxinoptix/PTKernel.cu
@@ -92,6 +92,47 @@ 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;
+}
+
+static __inline__ __device__
+ushort1 float_to_half(float in)
+{
+    half x = __float2half(in);
+    return reinterpret_cast<ushort1&>(x);
+}
+
+static __inline__ __device__
+float half_to_float(ushort1 in)
+{
+    half x = reinterpret_cast<half&>(in);
+    return __half2float(x);
+}
+
 extern "C" __global__ void __raygen__rg()
 {
 
@@ -357,12 +398,16 @@ extern "C" __global__ void __raygen__rg()
     if( subframe_index > 0 )
     {
         const float                 a = 1.0f / static_cast<float>( subframe_index+1 );
-        const float3 accum_color_prev = make_float3( params.accum_buffer[ image_index ]);
-        const float3 accum_color_prev_d = make_float3( params.accum_buffer_D[ image_index ]);
-        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_color_prev   = params.accum_buffer[ image_index ];
+        const float3 accum_color_prev_d = params.accum_buffer_D[ image_index ];
+        const float3 accum_color_prev_s = params.accum_buffer_S[ image_index ];
+        const float3 accum_color_prev_t = params.accum_buffer_T[ image_index ];
+        const float3 accum_color_prev_b = {
+                half_to_float(params.accum_buffer_B[ image_index ]),
+                half_to_float(params.accum_buffer_B[ image_index ]),
+                half_to_float(params.accum_buffer_B[ 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 );
@@ -380,31 +425,14 @@ extern "C" __global__ void __raygen__rg()
         }
     }
 
-    params.accum_buffer[ image_index ] = make_float4( accum_color.x, accum_color.y, accum_color.z, 1.0f);
-    params.accum_buffer_D[ image_index ] = make_float4( accum_color_d.x,accum_color_d.y,accum_color_d.z, 1.0f);
-    params.accum_buffer_S[ image_index ] = make_float4( accum_color_s.x,accum_color_s.y, accum_color_s.z, 1.0f);
-    params.accum_buffer_T[ image_index ] = make_float4( accum_color_t.x,accum_color_t.y,accum_color_t.z, 1.0f);
-    params.accum_buffer_B[ image_index ] = make_float4( accum_color_b, 1.0f);
-
-
-    vec3 rgb_mapped = PhysicalCamera(vec3(accum_color), aperture, shutter_speed, iso, midGray, false, false);
-    vec3 d_mapped = PhysicalCamera(vec3(accum_color_d), aperture, shutter_speed, iso, midGray, false, false);
-    vec3 s_mapped = PhysicalCamera(vec3(accum_color_s), aperture, shutter_speed, iso, midGray, false, false);
-    vec3 t_mapped = PhysicalCamera(vec3(accum_color_t), aperture, shutter_speed, iso, midGray, false, false);
-
-
-    float3 out_color = rgb_mapped;
-    float3 out_color_d = d_mapped;
-    float3 out_color_s = s_mapped;
-    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.accum_buffer[ image_index ]   = make_float3( accum_color.x, accum_color.y, accum_color.z);
+    params.accum_buffer_D[ image_index ] = make_float3( accum_color_d.x,accum_color_d.y,accum_color_d.z);
+    params.accum_buffer_S[ image_index ] = make_float3( accum_color_s.x,accum_color_s.y, accum_color_s.z);
+    params.accum_buffer_T[ image_index ] = make_float3( accum_color_t.x,accum_color_t.y,accum_color_t.z);
+    params.accum_buffer_B[ image_index ] = float_to_half(accum_color_b.x);
+
+    params.frame_buffer[ image_index ] = make_color ( accum_color );
+    params.frame_buffer_M[ image_index ] = float3_to_half3(accum_mask);
 
     if (params.denoise) {
         params.albedo_buffer[ image_index ] = tmp_albedo;

diff --git a/zenovis/xinxinoptix/TypeCaster.cpp b/zenovis/xinxinoptix/TypeCaster.cpp
@@ -18,4 +18,19 @@ ushort3 toHalf(float4 in)
 ushort3 toHalf(float3 in)
 {
   return toHalf({in.x, in.y, in.z, 0.0f});
+}
+
+float3 toFloat(ushort3 in) {
+    half x = reinterpret_cast<half&>(in.x);
+    half y = reinterpret_cast<half&>(in.y);
+    half z = reinterpret_cast<half&>(in.z);
+    return {
+        __half2float(x),
+        __half2float(y),
+        __half2float(z),
+    };
+}
+float toFloat(ushort1 in) {
+    half x = reinterpret_cast<half&>(in);
+    return __half2float(x);
 }
diff --git a/zenovis/xinxinoptix/TypeCaster.h b/zenovis/xinxinoptix/TypeCaster.h
@@ -3,4 +3,6 @@
 #include <vector_types.h>
 
 ushort3 toHalf(float4 in);
-ushort3 toHalf(float3 in);
+ushort3 toHalf(float3 in);
+float3  toFloat(ushort3 in);
+float  toFloat(ushort1 in);