Unfortunately adding the relocatble switch and then moving to 5.0.0 neither removed the problem.
The actual problem report follows;
Unknown error (Details: Function “_rtProgramCreateFromPTXFile” caught exception: Compile Error: Unitialized global found: at: [ GlobalValue: @17 = internal addrspace(1) global double, align 8, in module: Canonical__Z19camera_orthographicv from d:\andy.page3\CEM_ROOT\NT_MINGW14\x64\camera.ptx ])
The offending PTX file camera.ptx is included below, if i could determine the name of the global variable then might be sufficient to start understanding and solving the problem. There is no name given, it talks about @17 is that a line number ?
Any help would be much appreciated .
//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-21373419
// Cuda compilation tools, release 8.0, V8.0.55
// Based on LLVM 3.4svn
//
.version 5.0
.target sm_30
.address_size 64
// .globl _Z6cameraN5optix3RayEPN8stingray3rcs10RayDataRCSE
.extern .func (.param .b32 func_retval0) vprintf
(
.param .b64 vprintf_param_0,
.param .b64 vprintf_param_1
)
;
.func (.param .b64 func_retval0) __internal_trig_reduction_slowpathd
(
.param .b64 __internal_trig_reduction_slowpathd_param_0,
.param .b64 __internal_trig_reduction_slowpathd_param_1
)
;
.visible .global .align 4 .f32 scene_epsilon;
.visible .global .align 4 .f32 scene_max;
.visible .global .align 4 .u32 refl;
.visible .global .align 4 .u32 refl_max_depth;
.visible .global .align 4 .u32 refl_max_bounces;
.visible .global .align 4 .u32 refl_iterations;
.visible .global .align 4 .u32 refl_iterations_number;
.visible .global .align 4 .u32 refl_iterations_compact;
.visible .global .align 4 .u32 baverage;
.visible .global .align 4 .u32 idle_count;
.visible .global .align 4 .u32 NumLaunchRays;
.visible .global .align 4 .u32 front_and_back;
.visible .global .align 4 .u32 shading;
.visible .global .align 4 .u32 front_stop_recover;
.visible .global .align 4 .u32 color_refl;
.visible .global .align 4 .u32 color_refl_depth;
.visible .global .align 4 .u32 color_front_stop;
.visible .global .align 4 .b8 color_front_stop_color[12];
.visible .global .align 4 .b8 ambient_light_color[12];
.visible .global .align 4 .b8 exception_color[12];
.visible .global .align 4 .b8 bg_color[12];
.visible .global .align 4 .b8 ray_init_color[12];
.visible .global .align 4 .u32 cell_black_is_black;
.visible .global .align 16 .b8 perspective_half_fov[16];
.visible .global .align 1 .b8 count_buffer[1];
.visible .global .align 8 .b8 eye[24];
.visible .global .align 8 .b8 U[24];
.visible .global .align 8 .b8 V[24];
.visible .global .align 8 .b8 W_negative[24];
.visible .global .align 8 .b8 W_negative_launch_dir[24];
.visible .global .align 8 .b8 top_object[4];
.visible .global .align 4 .b8 top_shadower[4];
.visible .global .align 8 .f64 frequency_value;
.visible .global .align 4 .u32 measurement_plane;
.visible .global .align 4 .u32 calc_rcs_tube;
.visible .global .align 4 .u32 calc_rcs;
.visible .global .align 8 .f64 diamond_angle;
.visible .global .align 8 .f64 cone_angle;
.visible .global .align 4 .u32 save_ray_tubes;
.visible .global .align 4 .u32 ab3_mode;
.visible .global .align 4 .u32 btube_tri;
.visible .global .align 4 .u32 rcs_refl_lower;
.visible .global .align 4 .u32 rcs_refl_upper;
.visible .global .align 8 .f64 range_min;
.visible .global .align 8 .f64 range_max;
.visible .global .align 8 .f64 internal_pq_zero;
.visible .global .align 4 .u32 bmustbe_consistent;
.visible .global .align 8 .f64 internal_small_area;
.visible .global .align 8 .f64 internal_large_area;
.visible .global .align 4 .u32 buse_divergence_area;
.visible .global .align 4 .u32 bshape_expand;
.visible .global .align 4 .u32 bshape_func;
.visible .global .align 8 .f64 internal_area_ratio;
.visible .global .align 8 .f64 alpha;
.visible .global .align 8 .f64 power_ratio_limit;
.visible .global .align 4 .u32 bopen_geometry;
.visible .global .align 4 .u32 bram_filled;
.visible .global .align 8 .b8 up_vector[24];
.visible .global .align 8 .f64 three_lambda;
.visible .global .align 8 .f64 nep_max_area;
.visible .global .align 1 .u8 bswrti;
.visible .global .align 1 .u8 bpowa_correct;
.visible .global .align 4 .u32 bfields_checkup;
.visible .global .align 8 .f64 lengthtest;
.visible .global .align 8 .f64 cosangletest;
.visible .global .align 8 .f64 pointtest;
.visible .global .align 4 .u32 bbrti;
.visible .global .align 1 .b8 lights_buffer[1];
.visible .global .align 1 .b8 output_buffer[1];
.visible .global .align 1 .b8 ray_curvature_buffer[1];
.visible .global .align 1 .b8 h_e_field_buffer[1];
.visible .global .align 1 .b8 v_e_field_buffer[1];
.visible .global .align 1 .b8 df_mul_buffer[1];
.visible .global .align 1 .b8 measured_buffer[1];
.visible .global .align 1 .b8 thetaphi_buffer[1];
.visible .global .align 1 .b8 xy_buffer[1];
.visible .global .align 1 .b8 last_out_dir_buffer[1];
.visible .global .align 1 .b8 front_stop_buffer[1];
.visible .global .align 1 .b8 exception_buffer[1];
.visible .global .align 1 .b8 start_dir_buffer[1];
.visible .global .align 1 .b8 raybounces_buffer[1];
.visible .global .align 1 .b8 ray_tube_int_buffer[1];
.visible .global .align 1 .b8 raytube_endpoint_buffer[1];
.visible .global .align 1 .b8 raytube_bounces_buffer[1];
.visible .global .align 1 .b8 dist_tot_buffer[1];
.visible .global .align 1 .b8 gamma_buffer[1];
.visible .global .align 1 .b8 launchpoint_buffer[1];
.visible .global .align 1 .b8 raytube_launchpoint_buffer[1];
.visible .global .align 8 .b8 eye_m[24];
.visible .global .align 8 .b8 U_m[24];
.visible .global .align 8 .b8 V_m[24];
.visible .global .align 8 .b8 minpt[24];
.visible .global .align 8 .b8 maxpt[24];
.visible .global .align 8 .b8 launch_point_measurement[24];
.visible .global .align 8 .f64 raytube_area;
.visible .global .align 16 .b8 bistaticAngle[16];
.visible .global .align 1 .b8 atheta_h_contrib_buffer[1];
.visible .global .align 1 .b8 aphi_h_contrib_buffer[1];
.visible .global .align 1 .b8 atheta_v_contrib_buffer[1];
.visible .global .align 1 .b8 aphi_v_contrib_buffer[1];
.visible .global .align 1 .b8 shape_function_buffer[1];
.visible .global .align 1 .b8 area_buffer[1];
.visible .global .align 1 .b8 alt_raybounces_lookup_buffer[1];
.visible .global .align 1 .b8 alt_raybounces_buffer[1];
.visible .global .align 1 .b8 alt_count_buffer[1];
.visible .global .align 1 .b8 ram_definition_buffer[1];
.visible .global .align 1 .b8 cubic_segment_buffer[1];
.visible .global .align 1 .b8 cell_id_to_ram_id_buffer[1];
.visible .global .align 1 .b8 counterstate_buffer[1];
.visible .global .align 1 .b8 counterstate_field_buffer[1];
.visible .global .align 8 .b8 k0_unit[24];
.visible .global .align 8 .b8 xray_first[24];
.visible .global .align 8 .b8 yray_first[24];
.visible .global .align 1 .b8 powa_buffer[1];
.visible .global .align 8 .b8 launch_index[8];
.visible .global .align 8 .u64 _ZN21rti_internal_register20reg_bitness_detectorE;
.visible .global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail0E;
.visible .global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail1E;
.visible .global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail2E;
.visible .global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail3E;
.visible .global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail4E;
.visible .global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail5E;
.visible .global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail6E;
.visible .global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail7E;
.visible .global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail8E;
.visible .global .align 8 .u64 _ZN21rti_internal_register24reg_exception_64_detail9E;
.visible .global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail0E;
.visible .global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail1E;
.visible .global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail2E;
.visible .global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail3E;
.visible .global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail4E;
.visible .global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail5E;
.visible .global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail6E;
.visible .global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail7E;
.visible .global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail8E;
.visible .global .align 4 .u32 _ZN21rti_internal_register21reg_exception_detail9E;
.visible .global .align 4 .u32 _ZN21rti_internal_register14reg_rayIndex_xE;
.visible .global .align 4 .u32 _ZN21rti_internal_register14reg_rayIndex_yE;
.visible .global .align 4 .u32 _ZN21rti_internal_register14reg_rayIndex_zE;
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo13scene_epsilonE[8] = {82, 97, 121, 0, 4, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo9scene_maxE[8] = {82, 97, 121, 0, 4, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo4reflE[8] = {82, 97, 121, 0, 4, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo14refl_max_depthE[8] = {82, 97, 121, 0, 4, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo16refl_max_bouncesE[8] = {82, 97, 121, 0, 4, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo15refl_iterationsE[8] = {82, 97, 121, 0, 4, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo22refl_iterations_numberE[8] = {82, 97, 121, 0, 4, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo23refl_iterations_compactE[8] = {82, 97, 121, 0, 4, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo8baverageE[8] = {82, 97, 121, 0, 4, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo10idle_countE[8] = {82, 97, 121, 0, 4, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo13NumLaunchRaysE[8] = {82, 97, 121, 0, 4, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo14front_and_backE[8] = {82, 97, 121, 0, 4, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo7shadingE[8] = {82, 97, 121, 0, 4, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo18front_stop_recoverE[8] = {82, 97, 121, 0, 4, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo10color_reflE[8] = {82, 97, 121, 0, 4, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo16color_refl_depthE[8] = {82, 97, 121, 0, 4, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo16color_front_stopE[8] = {82, 97, 121, 0, 4, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo22color_front_stop_colorE[8] = {82, 97, 121, 0, 12, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo19ambient_light_colorE[8] = {82, 97, 121, 0, 12, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo15exception_colorE[8] = {82, 97, 121, 0, 12, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo8bg_colorE[8] = {82, 97, 121, 0, 12, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo14ray_init_colorE[8] = {82, 97, 121, 0, 12, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo19cell_black_is_blackE[8] = {82, 97, 121, 0, 4, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo20perspective_half_fovE[8] = {82, 97, 121, 0, 16, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo3eyeE[8] = {82, 97, 121, 0, 24, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo1UE[8] = {82, 97, 121, 0, 24, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo1VE[8] = {82, 97, 121, 0, 24, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo10W_negativeE[8] = {82, 97, 121, 0, 24, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo21W_negative_launch_dirE[8] = {82, 97, 121, 0, 24, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo10top_objectE[8] = {82, 97, 121, 0, 4, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo12top_shadowerE[8] = {82, 97, 121, 0, 4, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo15frequency_valueE[8] = {82, 97, 121, 0, 8, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo17measurement_planeE[8] = {82, 97, 121, 0, 4, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo13calc_rcs_tubeE[8] = {82, 97, 121, 0, 4, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo8calc_rcsE[8] = {82, 97, 121, 0, 4, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo13diamond_angleE[8] = {82, 97, 121, 0, 8, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo10cone_angleE[8] = {82, 97, 121, 0, 8, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo14save_ray_tubesE[8] = {82, 97, 121, 0, 4, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo8ab3_modeE[8] = {82, 97, 121, 0, 4, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo9btube_triE[8] = {82, 97, 121, 0, 4, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo14rcs_refl_lowerE[8] = {82, 97, 121, 0, 4, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo14rcs_refl_upperE[8] = {82, 97, 121, 0, 4, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo9range_minE[8] = {82, 97, 121, 0, 8, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo9range_maxE[8] = {82, 97, 121, 0, 8, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo16internal_pq_zeroE[8] = {82, 97, 121, 0, 8, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo18bmustbe_consistentE[8] = {82, 97, 121, 0, 4, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo19internal_small_areaE[8] = {82, 97, 121, 0, 8, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo19internal_large_areaE[8] = {82, 97, 121, 0, 8, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo20buse_divergence_areaE[8] = {82, 97, 121, 0, 4, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo13bshape_expandE[8] = {82, 97, 121, 0, 4, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo11bshape_funcE[8] = {82, 97, 121, 0, 4, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo19internal_area_ratioE[8] = {82, 97, 121, 0, 8, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo5alphaE[8] = {82, 97, 121, 0, 8, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo17power_ratio_limitE[8] = {82, 97, 121, 0, 8, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo14bopen_geometryE[8] = {82, 97, 121, 0, 4, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo11bram_filledE[8] = {82, 97, 121, 0, 4, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo9up_vectorE[8] = {82, 97, 121, 0, 24, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo12three_lambdaE[8] = {82, 97, 121, 0, 8, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo12nep_max_areaE[8] = {82, 97, 121, 0, 8, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo6bswrtiE[8] = {82, 97, 121, 0, 1, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo13bpowa_correctE[8] = {82, 97, 121, 0, 1, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo15bfields_checkupE[8] = {82, 97, 121, 0, 4, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo10lengthtestE[8] = {82, 97, 121, 0, 8, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo12cosangletestE[8] = {82, 97, 121, 0, 8, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo9pointtestE[8] = {82, 97, 121, 0, 8, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo5bbrtiE[8] = {82, 97, 121, 0, 4, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo5eye_mE[8] = {82, 97, 121, 0, 24, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo3U_mE[8] = {82, 97, 121, 0, 24, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo3V_mE[8] = {82, 97, 121, 0, 24, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo5minptE[8] = {82, 97, 121, 0, 24, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo5maxptE[8] = {82, 97, 121, 0, 24, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo24launch_point_measurementE[8] = {82, 97, 121, 0, 24, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo12raytube_areaE[8] = {82, 97, 121, 0, 8, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo13bistaticAngleE[8] = {82, 97, 121, 0, 16, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo7k0_unitE[8] = {82, 97, 121, 0, 24, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo10xray_firstE[8] = {82, 97, 121, 0, 24, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo10yray_firstE[8] = {82, 97, 121, 0, 24, 0, 0, 0};
.visible .global .align 4 .b8 _ZN21rti_internal_typeinfo12launch_indexE[8] = {82, 97, 121, 0, 8, 0, 0, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename13scene_epsilonE[6] = {102, 108, 111, 97, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename9scene_maxE[6] = {102, 108, 111, 97, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename4reflE[4] = {105, 110, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename14refl_max_depthE[4] = {105, 110, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename16refl_max_bouncesE[4] = {105, 110, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename15refl_iterationsE[4] = {105, 110, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename22refl_iterations_numberE[4] = {105, 110, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename23refl_iterations_compactE[4] = {105, 110, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename8baverageE[4] = {105, 110, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename10idle_countE[4] = {105, 110, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename13NumLaunchRaysE[4] = {105, 110, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename14front_and_backE[4] = {105, 110, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename7shadingE[4] = {105, 110, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename18front_stop_recoverE[4] = {105, 110, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename10color_reflE[4] = {105, 110, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename16color_refl_depthE[4] = {105, 110, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename16color_front_stopE[4] = {105, 110, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename22color_front_stop_colorE[7] = {102, 108, 111, 97, 116, 51, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename19ambient_light_colorE[7] = {102, 108, 111, 97, 116, 51, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename15exception_colorE[7] = {102, 108, 111, 97, 116, 51, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename8bg_colorE[7] = {102, 108, 111, 97, 116, 51, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename14ray_init_colorE[7] = {102, 108, 111, 97, 116, 51, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename19cell_black_is_blackE[4] = {105, 110, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename20perspective_half_fovE[18] = {115, 116, 105, 110, 103, 114, 97, 121, 58, 58, 116, 102, 108, 111, 97, 116, 50, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename3eyeE[18] = {115, 116, 105, 110, 103, 114, 97, 121, 58, 58, 116, 102, 108, 111, 97, 116, 51, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename1UE[18] = {115, 116, 105, 110, 103, 114, 97, 121, 58, 58, 116, 102, 108, 111, 97, 116, 51, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename1VE[18] = {115, 116, 105, 110, 103, 114, 97, 121, 58, 58, 116, 102, 108, 111, 97, 116, 51, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename10W_negativeE[18] = {115, 116, 105, 110, 103, 114, 97, 121, 58, 58, 116, 102, 108, 111, 97, 116, 51, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename21W_negative_launch_dirE[18] = {115, 116, 105, 110, 103, 114, 97, 121, 58, 58, 116, 102, 108, 111, 97, 116, 51, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename10top_objectE[9] = {114, 116, 79, 98, 106, 101, 99, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename12top_shadowerE[9] = {114, 116, 79, 98, 106, 101, 99, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename15frequency_valueE[17] = {115, 116, 105, 110, 103, 114, 97, 121, 58, 58, 116, 102, 108, 111, 97, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename17measurement_planeE[4] = {105, 110, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename13calc_rcs_tubeE[4] = {105, 110, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename8calc_rcsE[4] = {105, 110, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename13diamond_angleE[17] = {115, 116, 105, 110, 103, 114, 97, 121, 58, 58, 116, 102, 108, 111, 97, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename10cone_angleE[17] = {115, 116, 105, 110, 103, 114, 97, 121, 58, 58, 116, 102, 108, 111, 97, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename14save_ray_tubesE[4] = {105, 110, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename8ab3_modeE[4] = {105, 110, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename9btube_triE[4] = {105, 110, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename14rcs_refl_lowerE[4] = {105, 110, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename14rcs_refl_upperE[4] = {105, 110, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename9range_minE[17] = {115, 116, 105, 110, 103, 114, 97, 121, 58, 58, 116, 102, 108, 111, 97, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename9range_maxE[17] = {115, 116, 105, 110, 103, 114, 97, 121, 58, 58, 116, 102, 108, 111, 97, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename16internal_pq_zeroE[18] = {115, 116, 105, 110, 103, 114, 97, 121, 58, 58, 116, 100, 111, 117, 98, 108, 101, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename18bmustbe_consistentE[4] = {105, 110, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename19internal_small_areaE[18] = {115, 116, 105, 110, 103, 114, 97, 121, 58, 58, 116, 100, 111, 117, 98, 108, 101, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename19internal_large_areaE[18] = {115, 116, 105, 110, 103, 114, 97, 121, 58, 58, 116, 100, 111, 117, 98, 108, 101, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename20buse_divergence_areaE[4] = {105, 110, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename13bshape_expandE[4] = {105, 110, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename11bshape_funcE[4] = {105, 110, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename19internal_area_ratioE[18] = {115, 116, 105, 110, 103, 114, 97, 121, 58, 58, 116, 100, 111, 117, 98, 108, 101, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename5alphaE[17] = {115, 116, 105, 110, 103, 114, 97, 121, 58, 58, 116, 102, 108, 111, 97, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename17power_ratio_limitE[18] = {115, 116, 105, 110, 103, 114, 97, 121, 58, 58, 116, 100, 111, 117, 98, 108, 101, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename14bopen_geometryE[4] = {105, 110, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename11bram_filledE[4] = {105, 110, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename9up_vectorE[18] = {115, 116, 105, 110, 103, 114, 97, 121, 58, 58, 116, 102, 108, 111, 97, 116, 51, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename12three_lambdaE[17] = {115, 116, 105, 110, 103, 114, 97, 121, 58, 58, 116, 102, 108, 111, 97, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename12nep_max_areaE[17] = {115, 116, 105, 110, 103, 114, 97, 121, 58, 58, 116, 102, 108, 111, 97, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename6bswrtiE[5] = {98, 111, 111, 108, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename13bpowa_correctE[5] = {98, 111, 111, 108, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename15bfields_checkupE[4] = {105, 110, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename10lengthtestE[17] = {115, 116, 105, 110, 103, 114, 97, 121, 58, 58, 116, 102, 108, 111, 97, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename12cosangletestE[17] = {115, 116, 105, 110, 103, 114, 97, 121, 58, 58, 116, 102, 108, 111, 97, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename9pointtestE[17] = {115, 116, 105, 110, 103, 114, 97, 121, 58, 58, 116, 102, 108, 111, 97, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename5bbrtiE[4] = {105, 110, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename5eye_mE[18] = {115, 116, 105, 110, 103, 114, 97, 121, 58, 58, 116, 102, 108, 111, 97, 116, 51, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename3U_mE[18] = {115, 116, 105, 110, 103, 114, 97, 121, 58, 58, 116, 102, 108, 111, 97, 116, 51, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename3V_mE[18] = {115, 116, 105, 110, 103, 114, 97, 121, 58, 58, 116, 102, 108, 111, 97, 116, 51, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename5minptE[18] = {115, 116, 105, 110, 103, 114, 97, 121, 58, 58, 116, 102, 108, 111, 97, 116, 51, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename5maxptE[18] = {115, 116, 105, 110, 103, 114, 97, 121, 58, 58, 116, 102, 108, 111, 97, 116, 51, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename24launch_point_measurementE[18] = {115, 116, 105, 110, 103, 114, 97, 121, 58, 58, 116, 102, 108, 111, 97, 116, 51, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename12raytube_areaE[17] = {115, 116, 105, 110, 103, 114, 97, 121, 58, 58, 116, 102, 108, 111, 97, 116, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename13bistaticAngleE[18] = {115, 116, 105, 110, 103, 114, 97, 121, 58, 58, 116, 102, 108, 111, 97, 116, 50, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename7k0_unitE[18] = {115, 116, 105, 110, 103, 114, 97, 121, 58, 58, 116, 102, 108, 111, 97, 116, 51, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename10xray_firstE[18] = {115, 116, 105, 110, 103, 114, 97, 121, 58, 58, 116, 102, 108, 111, 97, 116, 51, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename10yray_firstE[18] = {115, 116, 105, 110, 103, 114, 97, 121, 58, 58, 116, 102, 108, 111, 97, 116, 51, 0};
.visible .global .align 1 .b8 _ZN21rti_internal_typename12launch_indexE[6] = {117, 105, 110, 116, 50, 0};
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum13scene_epsilonE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum9scene_maxE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum4reflE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum14refl_max_depthE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum16refl_max_bouncesE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum15refl_iterationsE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum22refl_iterations_numberE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum23refl_iterations_compactE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum8baverageE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum10idle_countE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum13NumLaunchRaysE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum14front_and_backE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum7shadingE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum18front_stop_recoverE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum10color_reflE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum16color_refl_depthE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum16color_front_stopE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum22color_front_stop_colorE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum19ambient_light_colorE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum15exception_colorE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum8bg_colorE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum14ray_init_colorE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum19cell_black_is_blackE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum20perspective_half_fovE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum3eyeE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum1UE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum1VE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum10W_negativeE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum21W_negative_launch_dirE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum10top_objectE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum12top_shadowerE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum15frequency_valueE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum17measurement_planeE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum13calc_rcs_tubeE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum8calc_rcsE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum13diamond_angleE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum10cone_angleE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum14save_ray_tubesE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum8ab3_modeE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum9btube_triE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum14rcs_refl_lowerE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum14rcs_refl_upperE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum9range_minE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum9range_maxE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum16internal_pq_zeroE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum18bmustbe_consistentE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum19internal_small_areaE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum19internal_large_areaE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum20buse_divergence_areaE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum13bshape_expandE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum11bshape_funcE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum19internal_area_ratioE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum5alphaE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum17power_ratio_limitE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum14bopen_geometryE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum11bram_filledE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum9up_vectorE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum12three_lambdaE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum12nep_max_areaE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum6bswrtiE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum13bpowa_correctE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum15bfields_checkupE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum10lengthtestE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum12cosangletestE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum9pointtestE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum5bbrtiE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum5eye_mE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum3U_mE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum3V_mE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum5minptE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum5maxptE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum24launch_point_measurementE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum12raytube_areaE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum13bistaticAngleE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum7k0_unitE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum10xray_firstE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum10yray_firstE = 4919;
.visible .global .align 4 .u32 _ZN21rti_internal_typeenum12launch_indexE = 4919;
.visible .global .align 1 .b8 _ZN21rti_internal_semantic13scene_epsilonE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic9scene_maxE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic4reflE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic14refl_max_depthE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic16refl_max_bouncesE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic15refl_iterationsE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic22refl_iterations_numberE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic23refl_iterations_compactE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic8baverageE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic10idle_countE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic13NumLaunchRaysE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic14front_and_backE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic7shadingE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic18front_stop_recoverE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic10color_reflE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic16color_refl_depthE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic16color_front_stopE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic22color_front_stop_colorE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic19ambient_light_colorE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic15exception_colorE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic8bg_colorE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic14ray_init_colorE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic19cell_black_is_blackE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic20perspective_half_fovE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic3eyeE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic1UE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic1VE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic10W_negativeE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic21W_negative_launch_dirE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic10top_objectE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic12top_shadowerE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic15frequency_valueE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic17measurement_planeE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic13calc_rcs_tubeE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic8calc_rcsE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic13diamond_angleE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic10cone_angleE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic14save_ray_tubesE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic8ab3_modeE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic9btube_triE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic14rcs_refl_lowerE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic14rcs_refl_upperE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic9range_minE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic9range_maxE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic16internal_pq_zeroE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic18bmustbe_consistentE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic19internal_small_areaE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic19internal_large_areaE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic20buse_divergence_areaE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic13bshape_expandE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic11bshape_funcE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic19internal_area_ratioE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic5alphaE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic17power_ratio_limitE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic14bopen_geometryE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic11bram_filledE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic9up_vectorE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic12three_lambdaE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic12nep_max_areaE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic6bswrtiE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic13bpowa_correctE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic15bfields_checkupE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic10lengthtestE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic12cosangletestE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic9pointtestE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic5bbrtiE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic5eye_mE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic3U_mE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic3V_mE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic5minptE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic5maxptE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic24launch_point_measurementE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic12raytube_areaE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic13bistaticAngleE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic7k0_unitE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic10xray_firstE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic10yray_firstE[1];
.visible .global .align 1 .b8 _ZN21rti_internal_semantic12launch_indexE[14] = {114, 116, 76, 97, 117, 110, 99, 104, 73, 110, 100, 101, 120, 0};
.visible .global .align 1 .b8 _ZN23rti_internal_annotation13scene_epsilonE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation9scene_maxE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation4reflE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation14refl_max_depthE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation16refl_max_bouncesE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation15refl_iterationsE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation22refl_iterations_numberE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation23refl_iterations_compactE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation8baverageE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation10idle_countE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation13NumLaunchRaysE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation14front_and_backE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation7shadingE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation18front_stop_recoverE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation10color_reflE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation16color_refl_depthE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation16color_front_stopE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation22color_front_stop_colorE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation19ambient_light_colorE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation15exception_colorE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation8bg_colorE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation14ray_init_colorE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation19cell_black_is_blackE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation20perspective_half_fovE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation3eyeE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation1UE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation1VE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation10W_negativeE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation21W_negative_launch_dirE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation10top_objectE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation12top_shadowerE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation15frequency_valueE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation17measurement_planeE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation13calc_rcs_tubeE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation8calc_rcsE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation13diamond_angleE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation10cone_angleE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation14save_ray_tubesE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation8ab3_modeE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation9btube_triE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation14rcs_refl_lowerE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation14rcs_refl_upperE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation9range_minE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation9range_maxE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation16internal_pq_zeroE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation18bmustbe_consistentE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation19internal_small_areaE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation19internal_large_areaE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation20buse_divergence_areaE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation13bshape_expandE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation11bshape_funcE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation19internal_area_ratioE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation5alphaE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation17power_ratio_limitE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation14bopen_geometryE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation11bram_filledE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation9up_vectorE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation12three_lambdaE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation12nep_max_areaE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation6bswrtiE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation13bpowa_correctE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation15bfields_checkupE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation10lengthtestE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation12cosangletestE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation9pointtestE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation5bbrtiE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation5eye_mE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation3U_mE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation3V_mE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation5minptE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation5maxptE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation24launch_point_measurementE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation12raytube_areaE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation13bistaticAngleE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation7k0_unitE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation10xray_firstE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation10yray_firstE[1];
.visible .global .align 1 .b8 _ZN23rti_internal_annotation12launch_indexE[1];
.global .align 1 .b8 $str[77] = {99, 97, 109, 101, 114, 97, 95, 101, 120, 99, 101, 112, 116, 105, 111, 110, 32, 58, 32, 67, 97, 117, 103, 104, 116, 32, 101, 120, 99, 101, 112, 116, 105, 111, 110, 32, 37, 117, 32, 40, 32, 48, 120, 37, 88, 32, 41, 32, 97, 116, 32, 108, 97, 117, 110, 99, 104, 32, 105, 110, 100, 101, 120, 32, 40, 32, 37, 100, 32, 44, 32, 37, 100, 32, 41, 10, 0};
.const .align 8 .b8 __cudart_i2opi_d[144] = {8, 93, 141, 31, 177, 95, 251, 107, 234, 146, 82, 138, 247, 57, 7, 61, 123, 241, 229, 235, 199, 186, 39, 117, 45, 234, 95, 158, 102, 63, 70, 79, 183, 9, 203, 39, 207, 126, 54, 109, 31, 109, 10, 90, 139, 17, 47, 239, 15, 152, 5, 222, 255, 151, 248, 31, 59, 40, 249, 189, 139, 95, 132, 156, 244, 57, 83, 131, 57, 214, 145, 57, 65, 126, 95, 180, 38, 112, 156, 233, 132, 68, 187, 46, 245, 53, 130, 232, 62, 167, 41, 177, 28, 235, 29, 254, 28, 146, 209, 9, 234, 46, 73, 6, 224, 210, 77, 66, 58, 110, 36, 183, 97, 197, 187, 222, 171, 99, 81, 254, 65, 144, 67, 60, 153, 149, 98, 219, 192, 221, 52, 245, 209, 87, 39, 252, 41, 21, 68, 78, 110, 131, 249, 162};
.const .align 8 .b8 __cudart_sin_cos_coeffs[128] = {186, 94, 120, 249, 101, 219, 229, 61, 70, 210, 176, 44, 241, 229, 90, 190, 146, 227, 172, 105, 227, 29, 199, 62, 161, 98, 219, 25, 160, 1, 42, 191, 24, 8, 17, 17, 17, 17, 129, 63, 84, 85, 85, 85, 85, 85, 197, 191, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 100, 129, 253, 32, 131, 255, 168, 189, 40, 133, 239, 193, 167, 238, 33, 62, 217, 230, 6, 142, 79, 126, 146, 190, 233, 188, 221, 25, 160, 1, 250, 62, 71, 93, 193, 22, 108, 193, 86, 191, 81, 85, 85, 85, 85, 85, 165, 63, 0, 0, 0, 0, 0, 0, 224, 191, 0, 0, 0, 0, 0, 0, 240, 63};
.visible .func _Z6cameraN5optix3RayEPN8stingray3rcs10RayDataRCSE(
.param .b64 _Z6cameraN5optix3RayEPN8stingray3rcs10RayDataRCSE_param_0,
.param .b64 _Z6cameraN5optix3RayEPN8stingray3rcs10RayDataRCSE_param_1
)
{
.local .align 8 .b8 __local_depot0[32];
.reg .b64 %SP;
.reg .b64 %SPL;
.reg .pred %p<62>;
.reg .b16 %rs<90>;
.reg .f32 %f<125>;
.reg .b32 %r<170>;
.reg .f64 %fd<27>;
.reg .b64 %rd<159>;
mov.u64 %rd158, __local_depot0;
cvta.local.u64 %SP, %rd158;
ld.param.u64 %rd11, [_Z6cameraN5optix3RayEPN8stingray3rcs10RayDataRCSE_param_0];
ld.param.u64 %rd12, [_Z6cameraN5optix3RayEPN8stingray3rcs10RayDataRCSE_param_1];
add.u64 %rd18, %SP, 0;
cvta.to.local.u64 %rd5, %rd18;
mov.u64 %rd19, output_buffer;
cvta.global.u64 %rd17, %rd19;
mov.u32 %r23, 2;
mov.u32 %r24, 4;
// inline asm
call (%rd13, %rd14, %rd15, %rd16), _rt_buffer_get_size_64, (%rd17, %r23, %r24);
// inline asm
setp.eq.s64 %p3, %rd12, 0;
@%p3 bra BB0_2;
ld.v2.u32 {%r25, %r26}, [%rd12];
ld.v2.u32 {%r27, %r28}, [%rd12+8];
ld.v2.u32 {%r29, %r30}, [%rd12+16];
ld.v4.u8 {%rs19, %rs20, %rs21, %rs22}, [%rd12+24];
ld.v4.u8 {%rs23, %rs24, %rs25, %rs26}, [%rd12+28];
st.local.v4.u8 [%rd5+28], {%rs23, %rs24, %rs25, %rs26};
st.local.v4.u8 [%rd5+24], {%rs19, %rs20, %rs21, %rs22};
st.local.v2.u32 [%rd5+16], {%r29, %r30};
st.local.v2.u32 [%rd5+8], {%r27, %r28};
st.local.v2.u32 [%rd5], {%r25, %r26};
bra.uni BB0_9;
BB0_2:
mov.u32 %r37, 1;
mov.u32 %r38, 0;
st.local.v2.u32 [%rd5], {%r38, %r37};
st.local.v2.u32 [%rd5+8], {%r38, %r38};
ld.global.v2.u32 {%r39, %r40}, [launch_index];
st.local.v2.u32 [%rd5+16], {%r39, %r40};
ld.global.f32 %f1, [ray_init_color+8];
mov.f32 %f58, 0f00000000;
setp.le.ftz.f32 %p4, %f1, 0f00000000;
mov.f32 %f106, %f58;
@%p4 bra BB0_4;
setp.ltu.ftz.f32 %p5, %f1, 0f3F800000;
mul.ftz.f32 %f59, %f1, 0f437FFD71;
selp.f32 %f2, %f59, 0f437FFD71, %p5;
mov.f32 %f106, %f2;
BB0_4:
mov.f32 %f3, %f106;
cvt.rzi.ftz.u32.f32 %r1, %f3;
ld.global.f32 %f4, [ray_init_color+4];
setp.le.ftz.f32 %p6, %f4, 0f00000000;
mov.f32 %f105, %f58;
@%p6 bra BB0_6;
setp.ltu.ftz.f32 %p7, %f4, 0f3F800000;
mul.ftz.f32 %f61, %f4, 0f437FFD71;
selp.f32 %f105, %f61, 0f437FFD71, %p7;
BB0_6:
cvt.rzi.ftz.u32.f32 %r2, %f105;
ld.global.f32 %f7, [ray_init_color];
setp.le.ftz.f32 %p8, %f7, 0f00000000;
mov.f32 %f104, %f58;
@%p8 bra BB0_8;
setp.ltu.ftz.f32 %p9, %f7, 0f3F800000;
mul.ftz.f32 %f63, %f7, 0f437FFD71;
selp.f32 %f104, %f63, 0f437FFD71, %p9;
BB0_8:
cvt.rzi.ftz.u32.f32 %r43, %f104;
cvt.u16.u32 %rs35, %r2;
cvt.u16.u32 %rs36, %r1;
cvt.u16.u32 %rs37, %r43;
mov.u16 %rs38, 255;
st.local.v4.u8 [%rd5+24], {%rs36, %rs35, %rs37, %rs38};
BB0_9:
ld.global.u32 %r168, [refl_max_depth];
setp.lt.s32 %p10, %r168, 1;
@%p10 bra BB0_21;
mov.u32 %r166, 0;
BB0_11:
mov.u32 %r4, %r166;
setp.eq.s32 %p11, %r4, 0;
@%p11 bra BB0_13;
bra.uni BB0_12;
BB0_13:
ld.f32 %f107, [%rd11];
ld.f32 %f108, [%rd11+4];
ld.f32 %f109, [%rd11+8];
ld.f32 %f110, [%rd11+12];
ld.f32 %f111, [%rd11+16];
ld.f32 %f112, [%rd11+20];
ld.u32 %r167, [%rd11+24];
ld.f32 %f113, [%rd11+28];
ld.f32 %f114, [%rd11+32];
bra.uni BB0_14;
BB0_12:
ld.global.f32 %f113, [scene_epsilon];
mov.u32 %r167, 0;
ld.global.f32 %f114, [scene_max];
BB0_14:
ld.global.u32 %r46, [top_object];
mov.u32 %r48, 32;
// inline asm
call _rt_trace_64, (%r46, %f107, %f108, %f109, %f110, %f111, %f112, %r167, %f113, %f114, %rd18, %r48);
// inline asm
setp.ne.s32 %p12, %r4, 0;
@%p12 bra BB0_16;
ld.local.v4.u8 {%rs55, %rs56, %rs57, %rs58}, [%rd5+24];
mov.u16 %rs88, %rs58;
mov.u16 %rs87, %rs57;
mov.u16 %rs86, %rs56;
mov.u16 %rs85, %rs55;
BB0_16:
ld.global.v2.u32 {%r51, %r52}, [launch_index];
cvt.u64.u32 %rd23, %r51;
cvt.u64.u32 %rd24, %r52;
mov.u64 %rd27, exception_buffer;
cvta.global.u64 %rd22, %rd27;
mov.u64 %rd26, 0;
// inline asm
call (%rd21), _rt_buffer_get_64, (%rd22, %r23, %r24, %rd23, %rd24, %rd26, %rd26);
// inline asm
ld.u32 %r55, [%rd21];
setp.ne.s32 %p13, %r55, 0;
@%p13 bra BB0_20;
ld.local.u32 %r166, [%rd5];
add.s32 %r56, %r4, 1;
setp.ne.s32 %p14, %r166, %r56;
@%p14 bra BB0_20;
ld.local.u32 %r57, [%rd5+12];
setp.ne.s32 %p15, %r57, 0;
@%p15 bra BB0_20;
add.s32 %r68, %r166, -1;
cvt.u64.u32 %rd55, %r68;
// inline asm
call (%rd28, %rd29, %rd30, %rd31), _rt_buffer_get_size_64, (%rd17, %r23, %r24);
// inline asm
// inline asm
call (%rd33, %rd34, %rd35, %rd36), _rt_buffer_get_size_64, (%rd17, %r23, %r24);
// inline asm
mul.lo.s64 %rd57, %rd55, %rd29;
mul.lo.s64 %rd58, %rd57, %rd33;
ld.global.u32 %rd59, [launch_index+4];
// inline asm
call (%rd38, %rd39, %rd40, %rd41), _rt_buffer_get_size_64, (%rd17, %r23, %r24);
// inline asm
mul.lo.s64 %rd60, %rd38, %rd59;
add.s64 %rd61, %rd60, %rd58;
ld.global.u32 %rd62, [launch_index];
add.s64 %rd63, %rd61, %rd62;
and.b64 %rd45, %rd63, 4294967295;
mov.u64 %rd64, raybounces_buffer;
cvta.global.u64 %rd44, %rd64;
mov.u32 %r64, 1;
mov.u32 %r65, 96;
// inline asm
call (%rd43), _rt_buffer_get_64, (%rd44, %r64, %r65, %rd45, %rd26, %rd26, %rd26);
// inline asm
ld.f64 %fd13, [%rd43+16];
ld.f64 %fd14, [%rd43+8];
ld.f64 %fd15, [%rd43];
ld.global.v2.u32 {%r69, %r70}, [launch_index];
cvt.u64.u32 %rd51, %r69;
cvt.u64.u32 %rd52, %r70;
mov.u64 %rd65, last_out_dir_buffer;
cvta.global.u64 %rd50, %rd65;
mov.u32 %r67, 24;
// inline asm
call (%rd49), _rt_buffer_get_64, (%rd50, %r23, %r67, %rd51, %rd52, %rd26, %rd26);
// inline asm
cvt.rn.ftz.f32.f64 %f107, %fd15;
cvt.rn.ftz.f32.f64 %f108, %fd14;
cvt.rn.ftz.f32.f64 %f109, %fd13;
ld.f64 %fd16, [%rd49+16];
ld.f64 %fd17, [%rd49+8];
ld.f64 %fd18, [%rd49];
cvt.rn.ftz.f32.f64 %f110, %fd18;
cvt.rn.ftz.f32.f64 %f111, %fd17;
cvt.rn.ftz.f32.f64 %f112, %fd16;
ld.global.u32 %r73, [refl_max_depth];
setp.lt.s32 %p16, %r166, %r73;
@%p16 bra BB0_11;
BB0_20:
ld.global.u32 %r168, [refl_max_depth];
BB0_21:
add.s32 %r74, %r168, 1;
ld.local.u32 %r169, [%rd5+4];
setp.ne.s32 %p17, %r169, %r74;
@%p17 bra BB0_23;
add.s32 %r169, %r169, -1;
st.local.u32 [%rd5+4], %r169;
BB0_23:
st.local.v4.u8 [%rd5+24], {%rs85, %rs86, %rs87, %rs88};
ld.global.u32 %r13, [refl_iterations_number];
mul.lo.s32 %r75, %r168, %r13;
selp.b32 %r76, 0, %r75, %p3;
add.s32 %r14, %r76, %r169;
ld.global.u32 %r15, [measurement_plane];
setp.ne.s32 %p19, %r15, 0;
selp.u32 %r77, 1, 0, %p19;
sub.s32 %r16, %r14, %r77;
ld.global.u32 %r78, ;
setp.eq.s32 %p20, %r78, 0;
@%p20 bra BB0_32;
ld.local.u32 %r79, [%rd5+8];
setp.eq.s32 %p21, %r79, 0;
@%p21 bra BB0_32;
ld.global.v2.u32 {%r82, %r83}, [launch_index];
cvt.u64.u32 %rd68, %r82;
cvt.u64.u32 %rd69, %r83;
mov.u64 %rd71, 0;
// inline asm
call (%rd66), _rt_buffer_get_64, (%rd17, %r23, %r24, %rd68, %rd69, %rd71, %rd71);
// inline asm
ld.global.f32 %f40, ;
mov.f32 %f73, 0f00000000;
setp.le.ftz.f32 %p22, %f40, 0f00000000;
mov.f32 %f119, %f73;
@%p22 bra BB0_27;
setp.ltu.ftz.f32 %p23, %f40, 0f3F800000;
mul.ftz.f32 %f74, %f40, 0f437FFD71;
selp.f32 %f41, %f74, 0f437FFD71, %p23;
mov.f32 %f119, %f41;
BB0_27:
mov.f32 %f42, %f119;
cvt.rzi.ftz.u32.f32 %r17, %f42;
ld.global.f32 %f43, ;
setp.le.ftz.f32 %p24, %f43, 0f00000000;
mov.f32 %f118, %f73;
@%p24 bra BB0_29;
setp.ltu.ftz.f32 %p25, %f43, 0f3F800000;
mul.ftz.f32 %f76, %f43, 0f437FFD71;
selp.f32 %f118, %f76, 0f437FFD71, %p25;
BB0_29:
cvt.rzi.ftz.u32.f32 %r18, %f118;
ld.global.f32 %f46, ;
setp.le.ftz.f32 %p26, %f46, 0f00000000;
mov.f32 %f117, %f73;
@%p26 bra BB0_31;
setp.ltu.ftz.f32 %p27, %f46, 0f3F800000;
mul.ftz.f32 %f78, %f46, 0f437FFD71;
selp.f32 %f117, %f78, 0f437FFD71, %p27;
BB0_31:
cvt.rzi.ftz.u32.f32 %r86, %f117;
cvt.u16.u32 %rs59, %r18;
cvt.u16.u32 %rs60, %r17;
cvt.u16.u32 %rs61, %r86;
mov.u16 %rs62, 255;
st.v4.u8 [%rd66], {%rs60, %rs59, %rs61, %rs62};
bra.uni BB0_52;
BB0_32:
setp.ne.s32 %p28, %r169, %r168;
@%p28 bra BB0_36;
setp.ne.s32 %p29, %r13, 0;
setp.ne.s32 %p30, %r168, 1;
or.pred %p31, %p29, %p30;
@!%p31 bra BB0_36;
bra.uni BB0_34;
BB0_34:
setp.eq.s32 %p32, %r15, 0;
@%p32 bra BB0_51;
ld.local.u32 %r87, [%rd5+12];
setp.eq.s32 %p33, %r87, 0;
@%p33 bra BB0_51;
BB0_36:
ld.global.f64 %fd19, [cone_angle];
setp.ge.f64 %p36, %fd19, 0d0000000000000000;
or.pred %p37, %p19, %p36;
mov.pred %p61, 0;
@!%p37 bra BB0_38;
bra.uni BB0_37;
BB0_37:
ld.local.u32 %r88, [%rd5+12];
setp.eq.s32 %p61, %r88, 0;
BB0_38:
setp.gt.s32 %p38, %r14, 0;
and.pred %p39, %p38, %p61;
@!%p39 bra BB0_40;
bra.uni BB0_39;
BB0_39:
setp.ne.s32 %p40, %r13, 0;
setp.ne.s32 %p41, %r168, 1;
setp.ne.s32 %p42, %r169, 1;
or.pred %p43, %p41, %p42;
or.pred %p44, %p40, %p43;
@%p44 bra BB0_50;
bra.uni BB0_40;
BB0_50:
ld.global.v2.u32 {%r108, %r109}, [launch_index];
cvt.u64.u32 %rd89, %r108;
cvt.u64.u32 %rd90, %r109;
mov.u64 %rd92, 0;
// inline asm
call (%rd87), _rt_buffer_get_64, (%rd17, %r23, %r24, %rd89, %rd90, %rd92, %rd92);
// inline asm
mov.f32 %f96, 0f00000000;
cvt.rzi.ftz.u32.f32 %r112, %f96;
mov.f32 %f97, 0f432664BC;
cvt.rzi.ftz.u32.f32 %r113, %f97;
mov.f32 %f98, 0f437FFD71;
cvt.rzi.ftz.u32.f32 %r114, %f98;
cvt.u16.u32 %rs75, %r114;
cvt.u16.u32 %rs76, %r113;
cvt.u16.u32 %rs77, %r112;
mov.u16 %rs78, 255;
st.v4.u8 [%rd87], {%rs77, %rs76, %rs75, %rs78};
bra.uni BB0_52;
BB0_40:
ld.global.u32 %r89, ;
ld.global.u32 %r19, ;
setp.ge.s32 %p45, %r16, %r19;
setp.ne.s32 %p46, %r89, 0;
and.pred %p47, %p45, %p46;
@%p47 bra BB0_43;
bra.uni BB0_41;
BB0_43:
sub.s32 %r98, %r16, %r19;
cvt.rn.f32.s32 %f80, %r98;
ld.global.u32 %r99, [refl_max_bounces];
sub.s32 %r100, %r99, %r19;
cvt.rn.f32.s32 %f81, %r100;
div.approx.ftz.f32 %f82, %f80, %f81;
ld.global.v2.u32 {%r101, %r102}, [launch_index];
cvt.u64.u32 %rd82, %r101;
cvt.u64.u32 %rd83, %r102;
mov.u64 %rd85, 0;
// inline asm
call (%rd80), _rt_buffer_get_64, (%rd17, %r23, %r24, %rd82, %rd83, %rd85, %rd85);
// inline asm
fma.rn.ftz.f32 %f83, %f82, 0f40000000, 0fBF800000;
setp.lt.ftz.f32 %p49, %f83, 0f00000000;
selp.f32 %f49, 0f00000000, %f83, %p49;
mov.f32 %f84, 0f3F800000;
sub.ftz.f32 %f85, %f84, %f82;
fma.rn.ftz.f32 %f86, %f85, 0f40000000, 0fBF800000;
setp.lt.ftz.f32 %p50, %f86, 0f00000000;
selp.f32 %f50, 0f00000000, %f86, %p50;
mov.f32 %f87, 0f3F000000;
sub.ftz.f32 %f88, %f87, %f82;
abs.ftz.f32 %f89, %f88;
sub.ftz.f32 %f90, %f87, %f89;
add.ftz.f32 %f51, %f90, %f90;
mov.f32 %f79, 0f00000000;
setp.le.ftz.f32 %p51, %f50, 0f00000000;
mov.f32 %f124, %f79;
@%p51 bra BB0_45;
setp.ltu.ftz.f32 %p52, %f50, 0f3F800000;
mul.ftz.f32 %f91, %f50, 0f437FFD71;
selp.f32 %f52, %f91, 0f437FFD71, %p52;
mov.f32 %f124, %f52;
BB0_45:
mov.f32 %f53, %f124;
cvt.rzi.ftz.u32.f32 %r20, %f53;
setp.le.ftz.f32 %p53, %f51, 0f00000000;
mov.f32 %f123, %f79;
@%p53 bra BB0_47;
setp.ltu.ftz.f32 %p54, %f51, 0f3F800000;
mul.ftz.f32 %f93, %f51, 0f437FFD71;
selp.f32 %f123, %f93, 0f437FFD71, %p54;
BB0_47:
cvt.rzi.ftz.u32.f32 %r21, %f123;
setp.le.ftz.f32 %p55, %f49, 0f00000000;
mov.f32 %f122, %f79;
@%p55 bra BB0_49;
setp.ltu.ftz.f32 %p56, %f49, 0f3F800000;
mul.ftz.f32 %f95, %f49, 0f437FFD71;
selp.f32 %f122, %f95, 0f437FFD71, %p56;
BB0_49:
cvt.rzi.ftz.u32.f32 %r105, %f122;
cvt.u16.u32 %rs71, %r21;
cvt.u16.u32 %rs72, %r20;
cvt.u16.u32 %rs73, %r105;
mov.u16 %rs74, 255;
st.v4.u8 [%rd80], {%rs72, %rs71, %rs73, %rs74};
bra.uni BB0_52;
BB0_41:
setp.ne.s32 %p48, %r13, 0;
@%p48 bra BB0_52;
ld.global.v2.u32 {%r92, %r93}, [launch_index];
cvt.u64.u32 %rd75, %r92;
cvt.u64.u32 %rd76, %r93;
mov.u64 %rd78, 0;
// inline asm
call (%rd73), _rt_buffer_get_64, (%rd17, %r23, %r24, %rd75, %rd76, %rd78, %rd78);
// inline asm
ld.local.v4.u8 {%rs63, %rs64, %rs65, %rs66}, [%rd5+24];
st.v4.u8 [%rd73], {%rs63, %rs64, %rs65, %rs66};
bra.uni BB0_52;
BB0_51:
ld.global.v2.u32 {%r117, %r118}, [launch_index];
cvt.u64.u32 %rd96, %r117;
cvt.u64.u32 %rd97, %r118;
mov.u64 %rd99, 0;
// inline asm
call (%rd94), _rt_buffer_get_64, (%rd17, %r23, %r24, %rd96, %rd97, %rd99, %rd99);
// inline asm
mov.f32 %f99, 0f432664BC;
cvt.rzi.ftz.u32.f32 %r121, %f99;
mov.f32 %f100, 0f4359976D;
cvt.rzi.ftz.u32.f32 %r122, %f100;
mov.f32 %f101, 0f437FFD71;
cvt.rzi.ftz.u32.f32 %r123, %f101;
cvt.u16.u32 %rs79, %r123;
cvt.u16.u32 %rs80, %r122;
cvt.u16.u32 %rs81, %r121;
mov.u16 %rs82, 255;
st.v4.u8 [%rd94], {%rs81, %rs80, %rs79, %rs82};
BB0_52:
ld.local.u32 %r126, [%rd5+4];
ld.global.v2.u32 {%r127, %r128}, [launch_index];
cvt.u64.u32 %rd103, %r127;
cvt.u64.u32 %rd104, %r128;
mov.u64 %rd107, count_buffer;
cvta.global.u64 %rd102, %rd107;
mov.u64 %rd106, 0;
// inline asm
call (%rd101), _rt_buffer_get_64, (%rd102, %r23, %r24, %rd103, %rd104, %rd106, %rd106);
// inline asm
st.u32 [%rd101], %r126;
ld.global.u64 %rd8, [launch_index];
ld.global.u32 %r131, [refl_max_depth];
ld.local.u32 %r22, [%rd5+4];
mov.u16 %rs89, 0;
setp.ne.s32 %p57, %r22, %r131;
@%p57 bra BB0_55;
ld.local.u64 %rd9, [%rd5+8];
cvt.u32.u64 %r132, %rd9;
setp.ne.s32 %p58, %r132, 0;
@%p58 bra BB0_55;
setp.lt.u64 %p59, %rd9, 4294967296;
selp.u16 %rs89, 1, 0, %p59;
BB0_55:
setp.eq.s16 %p60, %rs89, 0;
@%p60 bra BB0_57;
shr.u64 %rd120, %rd8, 32;
add.s32 %r137, %r22, -1;
cvt.u64.u32 %rd121, %r137;
mul.lo.s64 %rd122, %rd121, %rd14;
add.s64 %rd123, %rd122, %rd120;
mul.lo.s64 %rd124, %rd123, %rd13;
add.s64 %rd125, %rd124, %rd8;
and.b64 %rd110, %rd125, 4294967295;
mov.u64 %rd126, raybounces_buffer;
cvta.global.u64 %rd109, %rd126;
mov.u32 %r133, 1;
mov.u32 %r134, 96;
// inline asm
call (%rd108), _rt_buffer_get_64, (%rd109, %r133, %r134, %rd110, %rd106, %rd106, %rd106);
// inline asm
ld.f64 %fd21, [%rd108];
ld.f64 %fd22, [%rd108+8];
ld.f64 %fd23, [%rd108+16];
ld.global.v2.u32 {%r138, %r139}, [launch_index];
cvt.u64.u32 %rd116, %r138;
cvt.u64.u32 %rd117, %r139;
mov.u64 %rd127, last_out_dir_buffer;
cvta.global.u64 %rd115, %rd127;
mov.u32 %r136, 24;
// inline asm
call (%rd114), _rt_buffer_get_64, (%rd115, %r23, %r136, %rd116, %rd117, %rd106, %rd106);
// inline asm
ld.f64 %fd24, [%rd114];
ld.f64 %fd25, [%rd114+8];
ld.f64 %fd26, [%rd114+16];
BB0_57:
ld.global.u32 %rd152, [launch_index+4];
mul.lo.s64 %rd153, %rd152, %rd13;
add.s64 %rd154, %rd153, %rd8;
and.b64 %rd130, %rd154, 4294967295;
mov.u64 %rd155, launchpoint_buffer;
cvta.global.u64 %rd129, %rd155;
mov.u32 %r142, 1;
mov.u32 %r143, 56;
// inline asm
call (%rd128), _rt_buffer_get_64, (%rd129, %r142, %r143, %rd130, %rd106, %rd106, %rd106);
// inline asm
st.u8 [%rd128+48], %rs89;
st.f64 [%rd128+40], %fd26;
st.f64 [%rd128+32], %fd25;
st.f64 [%rd128+24], %fd24;
st.f64 [%rd128+16], %fd23;
st.f64 [%rd128+8], %fd22;
st.f64 [%rd128], %fd21;
ld.local.u32 %r150, [%rd5+8];
ld.global.v2.u32 {%r151, %r152}, [launch_index];
cvt.u64.u32 %rd136, %r151;
cvt.u64.u32 %rd137, %r152;
mov.u64 %rd156, front_stop_buffer;
cvta.global.u64 %rd135, %rd156;
// inline asm
call (%rd134), _rt_buffer_get_64, (%rd135, %r23, %r24, %rd136, %rd137, %rd106, %rd106);
// inline asm
st.u32 [%rd134], %r150;
ld.global.v2.u32 {%r155, %r156}, [launch_index];
cvt.u64.u32 %rd142, %r155;
cvt.u64.u32 %rd143, %r156;
mov.u64 %rd157, measured_buffer;
cvta.global.u64 %rd141, %rd157;
// inline asm
call (%rd140), _rt_buffer_get_64, (%rd141, %r23, %r24, %rd142, %rd143, %rd106, %rd106);
// inline asm
ld.local.u32 %r159, [%rd5+12];
ld.u32 %r160, [%rd140];
add.s32 %r161, %r159, %r160;
ld.global.v2.u32 {%r162, %r163}, [launch_index];
cvt.u64.u32 %rd148, %r162;
cvt.u64.u32 %rd149, %r163;
// inline asm
call (%rd146), _rt_buffer_get_64, (%rd141, %r23, %r24, %rd148, %rd149, %rd106, %rd106);
// inline asm
st.u32 [%rd146], %r161;
ret;
}
// .globl _Z14camera_raytubeN5optix3RayEj
.visible .func _Z14camera_raytubeN5optix3RayEj(
.param .b64 _Z14camera_raytubeN5optix3RayEj_param_0,
.param .b32 _Z14camera_raytubeN5optix3RayEj_param_1
)
{
.local .align 16 .b8 __local_depot1[96];
.reg .b64 %SP;
.reg .b64 %SPL;
.reg .pred %p<18>;
.reg .b16 %rs<6>;
.reg .f32 %f<51>;
.reg .b32 %r<59>;
.reg .f64 %fd<53>;
.reg .b64 %rd<62>;
mov.u64 %rd61, __local_depot1;
cvta.local.u64 %SP, %rd61;
ld.param.u64 %rd1, [_Z14camera_raytubeN5optix3RayEj_param_0];
ld.param.u32 %r12, [_Z14camera_raytubeN5optix3RayEj_param_1];
add.u64 %rd5, %SP, 0;
cvta.to.local.u64 %rd6, %rd5;
ld.f32 %f31, [%rd1];
ld.f32 %f32, [%rd1+4];
ld.f32 %f33, [%rd1+8];
cvt.ftz.f64.f32 %fd13, %f33;
add.s64 %rd2, %rd6, 72;
mov.u32 %r57, 1;
mov.u32 %r55, 0;
st.local.v2.u32 [%rd6+72], {%r55, %r57};
cvt.ftz.f64.f32 %fd14, %f32;
cvt.ftz.f64.f32 %fd15, %f31;
st.local.v2.u32 [%rd6+80], {%r55, %r12};
st.local.v2.f64 [%rd6+16], {%fd15, %fd14};
st.local.f64 [%rd6+32], %fd13;
mov.u64 %rd7, 0;
st.local.u64 [%rd6+64], %rd7;
ld.global.u32
camera_ptx.txt (201 KB)