diff --git a/zenovis/xinxinoptix/PTKernel.cu b/zenovis/xinxinoptix/PTKernel.cu index 2458338ac6..69e570881a 100644 --- a/zenovis/xinxinoptix/PTKernel.cu +++ b/zenovis/xinxinoptix/PTKernel.cu @@ -383,11 +383,11 @@ 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_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 = 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 ); @@ -406,11 +406,11 @@ 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); + 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 ] = accum_color_b; vec3 rgb_mapped = PhysicalCamera(vec3(accum_color), aperture, shutter_speed, iso, midGray, false, false); diff --git a/zenovis/xinxinoptix/optixPathTracer.cpp b/zenovis/xinxinoptix/optixPathTracer.cpp index 9eeb133087..47eca030a6 100644 --- a/zenovis/xinxinoptix/optixPathTracer.cpp +++ b/zenovis/xinxinoptix/optixPathTracer.cpp @@ -570,7 +570,7 @@ static void initLaunchParams( PathTracerState& state ) reinterpret_cast( &state.accum_buffer_p.reset() ), state.params.width * state.params.height * sizeof( float4 ) ) ); - state.params.accum_buffer = (float4*)(CUdeviceptr)state.accum_buffer_p; + state.params.accum_buffer = (float3*)(CUdeviceptr)state.accum_buffer_p; auto& params = state.params; @@ -628,25 +628,25 @@ static void handleResize( sutil::CUDAOutputBuffer& output_buffer, Params // 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_b .reset()), - params.width * params.height * sizeof( float4 ) + params.width * params.height * sizeof( float3 ) ) ); - 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 +660,10 @@ 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.accum_buffer_B = (float3*)(CUdeviceptr)state.accum_buffer_b; state.params.subframe_index = 0; } diff --git a/zenovis/xinxinoptix/optixPathTracer.h b/zenovis/xinxinoptix/optixPathTracer.h index 47386835dd..64a53857c2 100644 --- a/zenovis/xinxinoptix/optixPathTracer.h +++ b/zenovis/xinxinoptix/optixPathTracer.h @@ -154,11 +154,11 @@ 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; + float3* accum_buffer_B; uchar4* frame_buffer; ushort3* frame_buffer_C; ushort3* frame_buffer_D;