frame_buffer_M use half3

zenustech · Jun 20, 2024 · 0551d20 · 0551d20
1 parent 7f09426
commit 0551d20
Show file tree

Hide file tree

Showing 3 changed files with 17 additions and 17 deletions.
diff --git a/zenovis/xinxinoptix/PTKernel.cu b/zenovis/xinxinoptix/PTKernel.cu
@@ -407,7 +407,7 @@ extern "C" __global__ void __raygen__rg()
                 half_to_float(params.accum_buffer_B[ image_index ]),
                 half_to_float(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 );
@@ -432,7 +432,7 @@ extern "C" __global__ void __raygen__rg()
     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 ] = accum_mask;
+    params.frame_buffer_M[ image_index ] = float3_to_half3(accum_mask);
 
     if (params.denoise) {
         params.albedo_buffer[ image_index ] = float3_to_half3(tmp_albedo);

diff --git a/zenovis/xinxinoptix/optixPathTracer.cpp b/zenovis/xinxinoptix/optixPathTracer.cpp
@@ -253,7 +253,6 @@ ushort2 halfNormal(float4 in)
 #endif
 
 std::optional<sutil::CUDAOutputBuffer<uchar4>> output_buffer_o;
-std::optional<sutil::CUDAOutputBuffer<float3>> output_buffer_mask;
 using Vertex = float4;
 
 struct PathTracerState
@@ -300,6 +299,7 @@ struct PathTracerState
     raii<CUdeviceptr> accum_buffer_s;
     raii<CUdeviceptr> accum_buffer_t;
     raii<CUdeviceptr> accum_buffer_b;
+    raii<CUdeviceptr> accum_buffer_m;
 
     raii<CUdeviceptr> finite_lights_ptr;
 
@@ -613,7 +613,6 @@ static void handleResize( sutil::CUDAOutputBuffer<uchar4>& output_buffer, Params
     resize_dirty = false;
 
     output_buffer.resize( params.width, params.height );
-    (*output_buffer_mask).resize( params.width, params.height );
 
     // Realloc accumulation buffer
     CUDA_CHECK( cudaMalloc(
@@ -632,6 +631,10 @@ static void handleResize( sutil::CUDAOutputBuffer<uchar4>& output_buffer, Params
         reinterpret_cast<void**>( &state.accum_buffer_t .reset()),
         params.width * params.height * sizeof( float3 )
             ) );
+    CUDA_CHECK( cudaMalloc(
+        reinterpret_cast<void**>( &state.accum_buffer_m .reset()),
+        params.width * params.height * sizeof( ushort3 )
+            ) );
     CUDA_CHECK( cudaMalloc(
         reinterpret_cast<void**>( &state.accum_buffer_b .reset()),
         params.width * params.height * sizeof( ushort1 )
@@ -653,6 +656,7 @@ static void handleResize( sutil::CUDAOutputBuffer<uchar4>& output_buffer, Params
     state.params.accum_buffer_D = (float3*)(CUdeviceptr)state.accum_buffer_d;
     state.params.accum_buffer_S = (float3*)(CUdeviceptr)state.accum_buffer_s;
     state.params.accum_buffer_T = (float3*)(CUdeviceptr)state.accum_buffer_t;
+    state.params.frame_buffer_M = (ushort3*)(CUdeviceptr)state.accum_buffer_m;
     state.params.accum_buffer_B = (ushort1*)(CUdeviceptr)state.accum_buffer_b;
     state.params.subframe_index = 0;
 }
@@ -677,7 +681,6 @@ static void launchSubframe( sutil::CUDAOutputBuffer<uchar4>& output_buffer, Path
     // Launch
     uchar4* result_buffer_data = output_buffer.map();
     state.params.frame_buffer  = result_buffer_data;
-    state.params.frame_buffer_M = (*output_buffer_mask      ).map();
     state.params.num_lights = lightsWrapper.g_lights.size();
     state.params.denoise = denoise;
     for(int j=0;j<1;j++){
@@ -710,7 +713,6 @@ static void launchSubframe( sutil::CUDAOutputBuffer<uchar4>& output_buffer, Path
       }
     }
     output_buffer.unmap();
-    (*output_buffer_mask      ).unmap();
 
     try {
         CUDA_SYNC_CHECK();
@@ -1607,14 +1609,6 @@ void optixinit( int argc, char* argv[] )
       );
       output_buffer_o->setStream( 0 );
     }
-    if (!output_buffer_mask) {
-      output_buffer_mask.emplace(
-          output_buffer_type,
-          state.params.width,
-          state.params.height
-      );
-      output_buffer_mask->setStream( 0 );
-    }
 #ifdef OPTIX_BASE_GL
         if (!gl_display_o) {
             gl_display_o.emplace(sutil::BufferImageFormat::UNSIGNED_BYTE4);
@@ -3746,7 +3740,14 @@ std::vector<float> optixgetimg_extra2(std::string name, int w, int h) {
         }
     }
     else if (name == "mask") {
-        std::copy_n((float*) output_buffer_mask->getHostPointer(), tex_data.size(), tex_data.data());
+        std::vector<ushort3> temp_buffer(w * h);
+        cudaMemcpy(temp_buffer.data(), (void*)state.accum_buffer_m.handle, sizeof(ushort3) * temp_buffer.size(), cudaMemcpyDeviceToHost);
+        for (auto i = 0; i < temp_buffer.size(); i++) {
+            float3 v = toFloat(temp_buffer[i]);
+            tex_data[i * 3 + 0] = v.x;
+            tex_data[i * 3 + 1] = v.y;
+            tex_data[i * 3 + 2] = v.z;
+        }
     }
     else if (name == "color") {
         cudaMemcpy(tex_data.data(), (void*)state.accum_buffer_p.handle, sizeof(float) * tex_data.size(), cudaMemcpyDeviceToHost);
@@ -3978,7 +3979,6 @@ void optixDestroy() {
     OptixUtil::shaderCoreLUT.clear();
 
     output_buffer_o           .reset();
-    output_buffer_mask        .reset();
     g_StaticMeshPieces        .clear();
     g_meshPieces              .clear();
     state = {};

diff --git a/zenovis/xinxinoptix/optixPathTracer.h b/zenovis/xinxinoptix/optixPathTracer.h
@@ -160,7 +160,7 @@ struct Params
     float3*      accum_buffer_T;
     ushort1*     accum_buffer_B;
     uchar4*      frame_buffer;
-    float3*      frame_buffer_M;
+    ushort3*     frame_buffer_M;
 
     float3*      debug_buffer;
     ushort3*     albedo_buffer;