diff --git a/zenovis/xinxinoptix/ChiefDesignerEXR.h b/zenovis/xinxinoptix/ChiefDesignerEXR.h index d7c1835176..31e270a318 100644 --- 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 pixels + , int width + , int height + , std::vector 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 index e5ee736af3..843ce5bf47 100644 --- 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(x); + v.y = reinterpret_cast(y); + v.z = reinterpret_cast(z); + return v; +} + +static __inline__ __device__ +float3 half3_to_float3(ushort3 in) +{ + half x = reinterpret_cast(in.x); + half y = reinterpret_cast(in.y); + half z = reinterpret_cast(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(x); +} + +static __inline__ __device__ +float half_to_float(ushort1 in) +{ + half x = reinterpret_cast(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( 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 index eba49af49d..6f60aea0bc 100644 --- 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(in.x); + half y = reinterpret_cast(in.y); + half z = reinterpret_cast(in.z); + return { + __half2float(x), + __half2float(y), + __half2float(z), + }; +} +float toFloat(ushort1 in) { + half x = reinterpret_cast(in); + return __half2float(x); } \ No newline at end of file diff --git a/zenovis/xinxinoptix/TypeCaster.h b/zenovis/xinxinoptix/TypeCaster.h index 519fb09e8e..c208adc72a 100644 --- a/zenovis/xinxinoptix/TypeCaster.h +++ b/zenovis/xinxinoptix/TypeCaster.h @@ -3,4 +3,6 @@ #include ushort3 toHalf(float4 in); -ushort3 toHalf(float3 in); \ No newline at end of file +ushort3 toHalf(float3 in); +float3 toFloat(ushort3 in); +float toFloat(ushort1 in); \ No newline at end of file diff --git a/zenovis/xinxinoptix/optixPathTracer.cpp b/zenovis/xinxinoptix/optixPathTracer.cpp index ad38fddad9..ad0d2ffcc8 100644 --- a/zenovis/xinxinoptix/optixPathTracer.cpp +++ b/zenovis/xinxinoptix/optixPathTracer.cpp @@ -253,12 +253,6 @@ ushort2 halfNormal(float4 in) #endif std::optional> output_buffer_o; -std::optional> output_buffer_color; -std::optional> output_buffer_diffuse; -std::optional> output_buffer_specular; -std::optional> output_buffer_transmit; -std::optional> output_buffer_background; -std::optional> output_buffer_mask; using Vertex = float4; struct PathTracerState @@ -305,6 +299,7 @@ struct PathTracerState raii accum_buffer_s; raii accum_buffer_t; raii accum_buffer_b; + raii accum_buffer_m; raii finite_lights_ptr; @@ -568,9 +563,9 @@ static void initLaunchParams( PathTracerState& state ) CUDA_CHECK( cudaMalloc( reinterpret_cast( &state.accum_buffer_p.reset() ), - state.params.width * state.params.height * sizeof( float4 ) + state.params.width * state.params.height * sizeof( float3 ) ) ); - state.params.accum_buffer = (float4*)(CUdeviceptr)state.accum_buffer_p; + state.params.accum_buffer = (float3*)(CUdeviceptr)state.accum_buffer_p; auto& params = state.params; @@ -618,35 +613,33 @@ static void handleResize( sutil::CUDAOutputBuffer& output_buffer, Params resize_dirty = false; output_buffer.resize( params.width, params.height ); - (*output_buffer_color).resize( params.width, params.height ); - (*output_buffer_diffuse).resize( params.width, params.height ); - (*output_buffer_specular).resize( params.width, params.height ); - (*output_buffer_transmit).resize( params.width, params.height ); - (*output_buffer_background).resize( params.width, params.height ); - (*output_buffer_mask).resize( params.width, params.height ); // Realloc accumulation buffer CUDA_CHECK( cudaMalloc( reinterpret_cast( &state.accum_buffer_p .reset()), - params.width * params.height * sizeof( float4 ) + params.width * params.height * sizeof( float3 ) ) ); CUDA_CHECK( cudaMalloc( reinterpret_cast( &state.accum_buffer_d .reset()), - params.width * params.height * sizeof( float4 ) + params.width * params.height * sizeof( float3 ) ) ); CUDA_CHECK( cudaMalloc( reinterpret_cast( &state.accum_buffer_s .reset()), - params.width * params.height * sizeof( float4 ) + params.width * params.height * sizeof( float3 ) ) ); CUDA_CHECK( cudaMalloc( reinterpret_cast( &state.accum_buffer_t .reset()), - params.width * params.height * sizeof( float4 ) + params.width * params.height * sizeof( float3 ) + ) ); + CUDA_CHECK( cudaMalloc( + reinterpret_cast( &state.accum_buffer_m .reset()), + params.width * params.height * sizeof( ushort3 ) ) ); CUDA_CHECK( cudaMalloc( reinterpret_cast( &state.accum_buffer_b .reset()), - params.width * params.height * sizeof( float4 ) + params.width * params.height * sizeof( ushort1 ) ) ); - state.params.accum_buffer = (float4*)(CUdeviceptr)state.accum_buffer_p; + state.params.accum_buffer = (float3*)(CUdeviceptr)state.accum_buffer_p; CUDA_CHECK( cudaMallocManaged( reinterpret_cast( &state.albedo_buffer_p.reset()), @@ -660,10 +653,11 @@ static void handleResize( sutil::CUDAOutputBuffer& output_buffer, Params ) ); state.params.normal_buffer = (float3*)(CUdeviceptr)state.normal_buffer_p; - state.params.accum_buffer_D = (float4*)(CUdeviceptr)state.accum_buffer_d; - state.params.accum_buffer_S = (float4*)(CUdeviceptr)state.accum_buffer_s; - state.params.accum_buffer_T = (float4*)(CUdeviceptr)state.accum_buffer_t; - state.params.accum_buffer_B = (float4*)(CUdeviceptr)state.accum_buffer_b; + 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; } @@ -687,12 +681,6 @@ static void launchSubframe( sutil::CUDAOutputBuffer& output_buffer, Path // Launch uchar4* result_buffer_data = output_buffer.map(); state.params.frame_buffer = result_buffer_data; - state.params.frame_buffer_C = (*output_buffer_color ).map(); - state.params.frame_buffer_D = (*output_buffer_diffuse ).map(); - state.params.frame_buffer_S = (*output_buffer_specular ).map(); - state.params.frame_buffer_T = (*output_buffer_transmit ).map(); - state.params.frame_buffer_B = (*output_buffer_background).map(); - 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++){ @@ -725,12 +713,6 @@ static void launchSubframe( sutil::CUDAOutputBuffer& output_buffer, Path } } output_buffer.unmap(); - (*output_buffer_color ).unmap(); - (*output_buffer_diffuse ).unmap(); - (*output_buffer_specular ).unmap(); - (*output_buffer_transmit ).unmap(); - (*output_buffer_background).unmap(); - (*output_buffer_mask ).unmap(); try { CUDA_SYNC_CHECK(); @@ -1613,54 +1595,6 @@ void optixinit( int argc, char* argv[] ) ); output_buffer_o->setStream( 0 ); } - if (!output_buffer_color) { - output_buffer_color.emplace( - output_buffer_type, - state.params.width, - state.params.height - ); - output_buffer_color->setStream( 0 ); - } - if (!output_buffer_diffuse) { - output_buffer_diffuse.emplace( - output_buffer_type, - state.params.width, - state.params.height - ); - output_buffer_diffuse->setStream( 0 ); - } - if (!output_buffer_specular) { - output_buffer_specular.emplace( - output_buffer_type, - state.params.width, - state.params.height - ); - output_buffer_specular->setStream( 0 ); - } - if (!output_buffer_transmit) { - output_buffer_transmit.emplace( - output_buffer_type, - state.params.width, - state.params.height - ); - output_buffer_transmit->setStream( 0 ); - } - if (!output_buffer_background) { - output_buffer_background.emplace( - output_buffer_type, - state.params.width, - state.params.height - ); - output_buffer_background->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); @@ -3773,26 +3707,93 @@ void set_outside_random_number(int32_t outside_random_number) { state.params.outside_random_number = outside_random_number; } -void *optixgetimg_extra(std::string name) { +std::vector optixgetimg_extra2(std::string name, int w, int h) { + std::vector tex_data(w * h * 3); + if (name == "diffuse") { + cudaMemcpy(tex_data.data(), (void*)state.accum_buffer_d.handle, sizeof(float) * tex_data.size(), cudaMemcpyDeviceToHost); + } + else if (name == "specular") { + cudaMemcpy(tex_data.data(), (void*)state.accum_buffer_s.handle, sizeof(float) * tex_data.size(), cudaMemcpyDeviceToHost); + } + else if (name == "transmit") { + cudaMemcpy(tex_data.data(), (void*)state.accum_buffer_t.handle, sizeof(float) * tex_data.size(), cudaMemcpyDeviceToHost); + } + else if (name == "background") { + std::vector 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++) { + 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") { + std::vector 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); + } + else { + throw std::runtime_error("invalid optixgetimg_extra name: " + name); + } + return tex_data; +} + +std::vector optixgetimg_extra3(std::string name, int w, int h) { + std::vector tex_data(w * h * 3); if (name == "diffuse") { - return output_buffer_diffuse->getHostPointer(); + std::vector temp_buffer(w * h * 3); + cudaMemcpy(temp_buffer.data(), (void*)state.accum_buffer_d.handle, sizeof(temp_buffer[0]) * temp_buffer.size(), cudaMemcpyDeviceToHost); + for (auto i = 0; i < temp_buffer.size(); i++) { + tex_data[i] = temp_buffer[i]; + } } else if (name == "specular") { - return output_buffer_specular->getHostPointer(); + std::vector temp_buffer(w * h * 3); + cudaMemcpy(temp_buffer.data(), (void*)state.accum_buffer_s.handle, sizeof(temp_buffer[0]) * temp_buffer.size(), cudaMemcpyDeviceToHost); + for (auto i = 0; i < temp_buffer.size(); i++) { + tex_data[i] = temp_buffer[i]; + } } else if (name == "transmit") { - return output_buffer_transmit->getHostPointer(); + std::vector temp_buffer(w * h * 3); + cudaMemcpy(temp_buffer.data(), (void*)state.accum_buffer_t.handle, sizeof(temp_buffer[0]) * temp_buffer.size(), cudaMemcpyDeviceToHost); + for (auto i = 0; i < temp_buffer.size(); i++) { + tex_data[i] = temp_buffer[i]; + } } else if (name == "background") { - return output_buffer_background->getHostPointer(); + std::vector temp_buffer(w * h); + cudaMemcpy(temp_buffer.data(), (void*)state.accum_buffer_b.handle, sizeof(temp_buffer[0]) * temp_buffer.size(), cudaMemcpyDeviceToHost); + for (auto i = 0; i < temp_buffer.size(); i++) { + tex_data[i * 3 + 0] = temp_buffer[i]; + tex_data[i * 3 + 1] = temp_buffer[i]; + tex_data[i * 3 + 2] = temp_buffer[i]; + } } else if (name == "mask") { - return output_buffer_mask->getHostPointer(); + cudaMemcpy(tex_data.data(), (void*)state.accum_buffer_m.handle, sizeof(half) * tex_data.size(), cudaMemcpyDeviceToHost); } else if (name == "color") { - return output_buffer_color->getHostPointer(); + std::vector temp_buffer(w * h * 3); + cudaMemcpy(temp_buffer.data(), (void*)state.accum_buffer_p.handle, sizeof(temp_buffer[0]) * temp_buffer.size(), cudaMemcpyDeviceToHost); + for (auto i = 0; i < temp_buffer.size(); i++) { + tex_data[i] = temp_buffer[i]; + } + } + else { + throw std::runtime_error("invalid optixgetimg_extra name: " + name); } - throw std::runtime_error("invalid optixgetimg_extra name: " + name); + zeno::image_flip_vertical((ushort3*)tex_data.data(), w, h); + return tex_data; } static void save_exr(float3* ptr, int w, int h, std::string path) { std::vector data(w * h); @@ -3868,20 +3869,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_extra2("mask", w, h).data()); } // AOV if (enable_output_aov) { if (enable_output_exr) { zeno::create_directories_when_write_file(exr_path); - SaveMultiLayerEXR( + SaveMultiLayerEXR_half( { - (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_extra3("color", w, h).data(), + optixgetimg_extra3("diffuse", w, h).data(), + optixgetimg_extra3("specular", w, h).data(), + optixgetimg_extra3("transmit", w, h).data(), + optixgetimg_extra3("background", w, h).data(), + optixgetimg_extra3("mask", w, h).data(), }, w, h, @@ -3899,17 +3900,17 @@ void optixrender(int fbo, int samples, bool denoise, bool simpleRender) { } else { 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_extra2("diffuse", w, h).data()); + save_png_color(path + ".aov.specular.png", w, h, optixgetimg_extra2("specular", w, h).data()); + save_png_color(path + ".aov.transmit.png", w, h, optixgetimg_extra2("transmit", w, h).data()); + save_png_data(path + ".aov.background.png", w, h, optixgetimg_extra2("background", w, h).data()); + save_png_data(path + ".aov.mask.png", w, h, optixgetimg_extra2("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_extra2("color", w, h).data(), w, h, exr_path); } else { std::string jpg_native_path = zeno::create_directories_when_write_file(path); @@ -4018,12 +4019,6 @@ void optixDestroy() { occlusion_miss_group .handle=0; output_buffer_o .reset(); - output_buffer_diffuse .reset(); - output_buffer_specular .reset(); - output_buffer_transmit .reset(); - output_buffer_background .reset(); - output_buffer_mask .reset(); - state = {}; isPipelineCreated = false; } diff --git a/zenovis/xinxinoptix/optixPathTracer.h b/zenovis/xinxinoptix/optixPathTracer.h index 7df84d652e..636366da22 100644 --- a/zenovis/xinxinoptix/optixPathTracer.h +++ b/zenovis/xinxinoptix/optixPathTracer.h @@ -154,18 +154,13 @@ struct CameraInfo struct Params { unsigned int subframe_index; - float4* accum_buffer; - float4* accum_buffer_D; - float4* accum_buffer_S; - float4* accum_buffer_T; - float4* accum_buffer_B; + float3* accum_buffer; + float3* accum_buffer_D; + float3* accum_buffer_S; + float3* accum_buffer_T; + ushort1* 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_M; float3* debug_buffer; float3* albedo_buffer;