Working Optix 3 App fails Optix 4

Gaming and Visualization Technologies Visualization OptiX
1

I have developed an Optix App over the years and wanted to keep marginally up with the times and have just recompiled for Optix 4.1.1 from Optix 3.6.3 .
(Optix 4.1.1 => Using Visual Studio 2015 Update 3 and Cuda 8 sm_3 for GP100)

Now it compiled out of the box, but it fails when reading a PTX file with rtCreateProgramFromPTXFile, with a compile error mentioning an uninitialized data var in a module / program (i am writing this from home so don’t have the full error listing at hand).

My questions are ;

  • Why is the cu to ptx compile stage not telling me about obvious issues ?
    • At least for Optix 3.6.3, vaguely similar issues would come to light at the validate stage, and would be due to not issuing context->set values on a global GPU vars.
  • I have looked at the code for obvious issues, nothing jumps out at me and am now stuck
    • How can i debug this behind the scenes compile from PTX to whatever ?
  • Is there some incompatibilty in my choice development tool versions see opening statement?

Eagerly awaiting some tips / advice

Andy Page.

2

The CUDA source to PTX source compilation itself is outside of OptiX.
Issues about variables not declared in respective OptiX scopes cannot be detected at that point.
That would be be caught during the validaton and compilation steps inside OptiX which happens at latest at the first launch.

No, MSVS 2015 SP3, CUDA 8.0 compiling to SM 3.0 is perfectly fine and should work on all supported boards (Kepler and newer). That matches what I had been using with OptiX 4 versions.

Though if you recently upgraded to CUDA 8.0 an additional compiler option must be set to let CUDA generate PTX code for callable programs: --relocatable-device-code=true.
Please find more information in this post: https://devtalk.nvidia.com/default/topic/1000376/?comment=5111537

If you switch to a newer OptiX major version coming from 3.6.3, you should simply skip OptiX 4 versions altogether and use OptiX 5.0.0 instead. It supports the same range of GPU architectures as OptiX 4 while Volta GPU support is improved, it’s faster, contains bug fixes over OptiX 4, and added new features like motion blur and AI denoising: https://developer.nvidia.com/optix

If neither OptiX 5.0.0 nor the --relocatable-device-code=true affected the behavior, we would need some additional information to be able to analyze this.

3

Thanks, i will add that compile switch, and see what happens.

Optix 5.0.0 requires which CUDA and compiler version ? Is CUDA 8 and VS2015 SP3 okay ?

4

Yes, CUDA 8.0 and VS2015 SP3 is perfectly fine. CUDA 9.0 or 9.1 should work as well, mostly.
I have not experienced any issues when compiling for SM 3.0 to target all GPU architectures supported by OptiX with either CUDA 8.0 or 9.0, and haven’t switched to 9.1, yet.

If the error is about callable programs you would be able to see that effect inside the resulting PTX source code as well. Without the --relocatable-device-code=true option the functions are simply missing with CUDA 8.0 and newer.

5

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)

6

I think the forum didn’t accept the complete code. (I added code blocks.)
You can attach files to the forum on submitted posts with the small paper clip icon in the top right when hovering over a submitted post.

What have been the rtProgramCreateFromPTXFile() arguments (esp the program_name) when that happened?
There is no “camera_orthographic” in this part of the PTX code.

7

I have renamed the camera.ptx to camera_ptx.txt and attached to associated post.
It says scanning but hopefully all will complete okay.

8

Thanks, I can reproduce the error with OptiX 5.0.0 and file a bug report for analysis.

Sounds like an LLVM conversion problem for a double precision instruction, which would be a rare case and then could only be fixed in a newer OptiX version.
Normally OptiX programs avoid double precision calculations because of the involved cost.
I assume you need the precision, otherwise you could test if the problem goes away with single precision floating point code.

9

Some time in the past we thought we required double precision, maybe we will need to revisit this if the bug is not fixable.