accum_buffer_B use half

zenovis/xinxinoptix/PTKernel.cu

@@ -119,6 +119,20 @@ float3 half3_to_float3(ushort3 in)
     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()
 {
 
@@ -388,7 +402,11 @@ extern "C" __global__ void __raygen__rg()
         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 = half3_to_float3(params.accum_buffer_B[ 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    = 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 );
@@ -411,7 +429,7 @@ extern "C" __global__ void __raygen__rg()
     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 ] = float3_to_half3(accum_color_b);
+    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;

zenovis/xinxinoptix/TypeCaster.cpp

@@ -29,4 +29,8 @@ float3 toFloat(ushort3 in) {
         __half2float(y),
         __half2float(z),
     };
+}
+float toFloat(ushort1 in) {
+    half x = reinterpret_cast<half&>(in);
+    return __half2float(x);
 }

zenovis/xinxinoptix/TypeCaster.h

@@ -4,4 +4,5 @@
 
 ushort3 toHalf(float4 in);
 ushort3 toHalf(float3 in);
-float3  toFloat(ushort3 in);
+float3  toFloat(ushort3 in);
+float  toFloat(ushort1 in);

zenovis/xinxinoptix/optixPathTracer.cpp

@@ -634,7 +634,7 @@ static void handleResize( sutil::CUDAOutputBuffer<uchar4>& output_buffer, Params
             ) );
     CUDA_CHECK( cudaMalloc(
         reinterpret_cast<void**>( &state.accum_buffer_b .reset()),
-        params.width * params.height * sizeof( ushort3 )
+        params.width * params.height * sizeof( ushort1 )
             ) );
     state.params.accum_buffer = (float3*)(CUdeviceptr)state.accum_buffer_p;
 
@@ -653,7 +653,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.accum_buffer_B = (ushort3*)(CUdeviceptr)state.accum_buffer_b;
+    state.params.accum_buffer_B = (ushort1*)(CUdeviceptr)state.accum_buffer_b;
     state.params.subframe_index = 0;
 }
 
@@ -3736,13 +3736,13 @@ std::vector<float> optixgetimg_extra2(std::string name, int w, int h) {
         cudaMemcpy(tex_data.data(), (void*)state.accum_buffer_t.handle, sizeof(float) * tex_data.size(), cudaMemcpyDeviceToHost);
     }
     else if (name == "background") {
-        std::vector<ushort3> temp_buffer(w * h);
-        cudaMemcpy(temp_buffer.data(), (void*)state.accum_buffer_b.handle, sizeof(ushort3) * temp_buffer.size(), cudaMemcpyDeviceToHost);
+        std::vector<ushort1> temp_buffer(w * h);
+        cudaMemcpy(temp_buffer.data(), (void*)state.accum_buffer_b.handle, sizeof(ushort1) * 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;
+            float v = toFloat(temp_buffer[i]);
+            tex_data[i * 3 + 0] = v;
+            tex_data[i * 3 + 1] = v;
+            tex_data[i * 3 + 2] = v;
         }
     }
     else if (name == "mask") {

zenovis/xinxinoptix/optixPathTracer.h

@@ -158,7 +158,7 @@ struct Params
     float3*      accum_buffer_D;
     float3*      accum_buffer_S;
     float3*      accum_buffer_T;
-    ushort3*     accum_buffer_B;
+    ushort1*     accum_buffer_B;
     uchar4*      frame_buffer;
     float3*      frame_buffer_M;