ELF@ @8 @@@@nnnnnXOXO000( lRtd000PtdPPPQtd  , F"!+@Gj(EUrM v    L ] ~        6 H Z i x        Q a   ov%=LVn:Pa ,AUj*av%5<MT`p=V~!+DZn.(q:  8 ;$x  N$,QP) 8-$j 3    40pI(@,     \.8K $\  J0p  XF` HcH_0  p= P) dxp (mF Hf  P L$8  Lt* l9   0 2,  h`9o      jhV' <K8)     TL (  & U oX ! nh  6,%`[ 2P  HQP     v (|    pJ T1 |j4 WF |L ܂_wl?0'  8b yL0`  u Cs8E     ? Tt W@}  ` \A"  Hh   (X0 ,;ppZ 0|q  |/hi WX4TnPC  t  n Tj(gxV $Hk  )  T 5   { М7 HO@ h\ p 5{Xy L{ > Ew D\  h  <8@X_D  4P5PRPqPP_'8Ec2p .@B  @8@`  BD  D@QA @  K@D.(T@I  @!" @!@@ ! @RAH A!A@H  "@Hl  $ J  BB@! PHH@!@ Q@ #(+,.25:=FIKOQT3'6SߒI4@x!z#E~*;Wj; ƽCѿ3z>H#qɴՈ:@H. ? ioIn#+^쀙J$ b7d{E\-UF2pIj{zKߩx-?5|w)0W7v=+F^F[ k[+H!>x>@%㷐mt( q -D$x(}DW oC+Vέ! GJYvi %N|̣<積 ڗ*hJv~$/$(hκ˿ )M'UrT=>.x1:%i0x(ט[kh($4!=:C 5Un$H氒΀F j0;10OHe N GrSج)_Viz>gCEj^HO7/;߈NNk}0Kuʇ?犯e jv}H+D+᪶JP=4*~& aVx`NB{Ґ=H])`v}BB B%2/ {Mx bi߰ ;hY¥S(ӭ4VA[WFF_±k-I,@#s|@`Yۗy__gmon_start___ITM_deregisterTMCloneTable_ITM_registerTMCloneTable__cxa_finalizel4t2_init_handle0atm_nif_errorl4t2_nif_errorenif_get_stringenif_is_atomenif_get_local_pidatm_decoderatm_encodersnprintfenif_mutex_createenif_cond_createl4t2_resourceenif_alloc_resourcememsetenif_selfenif_monitor_processmemcpyenif_priv_datal4t2_component_initenif_release_resourceatm_l4t2_errorl4t2_errorenif_make_unique_integerenif_get_longenif_make_resourceatm_okenif_mutex_destroyenif_cond_destroyenif_make_tuplel4t2_get_composite_dataenif_get_tuplel4t2_get_rectenif_get_doubleenif_get_intl4t2_composite0enif_get_uintenif_get_list_lengthenif_get_list_cellnif_initenif_make_atomatm_erroratm_$transcoder_logatm_quietatm_panicatm_fatalatm_warningatm_yuvatm_infoatm_verboseatm_debugatm_unknownatm_eagainatm_nvbuf_erroratm_trueatm_falseatm_undefinedatm_destroy_meatm_waitatm_inputatm_outputatm_flushatm_flushedatm_fully_flushedatm_l4t2_eventatm_uyvy422atm_yuyv422atm_yuv420patm_yuv420p10atm_frame_widthatm_frame_heightatm_pixel_formatatm_fps_numatm_fps_denatm_qbufferatm_configatm_keyframeatm_frameatm_output_frameatm_output_frame_erroratm_EventPortSettingsChangedatm_EventPollCompleteatm_EventUnknownatm_output_buffer_readyatm_input_buffer_readyatm_decoder_buffer_readyatm_dec_bitstream_erroratm_dec_erroratm_dec_timeatm_no_statsatm_err_stratm_error_idatm_error_placeatm_downatm_l4t2_framerateatm_h264atm_hevcatm_mp2vatm_yadifatm_yadif_input_buffer_releaseatm_yadif_output_buffer_releaseatm_l4t2_bitrateatm_l4t2_profile_h264atm_l4t2_profile_h265atm_l4t2_level_h264atm_l4t2_presetatm_l4t2_idr_intervalatm_l4t2_num_bframesatm_l4t2_num_refsatm_l4t2_gop_sizeatm_l4t2_enable_audatm_l4t2_enable_vuiatm_l4t2_enable_sps_pps_at_idratm_l4t2_rate_control_modeatm_l4t2_enable_ext_colorformatatm_l4t2_virtual_buffer_sizeatm_l4t2_enable_error_reportingatm_l4t2_max_performanceatm_l4t2_frame_input_modeatm_l4t2_disable_dpbatm_l4t2_gop_closureatm_l4t2_skip_framesatm_l4t2_temporal_tradeoff_levelatm_l4t2_dec_drop_frame_intervalatm_l4t2_qp_rangeatm_l4t2_enc_two_pass_cbratm_l4t2_enc_peak_bitrateatm_l4t2_force_idratm_logoatm_alogoenif_open_resource_type_xyadif_resourcecu_egl_loadcu_egl_unloadenif_get_resourcel4t2_component_init_output_buffersl4t2_component_init_capture_buffersenif_make_listl4t2_component_get_bufferenif_make_ulongenif_make_list_celll4t2_atom2idl4t2_component_set_output_plane_formatl4t2_component_set_enabledenif_thread_createNvBufferSessionCreatel4t2_component_set_capture_plane_formatl4t2_plane_deinitNvBufferCreateExenif_make_intl4t2_plane_qbufferenif_inspect_binaryenif_mutex_lockenif_mutex_unlockl4t2_error_placereallocmp4_to_annexbenif_cond_broadcastenif_is_listenif_get_ulongenif_get_atoml4t2_component_set_parameterNvBufferTransformNvBufferGetParamsenif_cond_waitenif_alloc_envenif_make_longenif_sendenif_free_env__errno_locationl4t2_plane_dqbufferenif_thread_joinl4t2_component_poll_set_interruptl4t2_component_dq_eventusleepl4t2_plane_get_formatl4t2_component_get_sizel4t2_component_get_minimum_capture_buffersl4t2_plane_set_stream_statuscu_egl_initl4t2_plane_get_dma_bufNvBufferCompositecu_egl_compositecu_egl_freeannexb_to_mp4enif_make_binaryl4t2_plane_release_dma_bufcu_yadif_alloccu_yadif_sendcu_yadif_recvcu_yadif_freeNvBufferDestroyfreel4t2_component_destroyl4t2_logo_setupl4t2_logo_removel4t2_logo_scalev4l2_openv4l2_ioctll4t2_component_subscribe_eventl4t2_plane_set_formatl4t2_buffer_fill_planes_formatl4t2_plane_setup_buffersl4t2_plane_get_bufferl4t2_plane_get_cropl4t2_plane_set_stream_paramsl4t2_component_get_capture_plane_formatl4t2_plane_get_parml4t2_component_closeNvBufferSessionDestroyv4l2_closel4t2_component_poll_devicel4t2_component_get_input_metadatal4t2_component_get_output_metadatainitializer_mutexl4t2_plane_set_parml4t2_plane_get_stream_statusl4t2_plane_free_dma_bufsl4t2_buffer_free_mmapl4t2_buffer_free_dmal4t2_buffer_initl4t2_buffer_setup_mmapl4t2_buffer_setup_dmal4t2_buffer_setup_userptrl4t2_alloc_nvbufferl4t2_pix2pixl4t2_plane_flushl4t2_plane_send_eosmunmapmmapl4t2_buffer_setup_nvbuffermallocl4t2_error2strstrerrorerror_replyenif_make_badargav_frame_allocav_frame_get_bufferNvBuffer2Rawav_image_fill_arraysav_image_copyav_frame_apply_croppingsws_getContextsws_scaleav_image_copy_to_bufferRaw2NvBuffercu_egl_close_fdav_frame_freesws_freeContextff_load_image_binav_expr_parseav_expr_evalav_load_logsenif_set_pid_undefinedfetch_and_set_logger_pidatom_from_av_log_levelnif_log_handlerenif_is_pid_undefinedenif_make_stringenif_clear_envlog_callback_nullstrcmpav_log_get_levelav_log_format_linestderrfputsset_enable_ffmpeg_logsntohlenif_alloc_binaryav_logav_mallocavio_alloc_contextav_freeavformat_alloc_contextav_find_input_formatavformat_open_inputavformat_close_inputavformat_find_stream_infoavcodec_find_decoderavcodec_alloc_context3avcodec_free_contextavcodec_open2av_packet_allocav_read_frameav_packet_freeavcodec_send_packetavcodec_receive_frameav_frame_clonemake_encoded_framemake_raw_video_frameav_default_item_namecuStreamSynchronizecuCtxSynchronizecuLaunchKernelcuCtxCreate_v2scale_ptxcuModuleLoadDatacuModuleGetFunctionoverlay_ptxcuStreamCreatecuModuleUnloadcuCtxPopCurrent_v2cuCtxDestroy_v2cuInitcuDeviceGetCountcalloccuDeviceGetcuDeviceGetNamecuStreamDestroy_v2cuSurfObjectDestroycuGraphicsUnregisterResourceNvDestroyEGLImageNvEGLImageFromFdcuGraphicsEGLRegisterImagecuGraphicsResourceGetMappedEglFramecuSurfObjectCreatecu_array_from_framecuMemcpy2D_v2free_graphavfilter_graph_freeinit_hw_video_graphavfilter_get_by_nameavfilter_graph_allocavfilter_inout_allocavfilter_graph_create_filterav_buffersrc_parameters_setav_int_list_length_for_sizeav_opt_set_binav_strdupavfilter_graph_parse_ptravfilter_graph_configavfilter_inout_freeav_hwdevice_ctx_createav_hwframe_ctx_allocav_hwframe_ctx_initav_buffer_unrefav_buffersink_get_frameav_rescale_qav_hwframe_get_bufferav_buffersrc_add_frame_flagslibavcodec.so.60LIBAVCODEC_60libavutil.so.58LIBAVUTIL_58libavfilter.so.9LIBAVFILTER_9libavformat.so.60LIBAVFORMAT_60libc.so.6GLIBC_2.17/usr/lib/aarch64-linux-gnu/tegra:/opt/flussonic/lib/aarch64-linux-gnu/libcuda.so.1libasound.so.2libv4l2.so.0libnvbuf_utils.so.1.0.00o8`o@r}HyP}Xl`y}h}plxBGv}#|H4|Lh (r@yH}pXhoxvtwzyl(s8HmXܗhux{k00|C|k(q8HuXdhqxX}nP|0<u z(mHvjhpxhjm\sP8h0`x(8X ``p(H@0@PXh x  x   !#$8%X&P(0)+,H.p/h02@3H4567X8:=>@G(HKLQ T8UpV(08@HPX` h p x  ?OF (08@MHP X`h!pRx"#$%&'A()*+;,-. /(B0081@2H3P4X5`6h7px89-W:;SC9< =(0>8?@HPX`'hp*xD@AJ"BCDE1FG H(I0J8K@LHMPNXO`PhEpQxRSTUVWXYZ[]^_`abcd e(f0g8h@iHjPkXl`mhnpoxpqrtuvwxyz{|}~ (08@HPX`hpxsrc_comp_rectscale_logo0enomemtransform_dst_rectfatalinfofalsel4t2_num_refsL4T2_ERROR_CAPABILITYL4T2_ERROR_NVBUFFER_COMPOSITElogo_av_frame_apply_croppingempty_logo_dataRaw2NvBufferff_load_image: avcodec_open2 failedl4t2.%s_output_mtx[%s]monitorbad buffer ptrrelease_buffer0bad prop idundefinedl4t2_virtual_buffer_sizel4t2_qp_rangeL4T2_ERROR_PARAM_NO_MMAParg_8_must_binset_parameter0set_parameterdebugdec_timeL4T2_ERROR_NONEL4T2_ERROR_MMAPL4T2_ERROR_PARAM_ARGCL4T2_ERROR_CUEGL_UNSUPPORTED_PIXFMTarg_3_must_intpix_fmtsbad prop valueunknown_pix_fmtinputfps_numl4t2_max_performanceL4T2_ERROR_PARAM_ROLEHoverlay_hff_load_image_bin: av_find_input_format failedsurf_copy_p016_nv12DecOutputDecCaptureL4T2_ERROR_CUEGL_KERNEL_OVERLAYtranscoder_streamer_logsurf_resize_p016_nv12arg3component_initget_idinit_encoder_core0remove_logo0set_capture_plane_formatbad_resource_statedec_errorL4T2_ERROR_CUEGL_IMAGE_GETff_load_image: avcodec_send_packet failedsws_getContext fail %dx%dsurf_resize_nv12outcall_yadif0set_enabledyadif_down_mtx[%s]yuv420p10yadif_input_buffer_releasel4t2_enable_audL4T2_ERROR_OPEN_DEVICEsurf_fill_p016configure_decoder_input0transform_dec_bufferalloc_resourceunknownerror_placel4t2_disable_dpblogol4t2.%s_plane_capture_mtx[%s]L4T2_ERROR_CUEGL_FUNC_GETarg_2_must_intav_expr_parse_yEncOutputcompositetransform_src_rectverbosel4t2_bitratel4t2_gop_sizeL4T2_ERROR_NVBUFFER_APIL4T2_ERROR_BUFFER_BUSYL4T2_ERROR_CUEGL_STREAM_CREATEyadif_cuda=0:-1:0bad tuple lengthallocate_buffers0arg4uyvy422dec_bitstream_errorl4t2_frameratealogoyadif_l4t2_resourceff_load_image_bin: avio_alloc_context failedtransform0decoder_panicfps_denEventPollCompletel4t2_level_h264l4t2_skip_framesl4t2_enc_two_pass_cbrL4T2_ERROR_ENOMEML4T2_ERROR_CUEGL_IMAGE_MAParg_7_must_binlogo_load_failedff_load_image: av_read_frame failedindown_mtxinit_yadif0get buffertransform_flagl4t2_enable_ext_colorformatl4t2_frame_input_modeL4T2_ERROR_PARAM_NUM_PLANESL4T2_ERROR_CUEGL_CTX_CREATEarg_1_must_fdlogo_sws_scalesurf_copy_p016surf_copy_nv12l4t2.%s_down_mtx[%s]beam.vanillainit_output_buffersthread_createencoder_panicwarningencoderframeerr_strl4t2_enc_peak_bitrateL4T2_ERROR_UNKNOWNL4T2_ERROR_CUEGL_KERNEL_RESIZEav_expr_parse_x/root/flussonic/apps/transcoder/c_src/l4t2/l4t2_handler.cnif_loaded0bad directionfail init buffersset_output_plane_formatplane_deinitwaitkeyframel4t2_presetl4t2_enable_vuilogo_av_frame_get_bufferav_read_frame: avcodec_receive_frame failedsurf_copy_nv12_p016free_buffers0$transcoder_logtruepixel_formatl4t2_num_bframeslogo_enomemlogo_Raw2NvBufferalogo_destroy surf_copy_nv12_yuv420p10bityadif_cnd[%s]yadif_paniceagaininput_buffer_readyL4T2_ERROR_PARAM_PIXFMTL4T2_ERROR_ANNEXBL4T2_ERROR_CUEGL_SURF_CREATEff_load_image: av_frame_alloc failedencbad tuplepush_dec_frame0buffer_to_fdconfigure_fake_decoder0destroy_mel4t2_profile_h264L4T2_ERROR_CUEGL_STREAM_SYNCbad_logo_offsetxff_load_image_bin: av_malloc failedff_load_image_bin: avformat_alloc_context failedsurf_resize_nv12_yuv420pblend_rgba_nv12output_condconfigure_encoder0NvBufferCreateExl4t2_eventoutput_frameL4T2_ERROR_SUBSCRIBE_EVENTL4T2_ERROR_CUEGL_IMAGE_DESTROYL4T2_ERROR_CUEGL_KERNEL_FILLL4T2_ERROR_CUEGL_MODULE_LOADff_load_image: avcodec_alloc_context3 failed init_handle0param record namel4t2_errorl4t2_rate_control_model4t2_temporal_tradeoff_levelL4T2_ERROR_CUEGL_MAX_BUFFERSL4T2_ERROR_CUEGL_KERNEL_COPYff_load_image_bin: bad sar ptrsurf_copy_nv12_yuv420pblend_rgba_p016output_mtxyadif_mtx[%s]errorpanicyuvnif_erroroutputfully_flushedyadif_output_buffer_releaselogo_av_image_copy_to_bufferlogo_destroyff_load_image: failed no nb_streamsff_load_image: input_desc or time_base is NULLinval_fddecoderframe_widthl4t2_idr_intervall4t2_enable_error_reportingl4t2_force_idrL4T2_ERROR_CUEGL_ENC_PTRarg_4_must_intWerts-14.0init_blur0set_logo0argv1buffernvbuf_errordownpng_pipeff_load_image_bin: avformat_open_input failedargcdeccomponent_set_enabledNvBufferTransformquietarg_5_must_intarg_6_must_binsurf_fill_nv12bad fdrelease_handle0output_frame_errorno_statsl4t2_dec_drop_frame_intervalL4T2_ERROR_PARAM_STREAMSTATUSL4T2_ERROR_PARAM_DIRECTIONL4T2_ERROR_CUEGL_IMAGE_REGL4T2_ERROR_CUEGL_CTX_SYNCsurf_resize_nv12_yuv420p10bit0arg5argv0ctxyuyv422l4t2_gop_closureL4T2_ERROR_PARAM_NAMEL4T2_ERROR_COMPOSITEL4T2_ERROR_CUEGL_FRAME_TYPEbad alphal4t2_handlerrelease_buffer_fd0set_enabled0release_yadif0flushhevcl4t2_profile_h265L4T2_ERROR_STREAM_RUNNINGL4T2_ERROR_CUEGL_SURF_DESTROYL4T2_ERROR_EAGAINarg_0_must_encoderlogo_must_rgbabad_logo_sizeav_image_copymain_hbuffersinkdst_comp_rectmtxYADIFokerror_idl4t2_resourceL4T2_ERROR_NO_BUFFERSlogo_av_frame_allocmain_woverlay_wyarg0flushedarityyuv420ph264L4T2_ERROR_PARAM_ARGSlogo_NvBufferCreateExwff_load_image: find stream info failedff_load_image: frame_get_buffer fail %dx%dvideo_size=%dx%d:pix_fmt=%d:time_base=%d/%d:pixel_aspect=%d/%dsurf_resize_nv12_p016qbuffermp2vyadifl4t2_enable_sps_pps_at_idrL4T2_ERROR_PARAM_BUFTYPElogo_sws_getContextarg2init_capture_bufferscondsizeL4T2_ERROR_MAX_BUFFERSL4T2_ERROR_EOSlogo_NvBuffer2Rawff_load_image: avcodec_find_decoder failed surf_fill_yuv420pl4t2.%s_output_cnd[%s]arg6EncCapturein_fdconfigEventPortSettingsChangedEventUnknownoutput_buffer_readydecoder_buffer_readyL4T2_ERROR_PARAM_MEMTYPEL4T2_ERROR_CUEGL_IMAGE_UNREGlogo_av_image_fill_arrayshff_load_image_bin: no png image datablend_rgba_yuv420parg1push_enc_frame0init_fake_decoder0frame_heightL4T2_ERROR_UNSUPPORTED_PIXFMTL4T2_ERROR_CUEGL_NVBUFFERsurf_resize_p016// // Generated by NVIDIA NVVM Compiler // // Compiler Build ID: CL-29663091 // Cuda compilation tools, release 10.2, V10.2.300 // Based on LLVM 3.4svn // .version 6.5 .target sm_30 .address_size 64 // .globl surf_fill_yuv420p .visible .entry surf_fill_yuv420p( .param .u64 surf_fill_yuv420p_param_0, .param .u64 surf_fill_yuv420p_param_1, .param .u64 surf_fill_yuv420p_param_2, .param .u32 surf_fill_yuv420p_param_3, .param .u32 surf_fill_yuv420p_param_4, .param .u8 surf_fill_yuv420p_param_5, .param .u8 surf_fill_yuv420p_param_6, .param .u8 surf_fill_yuv420p_param_7 ) { .reg .pred %p<7>; .reg .b16 %rs<4>; .reg .b32 %r<13>; .reg .b64 %rd<4>; ld.param.u64 %rd1, [surf_fill_yuv420p_param_0]; ld.param.u64 %rd2, [surf_fill_yuv420p_param_1]; ld.param.u64 %rd3, [surf_fill_yuv420p_param_2]; ld.param.u32 %r3, [surf_fill_yuv420p_param_3]; ld.param.u32 %r4, [surf_fill_yuv420p_param_4]; ld.param.u8 %rs3, [surf_fill_yuv420p_param_7]; ld.param.u8 %rs2, [surf_fill_yuv420p_param_6]; ld.param.u8 %rs1, [surf_fill_yuv420p_param_5]; mov.u32 %r5, %ctaid.x; mov.u32 %r6, %ntid.x; mov.u32 %r7, %tid.x; mad.lo.s32 %r1, %r6, %r5, %r7; mov.u32 %r8, %ntid.y; mov.u32 %r9, %ctaid.y; mov.u32 %r10, %tid.y; mad.lo.s32 %r2, %r8, %r9, %r10; setp.ge.s32 %p1, %r2, %r4; setp.ge.s32 %p2, %r1, %r3; or.pred %p3, %p1, %p2; @%p3 bra BB0_3; sust.b.2d.b8.trap [%rd1, {%r1, %r2}], {%rs1}; shr.s32 %r11, %r4, 1; setp.ge.s32 %p4, %r2, %r11; shr.s32 %r12, %r3, 1; setp.ge.s32 %p5, %r1, %r12; or.pred %p6, %p4, %p5; @%p6 bra BB0_3; sust.b.2d.b8.trap [%rd2, {%r1, %r2}], {%rs2}; sust.b.2d.b8.trap [%rd3, {%r1, %r2}], {%rs3}; BB0_3: ret; } // .globl surf_fill_nv12 .visible .entry surf_fill_nv12( .param .u64 surf_fill_nv12_param_0, .param .u64 surf_fill_nv12_param_1, .param .u32 surf_fill_nv12_param_2, .param .u32 surf_fill_nv12_param_3, .param .u8 surf_fill_nv12_param_4, .param .u8 surf_fill_nv12_param_5, .param .u8 surf_fill_nv12_param_6 ) { .reg .pred %p<7>; .reg .b16 %rs<4>; .reg .b32 %r<14>; .reg .b64 %rd<3>; ld.param.u64 %rd1, [surf_fill_nv12_param_0]; ld.param.u64 %rd2, [surf_fill_nv12_param_1]; ld.param.u32 %r3, [surf_fill_nv12_param_2]; ld.param.u32 %r4, [surf_fill_nv12_param_3]; ld.param.u8 %rs3, [surf_fill_nv12_param_6]; ld.param.u8 %rs2, [surf_fill_nv12_param_5]; ld.param.u8 %rs1, [surf_fill_nv12_param_4]; mov.u32 %r5, %ctaid.x; mov.u32 %r6, %ntid.x; mov.u32 %r7, %tid.x; mad.lo.s32 %r1, %r6, %r5, %r7; mov.u32 %r8, %ntid.y; mov.u32 %r9, %ctaid.y; mov.u32 %r10, %tid.y; mad.lo.s32 %r2, %r8, %r9, %r10; setp.ge.s32 %p1, %r2, %r4; setp.ge.s32 %p2, %r1, %r3; or.pred %p3, %p1, %p2; @%p3 bra BB1_3; sust.b.2d.b8.trap [%rd1, {%r1, %r2}], {%rs1}; shr.s32 %r11, %r4, 1; setp.ge.s32 %p4, %r2, %r11; shr.s32 %r12, %r3, 1; setp.ge.s32 %p5, %r1, %r12; or.pred %p6, %p4, %p5; @%p6 bra BB1_3; shl.b32 %r13, %r1, 1; sust.b.2d.v2.b8.trap [%rd2, {%r13, %r2}], {%rs2, %rs3}; BB1_3: ret; } // .globl surf_copy_nv12_yuv420p .visible .entry surf_copy_nv12_yuv420p( .param .u64 surf_copy_nv12_yuv420p_param_0, .param .u64 surf_copy_nv12_yuv420p_param_1, .param .u64 surf_copy_nv12_yuv420p_param_2, .param .u64 surf_copy_nv12_yuv420p_param_3, .param .u64 surf_copy_nv12_yuv420p_param_4, .param .u32 surf_copy_nv12_yuv420p_param_5, .param .u32 surf_copy_nv12_yuv420p_param_6 ) { .reg .pred %p<7>; .reg .b16 %rs<7>; .reg .b32 %r<14>; .reg .b64 %rd<6>; ld.param.u64 %rd1, [surf_copy_nv12_yuv420p_param_0]; ld.param.u64 %rd2, [surf_copy_nv12_yuv420p_param_1]; ld.param.u64 %rd3, [surf_copy_nv12_yuv420p_param_2]; ld.param.u64 %rd4, [surf_copy_nv12_yuv420p_param_3]; ld.param.u64 %rd5, [surf_copy_nv12_yuv420p_param_4]; ld.param.u32 %r3, [surf_copy_nv12_yuv420p_param_5]; ld.param.u32 %r4, [surf_copy_nv12_yuv420p_param_6]; mov.u32 %r5, %ctaid.x; mov.u32 %r6, %ntid.x; mov.u32 %r7, %tid.x; mad.lo.s32 %r1, %r6, %r5, %r7; mov.u32 %r8, %ntid.y; mov.u32 %r9, %ctaid.y; mov.u32 %r10, %tid.y; mad.lo.s32 %r2, %r8, %r9, %r10; setp.ge.s32 %p1, %r2, %r4; setp.ge.s32 %p2, %r1, %r3; or.pred %p3, %p1, %p2; @%p3 bra BB2_3; suld.b.2d.b8.trap {%rs1}, [%rd1, {%r1, %r2}]; and.b16 %rs2, %rs1, 255; sust.b.2d.b8.trap [%rd3, {%r1, %r2}], {%rs2}; shr.s32 %r11, %r4, 1; setp.ge.s32 %p4, %r2, %r11; shr.s32 %r12, %r3, 1; setp.ge.s32 %p5, %r1, %r12; or.pred %p6, %p4, %p5; @%p6 bra BB2_3; shl.b32 %r13, %r1, 1; suld.b.2d.v2.b8.trap {%rs3, %rs4}, [%rd2, {%r13, %r2}]; and.b16 %rs5, %rs3, 255; sust.b.2d.b8.trap [%rd4, {%r1, %r2}], {%rs5}; and.b16 %rs6, %rs4, 255; sust.b.2d.b8.trap [%rd5, {%r1, %r2}], {%rs6}; BB2_3: ret; } // .globl surf_resize_nv12_yuv420p .visible .entry surf_resize_nv12_yuv420p( .param .u64 surf_resize_nv12_yuv420p_param_0, .param .u64 surf_resize_nv12_yuv420p_param_1, .param .u64 surf_resize_nv12_yuv420p_param_2, .param .u64 surf_resize_nv12_yuv420p_param_3, .param .u64 surf_resize_nv12_yuv420p_param_4, .param .u32 surf_resize_nv12_yuv420p_param_5, .param .u32 surf_resize_nv12_yuv420p_param_6, .param .f32 surf_resize_nv12_yuv420p_param_7, .param .f32 surf_resize_nv12_yuv420p_param_8, .param .u32 surf_resize_nv12_yuv420p_param_9, .param .u32 surf_resize_nv12_yuv420p_param_10, .param .u32 surf_resize_nv12_yuv420p_param_11, .param .u32 surf_resize_nv12_yuv420p_param_12 ) { .reg .pred %p<7>; .reg .b16 %rs<22>; .reg .f32 %f<23>; .reg .b32 %r<73>; .reg .b64 %rd<6>; ld.param.u64 %rd1, [surf_resize_nv12_yuv420p_param_0]; ld.param.u64 %rd2, [surf_resize_nv12_yuv420p_param_1]; ld.param.u64 %rd3, [surf_resize_nv12_yuv420p_param_2]; ld.param.u64 %rd4, [surf_resize_nv12_yuv420p_param_3]; ld.param.u64 %rd5, [surf_resize_nv12_yuv420p_param_4]; ld.param.u32 %r6, [surf_resize_nv12_yuv420p_param_5]; ld.param.u32 %r7, [surf_resize_nv12_yuv420p_param_6]; ld.param.f32 %f2, [surf_resize_nv12_yuv420p_param_7]; ld.param.f32 %f3, [surf_resize_nv12_yuv420p_param_8]; ld.param.u32 %r8, [surf_resize_nv12_yuv420p_param_9]; ld.param.u32 %r9, [surf_resize_nv12_yuv420p_param_10]; ld.param.u32 %r10, [surf_resize_nv12_yuv420p_param_11]; ld.param.u32 %r11, [surf_resize_nv12_yuv420p_param_12]; mov.u32 %r12, %ctaid.x; mov.u32 %r13, %ntid.x; mov.u32 %r14, %tid.x; mad.lo.s32 %r1, %r13, %r12, %r14; mov.u32 %r15, %ntid.y; mov.u32 %r16, %ctaid.y; mov.u32 %r17, %tid.y; mad.lo.s32 %r2, %r15, %r16, %r17; setp.ge.s32 %p1, %r2, %r7; setp.ge.s32 %p2, %r1, %r6; or.pred %p3, %p1, %p2; @%p3 bra BB3_3; cvt.rn.f32.s32 %f4, %r1; add.f32 %f5, %f4, 0f3F000000; mul.f32 %f6, %f5, %f2; cvt.rn.f32.s32 %f7, %r2; add.f32 %f8, %f7, 0f3F000000; mul.f32 %f9, %f8, %f3; add.f32 %f10, %f2, 0fBF800000; mul.f32 %f11, %f10, 0f3F000000; cvt.sat.f32.f32 %f12, %f11; add.f32 %f13, %f3, 0fBF800000; mul.f32 %f14, %f13, 0f3F000000; cvt.sat.f32.f32 %f15, %f14; add.f32 %f16, %f12, 0f3F000000; div.rn.f32 %f17, %f12, %f16; add.f32 %f18, %f15, 0f3F000000; div.rn.f32 %f19, %f15, %f18; sub.f32 %f20, %f6, %f17; cvt.rzi.s32.f32 %r3, %f20; add.f32 %f21, %f6, %f17; cvt.rzi.s32.f32 %r4, %f21; sub.f32 %f1, %f9, %f19; add.f32 %f22, %f9, %f19; cvt.rzi.s32.f32 %r5, %f22; add.s32 %r18, %r3, %r10; add.s32 %r19, %r5, %r11; suld.b.2d.b8.trap {%rs1}, [%rd1, {%r18, %r19}]; cvt.u32.u16 %r20, %rs1; suld.b.2d.b8.trap {%rs2}, [%rd1, {%r18, %r19}]; cvt.u32.u16 %r21, %rs2; suld.b.2d.b8.trap {%rs3}, [%rd1, {%r18, %r19}]; cvt.u32.u16 %r22, %rs3; suld.b.2d.b8.trap {%rs4}, [%rd1, {%r18, %r19}]; cvt.u32.u16 %r23, %rs4; and.b32 %r24, %r20, 255; and.b32 %r25, %r21, 255; and.b32 %r26, %r22, 255; and.b32 %r27, %r23, 255; add.s32 %r28, %r24, %r25; add.s32 %r29, %r28, %r26; add.s32 %r30, %r29, %r27; add.s32 %r31, %r30, 2; cvt.u16.u32 %rs5, %r31; shr.u16 %rs6, %rs5, 2; and.b16 %rs7, %rs6, 255; add.s32 %r32, %r2, %r9; add.s32 %r33, %r1, %r8; sust.b.2d.b8.trap [%rd3, {%r33, %r32}], {%rs7}; shr.s32 %r34, %r7, 1; setp.ge.s32 %p4, %r2, %r34; shr.s32 %r35, %r6, 1; setp.ge.s32 %p5, %r1, %r35; or.pred %p6, %p4, %p5; @%p6 bra BB3_3; shr.s32 %r36, %r10, 1; shr.s32 %r37, %r11, 1; add.s32 %r38, %r5, %r37; cvt.rzi.s32.f32 %r39, %f1; add.s32 %r40, %r39, %r37; add.s32 %r41, %r3, %r36; shl.b32 %r42, %r41, 1; suld.b.2d.v2.b8.trap {%rs8, %rs9}, [%rd2, {%r42, %r40}]; cvt.u32.u16 %r43, %rs8; cvt.u32.u16 %r44, %rs9; add.s32 %r45, %r4, %r36; shl.b32 %r46, %r45, 1; suld.b.2d.v2.b8.trap {%rs10, %rs11}, [%rd2, {%r46, %r40}]; cvt.u32.u16 %r47, %rs10; cvt.u32.u16 %r48, %rs11; suld.b.2d.v2.b8.trap {%rs12, %rs13}, [%rd2, {%r42, %r38}]; cvt.u32.u16 %r49, %rs12; cvt.u32.u16 %r50, %rs13; suld.b.2d.v2.b8.trap {%rs14, %rs15}, [%rd2, {%r46, %r38}]; cvt.u32.u16 %r51, %rs14; cvt.u32.u16 %r52, %rs15; and.b32 %r53, %r43, 255; and.b32 %r54, %r47, 255; and.b32 %r55, %r49, 255; and.b32 %r56, %r51, 255; add.s32 %r57, %r53, %r54; add.s32 %r58, %r57, %r55; add.s32 %r59, %r58, %r56; add.s32 %r60, %r59, 2; and.b32 %r61, %r44, 255; and.b32 %r62, %r48, 255; and.b32 %r63, %r50, 255; and.b32 %r64, %r52, 255; add.s32 %r65, %r61, %r62; add.s32 %r66, %r65, %r63; add.s32 %r67, %r66, %r64; add.s32 %r68, %r67, 2; cvt.u16.u32 %rs16, %r60; shr.u16 %rs17, %rs16, 2; shr.s32 %r69, %r8, 1; add.s32 %r70, %r1, %r69; shr.s32 %r71, %r9, 1; add.s32 %r72, %r2, %r71; and.b16 %rs18, %rs17, 255; sust.b.2d.b8.trap [%rd4, {%r70, %r72}], {%rs18}; cvt.u16.u32 %rs19, %r68; shr.u16 %rs20, %rs19, 2; and.b16 %rs21, %rs20, 255; sust.b.2d.b8.trap [%rd5, {%r70, %r72}], {%rs21}; BB3_3: ret; } // .globl surf_resize_nv12 .visible .entry surf_resize_nv12( .param .u64 surf_resize_nv12_param_0, .param .u64 surf_resize_nv12_param_1, .param .u64 surf_resize_nv12_param_2, .param .u64 surf_resize_nv12_param_3, .param .u32 surf_resize_nv12_param_4, .param .u32 surf_resize_nv12_param_5, .param .f32 surf_resize_nv12_param_6, .param .f32 surf_resize_nv12_param_7, .param .u32 surf_resize_nv12_param_8, .param .u32 surf_resize_nv12_param_9, .param .u32 surf_resize_nv12_param_10, .param .u32 surf_resize_nv12_param_11 ) { .reg .pred %p<7>; .reg .b16 %rs<22>; .reg .f32 %f<23>; .reg .b32 %r<74>; .reg .b64 %rd<5>; ld.param.u64 %rd1, [surf_resize_nv12_param_0]; ld.param.u64 %rd2, [surf_resize_nv12_param_1]; ld.param.u64 %rd3, [surf_resize_nv12_param_2]; ld.param.u64 %rd4, [surf_resize_nv12_param_3]; ld.param.u32 %r6, [surf_resize_nv12_param_4]; ld.param.u32 %r7, [surf_resize_nv12_param_5]; ld.param.f32 %f2, [surf_resize_nv12_param_6]; ld.param.f32 %f3, [surf_resize_nv12_param_7]; ld.param.u32 %r8, [surf_resize_nv12_param_8]; ld.param.u32 %r9, [surf_resize_nv12_param_9]; ld.param.u32 %r10, [surf_resize_nv12_param_10]; ld.param.u32 %r11, [surf_resize_nv12_param_11]; mov.u32 %r12, %ctaid.x; mov.u32 %r13, %ntid.x; mov.u32 %r14, %tid.x; mad.lo.s32 %r1, %r13, %r12, %r14; mov.u32 %r15, %ntid.y; mov.u32 %r16, %ctaid.y; mov.u32 %r17, %tid.y; mad.lo.s32 %r2, %r15, %r16, %r17; setp.ge.s32 %p1, %r2, %r7; setp.ge.s32 %p2, %r1, %r6; or.pred %p3, %p1, %p2; @%p3 bra BB4_3; cvt.rn.f32.s32 %f4, %r1; add.f32 %f5, %f4, 0f3F000000; mul.f32 %f6, %f5, %f2; cvt.rn.f32.s32 %f7, %r2; add.f32 %f8, %f7, 0f3F000000; mul.f32 %f9, %f8, %f3; add.f32 %f10, %f2, 0fBF800000; mul.f32 %f11, %f10, 0f3F000000; cvt.sat.f32.f32 %f12, %f11; add.f32 %f13, %f3, 0fBF800000; mul.f32 %f14, %f13, 0f3F000000; cvt.sat.f32.f32 %f15, %f14; add.f32 %f16, %f12, 0f3F000000; div.rn.f32 %f17, %f12, %f16; add.f32 %f18, %f15, 0f3F000000; div.rn.f32 %f19, %f15, %f18; sub.f32 %f20, %f6, %f17; cvt.rzi.s32.f32 %r3, %f20; add.f32 %f21, %f6, %f17; cvt.rzi.s32.f32 %r4, %f21; sub.f32 %f1, %f9, %f19; add.f32 %f22, %f9, %f19; cvt.rzi.s32.f32 %r5, %f22; add.s32 %r18, %r3, %r10; add.s32 %r19, %r5, %r11; suld.b.2d.b8.trap {%rs1}, [%rd1, {%r18, %r19}]; cvt.u32.u16 %r20, %rs1; suld.b.2d.b8.trap {%rs2}, [%rd1, {%r18, %r19}]; cvt.u32.u16 %r21, %rs2; suld.b.2d.b8.trap {%rs3}, [%rd1, {%r18, %r19}]; cvt.u32.u16 %r22, %rs3; suld.b.2d.b8.trap {%rs4}, [%rd1, {%r18, %r19}]; cvt.u32.u16 %r23, %rs4; and.b32 %r24, %r20, 255; and.b32 %r25, %r21, 255; and.b32 %r26, %r22, 255; and.b32 %r27, %r23, 255; add.s32 %r28, %r24, %r25; add.s32 %r29, %r28, %r26; add.s32 %r30, %r29, %r27; add.s32 %r31, %r30, 2; cvt.u16.u32 %rs5, %r31; shr.u16 %rs6, %rs5, 2; and.b16 %rs7, %rs6, 255; add.s32 %r32, %r2, %r9; add.s32 %r33, %r1, %r8; sust.b.2d.b8.trap [%rd3, {%r33, %r32}], {%rs7}; shr.s32 %r34, %r7, 1; setp.ge.s32 %p4, %r2, %r34; shr.s32 %r35, %r6, 1; setp.ge.s32 %p5, %r1, %r35; or.pred %p6, %p4, %p5; @%p6 bra BB4_3; shr.s32 %r36, %r10, 1; shr.s32 %r37, %r11, 1; add.s32 %r38, %r5, %r37; cvt.rzi.s32.f32 %r39, %f1; add.s32 %r40, %r39, %r37; add.s32 %r41, %r3, %r36; shl.b32 %r42, %r41, 1; suld.b.2d.v2.b8.trap {%rs8, %rs9}, [%rd2, {%r42, %r40}]; cvt.u32.u16 %r43, %rs8; cvt.u32.u16 %r44, %rs9; add.s32 %r45, %r4, %r36; shl.b32 %r46, %r45, 1; suld.b.2d.v2.b8.trap {%rs10, %rs11}, [%rd2, {%r46, %r40}]; cvt.u32.u16 %r47, %rs10; cvt.u32.u16 %r48, %rs11; suld.b.2d.v2.b8.trap {%rs12, %rs13}, [%rd2, {%r42, %r38}]; cvt.u32.u16 %r49, %rs12; cvt.u32.u16 %r50, %rs13; suld.b.2d.v2.b8.trap {%rs14, %rs15}, [%rd2, {%r46, %r38}]; cvt.u32.u16 %r51, %rs14; cvt.u32.u16 %r52, %rs15; and.b32 %r53, %r43, 255; and.b32 %r54, %r47, 255; and.b32 %r55, %r49, 255; and.b32 %r56, %r51, 255; add.s32 %r57, %r53, %r54; add.s32 %r58, %r57, %r55; add.s32 %r59, %r58, %r56; add.s32 %r60, %r59, 2; and.b32 %r61, %r44, 255; and.b32 %r62, %r48, 255; and.b32 %r63, %r50, 255; and.b32 %r64, %r52, 255; add.s32 %r65, %r61, %r62; add.s32 %r66, %r65, %r63; add.s32 %r67, %r66, %r64; add.s32 %r68, %r67, 2; cvt.u16.u32 %rs16, %r60; shr.u16 %rs17, %rs16, 2; cvt.u16.u32 %rs18, %r68; shr.u16 %rs19, %rs18, 2; shl.b32 %r69, %r1, 1; add.s32 %r70, %r69, %r8; and.b32 %r71, %r70, -2; shr.s32 %r72, %r9, 1; add.s32 %r73, %r2, %r72; and.b16 %rs20, %rs17, 255; and.b16 %rs21, %rs19, 255; sust.b.2d.v2.b8.trap [%rd4, {%r71, %r73}], {%rs20, %rs21}; BB4_3: ret; } // .globl surf_copy_nv12_yuv420p10bit .visible .entry surf_copy_nv12_yuv420p10bit( .param .u64 surf_copy_nv12_yuv420p10bit_param_0, .param .u64 surf_copy_nv12_yuv420p10bit_param_1, .param .u64 surf_copy_nv12_yuv420p10bit_param_2, .param .u64 surf_copy_nv12_yuv420p10bit_param_3, .param .u64 surf_copy_nv12_yuv420p10bit_param_4, .param .u32 surf_copy_nv12_yuv420p10bit_param_5, .param .u32 surf_copy_nv12_yuv420p10bit_param_6 ) { .reg .pred %p<7>; .reg .b16 %rs<7>; .reg .b32 %r<15>; .reg .b64 %rd<6>; ld.param.u64 %rd1, [surf_copy_nv12_yuv420p10bit_param_0]; ld.param.u64 %rd2, [surf_copy_nv12_yuv420p10bit_param_1]; ld.param.u64 %rd3, [surf_copy_nv12_yuv420p10bit_param_2]; ld.param.u64 %rd4, [surf_copy_nv12_yuv420p10bit_param_3]; ld.param.u64 %rd5, [surf_copy_nv12_yuv420p10bit_param_4]; ld.param.u32 %r3, [surf_copy_nv12_yuv420p10bit_param_5]; ld.param.u32 %r4, [surf_copy_nv12_yuv420p10bit_param_6]; mov.u32 %r5, %ctaid.x; mov.u32 %r6, %ntid.x; mov.u32 %r7, %tid.x; mad.lo.s32 %r1, %r6, %r5, %r7; mov.u32 %r8, %ntid.y; mov.u32 %r9, %ctaid.y; mov.u32 %r10, %tid.y; mad.lo.s32 %r2, %r8, %r9, %r10; setp.ge.s32 %p1, %r2, %r4; setp.ge.s32 %p2, %r1, %r3; or.pred %p3, %p1, %p2; @%p3 bra BB5_3; shl.b32 %r11, %r1, 1; suld.b.2d.b16.trap {%rs1}, [%rd1, {%r11, %r2}]; shr.u16 %rs2, %rs1, 8; sust.b.2d.b8.trap [%rd3, {%r1, %r2}], {%rs2}; shr.s32 %r12, %r4, 1; setp.ge.s32 %p4, %r2, %r12; shr.s32 %r13, %r3, 1; setp.ge.s32 %p5, %r1, %r13; or.pred %p6, %p4, %p5; @%p6 bra BB5_3; shl.b32 %r14, %r1, 2; suld.b.2d.v2.b16.trap {%rs3, %rs4}, [%rd2, {%r14, %r2}]; shr.u16 %rs5, %rs3, 8; sust.b.2d.b8.trap [%rd4, {%r1, %r2}], {%rs5}; shr.u16 %rs6, %rs4, 8; sust.b.2d.b8.trap [%rd5, {%r1, %r2}], {%rs6}; BB5_3: ret; } // .globl surf_resize_nv12_yuv420p10bit .visible .entry surf_resize_nv12_yuv420p10bit( .param .u64 surf_resize_nv12_yuv420p10bit_param_0, .param .u64 surf_resize_nv12_yuv420p10bit_param_1, .param .u64 surf_resize_nv12_yuv420p10bit_param_2, .param .u64 surf_resize_nv12_yuv420p10bit_param_3, .param .u64 surf_resize_nv12_yuv420p10bit_param_4, .param .u32 surf_resize_nv12_yuv420p10bit_param_5, .param .u32 surf_resize_nv12_yuv420p10bit_param_6, .param .f32 surf_resize_nv12_yuv420p10bit_param_7, .param .f32 surf_resize_nv12_yuv420p10bit_param_8, .param .u32 surf_resize_nv12_yuv420p10bit_param_9, .param .u32 surf_resize_nv12_yuv420p10bit_param_10, .param .u32 surf_resize_nv12_yuv420p10bit_param_11, .param .u32 surf_resize_nv12_yuv420p10bit_param_12 ) { .reg .pred %p<7>; .reg .b16 %rs<19>; .reg .f32 %f<23>; .reg .b32 %r<65>; .reg .b64 %rd<6>; ld.param.u64 %rd1, [surf_resize_nv12_yuv420p10bit_param_0]; ld.param.u64 %rd2, [surf_resize_nv12_yuv420p10bit_param_1]; ld.param.u64 %rd3, [surf_resize_nv12_yuv420p10bit_param_2]; ld.param.u64 %rd4, [surf_resize_nv12_yuv420p10bit_param_3]; ld.param.u64 %rd5, [surf_resize_nv12_yuv420p10bit_param_4]; ld.param.u32 %r6, [surf_resize_nv12_yuv420p10bit_param_5]; ld.param.u32 %r7, [surf_resize_nv12_yuv420p10bit_param_6]; ld.param.f32 %f2, [surf_resize_nv12_yuv420p10bit_param_7]; ld.param.f32 %f3, [surf_resize_nv12_yuv420p10bit_param_8]; ld.param.u32 %r8, [surf_resize_nv12_yuv420p10bit_param_9]; ld.param.u32 %r9, [surf_resize_nv12_yuv420p10bit_param_10]; ld.param.u32 %r10, [surf_resize_nv12_yuv420p10bit_param_11]; ld.param.u32 %r11, [surf_resize_nv12_yuv420p10bit_param_12]; mov.u32 %r12, %ctaid.x; mov.u32 %r13, %ntid.x; mov.u32 %r14, %tid.x; mad.lo.s32 %r1, %r13, %r12, %r14; mov.u32 %r15, %ntid.y; mov.u32 %r16, %ctaid.y; mov.u32 %r17, %tid.y; mad.lo.s32 %r2, %r15, %r16, %r17; setp.ge.s32 %p1, %r2, %r7; setp.ge.s32 %p2, %r1, %r6; or.pred %p3, %p1, %p2; @%p3 bra BB6_3; cvt.rn.f32.s32 %f4, %r1; add.f32 %f5, %f4, 0f3F000000; mul.f32 %f6, %f5, %f2; cvt.rn.f32.s32 %f7, %r2; add.f32 %f8, %f7, 0f3F000000; mul.f32 %f9, %f8, %f3; add.f32 %f10, %f2, 0fBF800000; mul.f32 %f11, %f10, 0f3F000000; cvt.sat.f32.f32 %f12, %f11; add.f32 %f13, %f3, 0fBF800000; mul.f32 %f14, %f13, 0f3F000000; cvt.sat.f32.f32 %f15, %f14; add.f32 %f16, %f12, 0f3F000000; div.rn.f32 %f17, %f12, %f16; add.f32 %f18, %f15, 0f3F000000; div.rn.f32 %f19, %f15, %f18; sub.f32 %f20, %f6, %f17; cvt.rzi.s32.f32 %r3, %f20; add.f32 %f21, %f6, %f17; cvt.rzi.s32.f32 %r4, %f21; sub.f32 %f1, %f9, %f19; add.f32 %f22, %f9, %f19; cvt.rzi.s32.f32 %r5, %f22; add.s32 %r18, %r3, %r10; add.s32 %r19, %r5, %r11; shl.b32 %r20, %r18, 1; suld.b.2d.b16.trap {%rs1}, [%rd1, {%r20, %r19}]; suld.b.2d.b16.trap {%rs2}, [%rd1, {%r20, %r19}]; suld.b.2d.b16.trap {%rs3}, [%rd1, {%r20, %r19}]; suld.b.2d.b16.trap {%rs4}, [%rd1, {%r20, %r19}]; cvt.u32.u16 %r21, %rs1; cvt.u32.u16 %r22, %rs2; cvt.u32.u16 %r23, %rs3; cvt.u32.u16 %r24, %rs4; add.s32 %r25, %r21, %r22; add.s32 %r26, %r25, %r23; add.s32 %r27, %r26, %r24; add.s32 %r28, %r27, 2; shr.u32 %r29, %r28, 10; cvt.u16.u32 %rs5, %r29; and.b16 %rs6, %rs5, 255; add.s32 %r30, %r2, %r9; add.s32 %r31, %r1, %r8; sust.b.2d.b8.trap [%rd3, {%r31, %r30}], {%rs6}; shr.s32 %r32, %r7, 1; setp.ge.s32 %p4, %r2, %r32; shr.s32 %r33, %r6, 1; setp.ge.s32 %p5, %r1, %r33; or.pred %p6, %p4, %p5; @%p6 bra BB6_3; shr.s32 %r34, %r10, 1; shr.s32 %r35, %r11, 1; add.s32 %r36, %r5, %r35; cvt.rzi.s32.f32 %r37, %f1; add.s32 %r38, %r37, %r35; add.s32 %r39, %r3, %r34; shl.b32 %r40, %r39, 2; suld.b.2d.v2.b16.trap {%rs7, %rs8}, [%rd2, {%r40, %r38}]; add.s32 %r41, %r4, %r34; shl.b32 %r42, %r41, 2; suld.b.2d.v2.b16.trap {%rs9, %rs10}, [%rd2, {%r42, %r38}]; suld.b.2d.v2.b16.trap {%rs11, %rs12}, [%rd2, {%r40, %r36}]; suld.b.2d.v2.b16.trap {%rs13, %rs14}, [%rd2, {%r42, %r36}]; cvt.u32.u16 %r43, %rs7; cvt.u32.u16 %r44, %rs9; cvt.u32.u16 %r45, %rs11; cvt.u32.u16 %r46, %rs13; add.s32 %r47, %r43, %r44; add.s32 %r48, %r47, %r45; add.s32 %r49, %r48, %r46; add.s32 %r50, %r49, 2; shr.u32 %r51, %r50, 10; cvt.u32.u16 %r52, %rs8; cvt.u32.u16 %r53, %rs10; cvt.u32.u16 %r54, %rs12; cvt.u32.u16 %r55, %rs14; add.s32 %r56, %r52, %r53; add.s32 %r57, %r56, %r54; add.s32 %r58, %r57, %r55; add.s32 %r59, %r58, 2; shr.u32 %r60, %r59, 10; cvt.u16.u32 %rs15, %r51; shr.s32 %r61, %r8, 1; add.s32 %r62, %r1, %r61; shr.s32 %r63, %r9, 1; add.s32 %r64, %r2, %r63; and.b16 %rs16, %rs15, 255; sust.b.2d.b8.trap [%rd4, {%r62, %r64}], {%rs16}; cvt.u16.u32 %rs17, %r60; and.b16 %rs18, %rs17, 255; sust.b.2d.b8.trap [%rd5, {%r62, %r64}], {%rs18}; BB6_3: ret; } // .globl surf_resize_p016_nv12 .visible .entry surf_resize_p016_nv12( .param .u64 surf_resize_p016_nv12_param_0, .param .u64 surf_resize_p016_nv12_param_1, .param .u64 surf_resize_p016_nv12_param_2, .param .u64 surf_resize_p016_nv12_param_3, .param .u32 surf_resize_p016_nv12_param_4, .param .u32 surf_resize_p016_nv12_param_5, .param .f32 surf_resize_p016_nv12_param_6, .param .f32 surf_resize_p016_nv12_param_7, .param .u32 surf_resize_p016_nv12_param_8, .param .u32 surf_resize_p016_nv12_param_9, .param .u32 surf_resize_p016_nv12_param_10, .param .u32 surf_resize_p016_nv12_param_11 ) { .reg .pred %p<7>; .reg .b16 %rs<19>; .reg .f32 %f<23>; .reg .b32 %r<66>; .reg .b64 %rd<5>; ld.param.u64 %rd1, [surf_resize_p016_nv12_param_0]; ld.param.u64 %rd2, [surf_resize_p016_nv12_param_1]; ld.param.u64 %rd3, [surf_resize_p016_nv12_param_2]; ld.param.u64 %rd4, [surf_resize_p016_nv12_param_3]; ld.param.u32 %r6, [surf_resize_p016_nv12_param_4]; ld.param.u32 %r7, [surf_resize_p016_nv12_param_5]; ld.param.f32 %f2, [surf_resize_p016_nv12_param_6]; ld.param.f32 %f3, [surf_resize_p016_nv12_param_7]; ld.param.u32 %r8, [surf_resize_p016_nv12_param_8]; ld.param.u32 %r9, [surf_resize_p016_nv12_param_9]; ld.param.u32 %r10, [surf_resize_p016_nv12_param_10]; ld.param.u32 %r11, [surf_resize_p016_nv12_param_11]; mov.u32 %r12, %ctaid.x; mov.u32 %r13, %ntid.x; mov.u32 %r14, %tid.x; mad.lo.s32 %r1, %r13, %r12, %r14; mov.u32 %r15, %ntid.y; mov.u32 %r16, %ctaid.y; mov.u32 %r17, %tid.y; mad.lo.s32 %r2, %r15, %r16, %r17; setp.ge.s32 %p1, %r2, %r7; setp.ge.s32 %p2, %r1, %r6; or.pred %p3, %p1, %p2; @%p3 bra BB7_3; cvt.rn.f32.s32 %f4, %r1; add.f32 %f5, %f4, 0f3F000000; mul.f32 %f6, %f5, %f2; cvt.rn.f32.s32 %f7, %r2; add.f32 %f8, %f7, 0f3F000000; mul.f32 %f9, %f8, %f3; add.f32 %f10, %f2, 0fBF800000; mul.f32 %f11, %f10, 0f3F000000; cvt.sat.f32.f32 %f12, %f11; add.f32 %f13, %f3, 0fBF800000; mul.f32 %f14, %f13, 0f3F000000; cvt.sat.f32.f32 %f15, %f14; add.f32 %f16, %f12, 0f3F000000; div.rn.f32 %f17, %f12, %f16; add.f32 %f18, %f15, 0f3F000000; div.rn.f32 %f19, %f15, %f18; sub.f32 %f20, %f6, %f17; cvt.rzi.s32.f32 %r3, %f20; add.f32 %f21, %f6, %f17; cvt.rzi.s32.f32 %r4, %f21; sub.f32 %f1, %f9, %f19; add.f32 %f22, %f9, %f19; cvt.rzi.s32.f32 %r5, %f22; add.s32 %r18, %r3, %r10; add.s32 %r19, %r5, %r11; shl.b32 %r20, %r18, 1; suld.b.2d.b16.trap {%rs1}, [%rd1, {%r20, %r19}]; suld.b.2d.b16.trap {%rs2}, [%rd1, {%r20, %r19}]; suld.b.2d.b16.trap {%rs3}, [%rd1, {%r20, %r19}]; suld.b.2d.b16.trap {%rs4}, [%rd1, {%r20, %r19}]; cvt.u32.u16 %r21, %rs1; cvt.u32.u16 %r22, %rs2; cvt.u32.u16 %r23, %rs3; cvt.u32.u16 %r24, %rs4; add.s32 %r25, %r21, %r22; add.s32 %r26, %r25, %r23; add.s32 %r27, %r26, %r24; add.s32 %r28, %r27, 2; shr.u32 %r29, %r28, 10; cvt.u16.u32 %rs5, %r29; and.b16 %rs6, %rs5, 255; add.s32 %r30, %r2, %r9; add.s32 %r31, %r1, %r8; sust.b.2d.b8.trap [%rd3, {%r31, %r30}], {%rs6}; shr.s32 %r32, %r7, 1; setp.ge.s32 %p4, %r2, %r32; shr.s32 %r33, %r6, 1; setp.ge.s32 %p5, %r1, %r33; or.pred %p6, %p4, %p5; @%p6 bra BB7_3; shr.s32 %r34, %r10, 1; shr.s32 %r35, %r11, 1; add.s32 %r36, %r5, %r35; cvt.rzi.s32.f32 %r37, %f1; add.s32 %r38, %r37, %r35; add.s32 %r39, %r3, %r34; shl.b32 %r40, %r39, 2; suld.b.2d.v2.b16.trap {%rs7, %rs8}, [%rd2, {%r40, %r38}]; add.s32 %r41, %r4, %r34; shl.b32 %r42, %r41, 2; suld.b.2d.v2.b16.trap {%rs9, %rs10}, [%rd2, {%r42, %r38}]; suld.b.2d.v2.b16.trap {%rs11, %rs12}, [%rd2, {%r40, %r36}]; suld.b.2d.v2.b16.trap {%rs13, %rs14}, [%rd2, {%r42, %r36}]; cvt.u32.u16 %r43, %rs7; cvt.u32.u16 %r44, %rs9; cvt.u32.u16 %r45, %rs11; cvt.u32.u16 %r46, %rs13; add.s32 %r47, %r43, %r44; add.s32 %r48, %r47, %r45; add.s32 %r49, %r48, %r46; add.s32 %r50, %r49, 2; shr.u32 %r51, %r50, 10; cvt.u32.u16 %r52, %rs8; cvt.u32.u16 %r53, %rs10; cvt.u32.u16 %r54, %rs12; cvt.u32.u16 %r55, %rs14; add.s32 %r56, %r52, %r53; add.s32 %r57, %r56, %r54; add.s32 %r58, %r57, %r55; add.s32 %r59, %r58, 2; shr.u32 %r60, %r59, 10; cvt.u16.u32 %rs15, %r51; cvt.u16.u32 %rs16, %r60; shl.b32 %r61, %r1, 1; add.s32 %r62, %r61, %r8; and.b32 %r63, %r62, -2; shr.s32 %r64, %r9, 1; add.s32 %r65, %r2, %r64; and.b16 %rs17, %rs15, 255; and.b16 %rs18, %rs16, 255; sust.b.2d.v2.b8.trap [%rd4, {%r63, %r65}], {%rs17, %rs18}; BB7_3: ret; } // .globl surf_resize_p016 .visible .entry surf_resize_p016( .param .u64 surf_resize_p016_param_0, .param .u64 surf_resize_p016_param_1, .param .u64 surf_resize_p016_param_2, .param .u64 surf_resize_p016_param_3, .param .u32 surf_resize_p016_param_4, .param .u32 surf_resize_p016_param_5, .param .f32 surf_resize_p016_param_6, .param .f32 surf_resize_p016_param_7, .param .u32 surf_resize_p016_param_8, .param .u32 surf_resize_p016_param_9, .param .u32 surf_resize_p016_param_10, .param .u32 surf_resize_p016_param_11 ) { .reg .pred %p<7>; .reg .b16 %rs<16>; .reg .f32 %f<23>; .reg .b32 %r<67>; .reg .b64 %rd<5>; ld.param.u64 %rd1, [surf_resize_p016_param_0]; ld.param.u64 %rd2, [surf_resize_p016_param_1]; ld.param.u64 %rd3, [surf_resize_p016_param_2]; ld.param.u64 %rd4, [surf_resize_p016_param_3]; ld.param.u32 %r6, [surf_resize_p016_param_4]; ld.param.u32 %r7, [surf_resize_p016_param_5]; ld.param.f32 %f2, [surf_resize_p016_param_6]; ld.param.f32 %f3, [surf_resize_p016_param_7]; ld.param.u32 %r8, [surf_resize_p016_param_8]; ld.param.u32 %r9, [surf_resize_p016_param_9]; ld.param.u32 %r10, [surf_resize_p016_param_10]; ld.param.u32 %r11, [surf_resize_p016_param_11]; mov.u32 %r12, %ctaid.x; mov.u32 %r13, %ntid.x; mov.u32 %r14, %tid.x; mad.lo.s32 %r1, %r13, %r12, %r14; mov.u32 %r15, %ntid.y; mov.u32 %r16, %ctaid.y; mov.u32 %r17, %tid.y; mad.lo.s32 %r2, %r15, %r16, %r17; setp.ge.s32 %p1, %r2, %r7; setp.ge.s32 %p2, %r1, %r6; or.pred %p3, %p1, %p2; @%p3 bra BB8_3; cvt.rn.f32.s32 %f4, %r1; add.f32 %f5, %f4, 0f3F000000; mul.f32 %f6, %f5, %f2; cvt.rn.f32.s32 %f7, %r2; add.f32 %f8, %f7, 0f3F000000; mul.f32 %f9, %f8, %f3; add.f32 %f10, %f2, 0fBF800000; mul.f32 %f11, %f10, 0f3F000000; cvt.sat.f32.f32 %f12, %f11; add.f32 %f13, %f3, 0fBF800000; mul.f32 %f14, %f13, 0f3F000000; cvt.sat.f32.f32 %f15, %f14; add.f32 %f16, %f12, 0f3F000000; div.rn.f32 %f17, %f12, %f16; add.f32 %f18, %f15, 0f3F000000; div.rn.f32 %f19, %f15, %f18; sub.f32 %f20, %f6, %f17; cvt.rzi.s32.f32 %r3, %f20; add.f32 %f21, %f6, %f17; cvt.rzi.s32.f32 %r4, %f21; sub.f32 %f1, %f9, %f19; add.f32 %f22, %f9, %f19; cvt.rzi.s32.f32 %r5, %f22; add.s32 %r18, %r3, %r10; add.s32 %r19, %r5, %r11; shl.b32 %r20, %r18, 1; suld.b.2d.b16.trap {%rs1}, [%rd1, {%r20, %r19}]; suld.b.2d.b16.trap {%rs2}, [%rd1, {%r20, %r19}]; suld.b.2d.b16.trap {%rs3}, [%rd1, {%r20, %r19}]; suld.b.2d.b16.trap {%rs4}, [%rd1, {%r20, %r19}]; cvt.u32.u16 %r21, %rs1; cvt.u32.u16 %r22, %rs2; cvt.u32.u16 %r23, %rs3; cvt.u32.u16 %r24, %rs4; add.s32 %r25, %r21, %r22; add.s32 %r26, %r25, %r23; add.s32 %r27, %r26, %r24; add.s32 %r28, %r27, 2; shr.u32 %r29, %r28, 2; cvt.u16.u32 %rs5, %r29; add.s32 %r30, %r1, %r8; shl.b32 %r31, %r30, 1; add.s32 %r32, %r2, %r9; sust.b.2d.b16.trap [%rd3, {%r31, %r32}], {%rs5}; shr.s32 %r33, %r7, 1; setp.ge.s32 %p4, %r2, %r33; shr.s32 %r34, %r6, 1; setp.ge.s32 %p5, %r1, %r34; or.pred %p6, %p4, %p5; @%p6 bra BB8_3; shr.s32 %r35, %r10, 1; shr.s32 %r36, %r11, 1; add.s32 %r37, %r5, %r36; cvt.rzi.s32.f32 %r38, %f1; add.s32 %r39, %r38, %r36; add.s32 %r40, %r3, %r35; shl.b32 %r41, %r40, 2; suld.b.2d.v2.b16.trap {%rs6, %rs7}, [%rd2, {%r41, %r39}]; add.s32 %r42, %r4, %r35; shl.b32 %r43, %r42, 2; suld.b.2d.v2.b16.trap {%rs8, %rs9}, [%rd2, {%r43, %r39}]; suld.b.2d.v2.b16.trap {%rs10, %rs11}, [%rd2, {%r41, %r37}]; suld.b.2d.v2.b16.trap {%rs12, %rs13}, [%rd2, {%r43, %r37}]; cvt.u32.u16 %r44, %rs6; cvt.u32.u16 %r45, %rs8; cvt.u32.u16 %r46, %rs10; cvt.u32.u16 %r47, %rs12; add.s32 %r48, %r44, %r45; add.s32 %r49, %r48, %r46; add.s32 %r50, %r49, %r47; add.s32 %r51, %r50, 2; shr.u32 %r52, %r51, 2; cvt.u32.u16 %r53, %rs7; cvt.u32.u16 %r54, %rs9; cvt.u32.u16 %r55, %rs11; cvt.u32.u16 %r56, %rs13; add.s32 %r57, %r53, %r54; add.s32 %r58, %r57, %r55; add.s32 %r59, %r58, %r56; add.s32 %r60, %r59, 2; shr.u32 %r61, %r60, 2; cvt.u16.u32 %rs14, %r52; cvt.u16.u32 %rs15, %r61; shr.u32 %r62, %r8, 1; add.s32 %r63, %r1, %r62; shl.b32 %r64, %r63, 2; shr.s32 %r65, %r9, 1; add.s32 %r66, %r2, %r65; sust.b.2d.v2.b16.trap [%rd4, {%r64, %r66}], {%rs14, %rs15}; BB8_3: ret; } // .globl surf_copy_p016 .visible .entry surf_copy_p016( .param .u64 surf_copy_p016_param_0, .param .u64 surf_copy_p016_param_1, .param .u64 surf_copy_p016_param_2, .param .u64 surf_copy_p016_param_3, .param .u64 surf_copy_p016_param_4, .param .u32 surf_copy_p016_param_5, .param .u32 surf_copy_p016_param_6 ) { .reg .pred %p<7>; .reg .b16 %rs<4>; .reg .b32 %r<15>; .reg .b64 %rd<5>; ld.param.u64 %rd1, [surf_copy_p016_param_0]; ld.param.u64 %rd2, [surf_copy_p016_param_1]; ld.param.u64 %rd3, [surf_copy_p016_param_2]; ld.param.u64 %rd4, [surf_copy_p016_param_3]; ld.param.u32 %r3, [surf_copy_p016_param_5]; ld.param.u32 %r4, [surf_copy_p016_param_6]; mov.u32 %r5, %ctaid.x; mov.u32 %r6, %ntid.x; mov.u32 %r7, %tid.x; mad.lo.s32 %r1, %r6, %r5, %r7; mov.u32 %r8, %ntid.y; mov.u32 %r9, %ctaid.y; mov.u32 %r10, %tid.y; mad.lo.s32 %r2, %r8, %r9, %r10; setp.ge.s32 %p1, %r2, %r4; setp.ge.s32 %p2, %r1, %r3; or.pred %p3, %p1, %p2; @%p3 bra BB9_3; shl.b32 %r11, %r1, 1; suld.b.2d.b16.trap {%rs1}, [%rd1, {%r11, %r2}]; sust.b.2d.b16.trap [%rd3, {%r11, %r2}], {%rs1}; shr.s32 %r12, %r4, 1; setp.ge.s32 %p4, %r2, %r12; shr.s32 %r13, %r3, 1; setp.ge.s32 %p5, %r1, %r13; or.pred %p6, %p4, %p5; @%p6 bra BB9_3; shl.b32 %r14, %r1, 2; suld.b.2d.v2.b16.trap {%rs2, %rs3}, [%rd2, {%r14, %r2}]; sust.b.2d.v2.b16.trap [%rd4, {%r14, %r2}], {%rs2, %rs3}; BB9_3: ret; } // .globl surf_copy_nv12 .visible .entry surf_copy_nv12( .param .u64 surf_copy_nv12_param_0, .param .u64 surf_copy_nv12_param_1, .param .u64 surf_copy_nv12_param_2, .param .u64 surf_copy_nv12_param_3, .param .u64 surf_copy_nv12_param_4, .param .u32 surf_copy_nv12_param_5, .param .u32 surf_copy_nv12_param_6 ) { .reg .pred %p<7>; .reg .b16 %rs<7>; .reg .b32 %r<14>; .reg .b64 %rd<5>; ld.param.u64 %rd1, [surf_copy_nv12_param_0]; ld.param.u64 %rd2, [surf_copy_nv12_param_1]; ld.param.u64 %rd3, [surf_copy_nv12_param_2]; ld.param.u64 %rd4, [surf_copy_nv12_param_3]; ld.param.u32 %r3, [surf_copy_nv12_param_5]; ld.param.u32 %r4, [surf_copy_nv12_param_6]; mov.u32 %r5, %ctaid.x; mov.u32 %r6, %ntid.x; mov.u32 %r7, %tid.x; mad.lo.s32 %r1, %r6, %r5, %r7; mov.u32 %r8, %ntid.y; mov.u32 %r9, %ctaid.y; mov.u32 %r10, %tid.y; mad.lo.s32 %r2, %r8, %r9, %r10; setp.ge.s32 %p1, %r2, %r4; setp.ge.s32 %p2, %r1, %r3; or.pred %p3, %p1, %p2; @%p3 bra BB10_3; suld.b.2d.b8.trap {%rs1}, [%rd1, {%r1, %r2}]; and.b16 %rs2, %rs1, 255; sust.b.2d.b8.trap [%rd3, {%r1, %r2}], {%rs2}; shr.s32 %r11, %r4, 1; setp.ge.s32 %p4, %r2, %r11; shr.s32 %r12, %r3, 1; setp.ge.s32 %p5, %r1, %r12; or.pred %p6, %p4, %p5; @%p6 bra BB10_3; shl.b32 %r13, %r1, 1; suld.b.2d.v2.b8.trap {%rs3, %rs4}, [%rd2, {%r13, %r2}]; and.b16 %rs5, %rs3, 255; and.b16 %rs6, %rs4, 255; sust.b.2d.v2.b8.trap [%rd4, {%r13, %r2}], {%rs5, %rs6}; BB10_3: ret; } // .globl surf_resize_nv12_p016 .visible .entry surf_resize_nv12_p016( .param .u64 surf_resize_nv12_p016_param_0, .param .u64 surf_resize_nv12_p016_param_1, .param .u64 surf_resize_nv12_p016_param_2, .param .u64 surf_resize_nv12_p016_param_3, .param .u32 surf_resize_nv12_p016_param_4, .param .u32 surf_resize_nv12_p016_param_5, .param .f32 surf_resize_nv12_p016_param_6, .param .f32 surf_resize_nv12_p016_param_7, .param .u32 surf_resize_nv12_p016_param_8, .param .u32 surf_resize_nv12_p016_param_9, .param .u32 surf_resize_nv12_p016_param_10, .param .u32 surf_resize_nv12_p016_param_11 ) { .reg .pred %p<7>; .reg .b16 %rs<16>; .reg .f32 %f<23>; .reg .b32 %r<78>; .reg .b64 %rd<5>; ld.param.u64 %rd1, [surf_resize_nv12_p016_param_0]; ld.param.u64 %rd2, [surf_resize_nv12_p016_param_1]; ld.param.u64 %rd3, [surf_resize_nv12_p016_param_2]; ld.param.u64 %rd4, [surf_resize_nv12_p016_param_3]; ld.param.u32 %r6, [surf_resize_nv12_p016_param_4]; ld.param.u32 %r7, [surf_resize_nv12_p016_param_5]; ld.param.f32 %f2, [surf_resize_nv12_p016_param_6]; ld.param.f32 %f3, [surf_resize_nv12_p016_param_7]; ld.param.u32 %r8, [surf_resize_nv12_p016_param_8]; ld.param.u32 %r9, [surf_resize_nv12_p016_param_9]; ld.param.u32 %r10, [surf_resize_nv12_p016_param_10]; ld.param.u32 %r11, [surf_resize_nv12_p016_param_11]; mov.u32 %r12, %ctaid.x; mov.u32 %r13, %ntid.x; mov.u32 %r14, %tid.x; mad.lo.s32 %r1, %r13, %r12, %r14; mov.u32 %r15, %ntid.y; mov.u32 %r16, %ctaid.y; mov.u32 %r17, %tid.y; mad.lo.s32 %r2, %r15, %r16, %r17; setp.ge.s32 %p1, %r2, %r7; setp.ge.s32 %p2, %r1, %r6; or.pred %p3, %p1, %p2; @%p3 bra BB11_3; cvt.rn.f32.s32 %f4, %r1; add.f32 %f5, %f4, 0f3F000000; mul.f32 %f6, %f5, %f2; cvt.rn.f32.s32 %f7, %r2; add.f32 %f8, %f7, 0f3F000000; mul.f32 %f9, %f8, %f3; add.f32 %f10, %f2, 0fBF800000; mul.f32 %f11, %f10, 0f3F000000; cvt.sat.f32.f32 %f12, %f11; add.f32 %f13, %f3, 0fBF800000; mul.f32 %f14, %f13, 0f3F000000; cvt.sat.f32.f32 %f15, %f14; add.f32 %f16, %f12, 0f3F000000; div.rn.f32 %f17, %f12, %f16; add.f32 %f18, %f15, 0f3F000000; div.rn.f32 %f19, %f15, %f18; sub.f32 %f20, %f6, %f17; cvt.rzi.s32.f32 %r3, %f20; add.f32 %f21, %f6, %f17; cvt.rzi.s32.f32 %r4, %f21; sub.f32 %f1, %f9, %f19; add.f32 %f22, %f9, %f19; cvt.rzi.s32.f32 %r5, %f22; add.s32 %r18, %r3, %r10; add.s32 %r19, %r5, %r11; suld.b.2d.b8.trap {%rs1}, [%rd1, {%r18, %r19}]; cvt.u32.u16 %r20, %rs1; suld.b.2d.b8.trap {%rs2}, [%rd1, {%r18, %r19}]; cvt.u32.u16 %r21, %rs2; suld.b.2d.b8.trap {%rs3}, [%rd1, {%r18, %r19}]; cvt.u32.u16 %r22, %rs3; suld.b.2d.b8.trap {%rs4}, [%rd1, {%r18, %r19}]; cvt.u32.u16 %r23, %rs4; and.b32 %r24, %r20, 255; and.b32 %r25, %r21, 255; add.s32 %r26, %r25, %r24; and.b32 %r27, %r22, 255; add.s32 %r28, %r26, %r27; and.b32 %r29, %r23, 255; add.s32 %r30, %r28, %r29; shl.b32 %r31, %r30, 6; add.s32 %r32, %r31, 128; cvt.u16.u32 %rs5, %r32; add.s32 %r33, %r1, %r8; shl.b32 %r34, %r33, 1; add.s32 %r35, %r2, %r9; sust.b.2d.b16.trap [%rd3, {%r34, %r35}], {%rs5}; shr.s32 %r36, %r7, 1; setp.ge.s32 %p4, %r2, %r36; shr.s32 %r37, %r6, 1; setp.ge.s32 %p5, %r1, %r37; or.pred %p6, %p4, %p5; @%p6 bra BB11_3; shr.s32 %r38, %r10, 1; shr.s32 %r39, %r11, 1; add.s32 %r40, %r5, %r39; cvt.rzi.s32.f32 %r41, %f1; add.s32 %r42, %r41, %r39; add.s32 %r43, %r3, %r38; shl.b32 %r44, %r43, 1; suld.b.2d.v2.b8.trap {%rs6, %rs7}, [%rd2, {%r44, %r42}]; cvt.u32.u16 %r45, %rs6; cvt.u32.u16 %r46, %rs7; add.s32 %r47, %r4, %r38; shl.b32 %r48, %r47, 1; suld.b.2d.v2.b8.trap {%rs8, %rs9}, [%rd2, {%r48, %r42}]; cvt.u32.u16 %r49, %rs8; cvt.u32.u16 %r50, %rs9; suld.b.2d.v2.b8.trap {%rs10, %rs11}, [%rd2, {%r44, %r40}]; cvt.u32.u16 %r51, %rs10; cvt.u32.u16 %r52, %rs11; suld.b.2d.v2.b8.trap {%rs12, %rs13}, [%rd2, {%r48, %r40}]; cvt.u32.u16 %r53, %rs12; cvt.u32.u16 %r54, %rs13; and.b32 %r55, %r45, 255; and.b32 %r56, %r49, 255; add.s32 %r57, %r56, %r55; and.b32 %r58, %r51, 255; add.s32 %r59, %r57, %r58; and.b32 %r60, %r53, 255; add.s32 %r61, %r59, %r60; and.b32 %r62, %r46, 255; and.b32 %r63, %r50, 255; add.s32 %r64, %r63, %r62; and.b32 %r65, %r52, 255; add.s32 %r66, %r64, %r65; and.b32 %r67, %r54, 255; add.s32 %r68, %r66, %r67; shl.b32 %r69, %r61, 6; add.s32 %r70, %r69, 128; cvt.u16.u32 %rs14, %r70; shl.b32 %r71, %r68, 6; add.s32 %r72, %r71, 128; cvt.u16.u32 %rs15, %r72; shl.b32 %r73, %r1, 2; add.s32 %r74, %r73, %r8; and.b32 %r75, %r74, -4; shr.s32 %r76, %r9, 1; add.s32 %r77, %r2, %r76; sust.b.2d.v2.b16.trap [%rd4, {%r75, %r77}], {%rs14, %rs15}; BB11_3: ret; } // .globl surf_copy_nv12_p016 .visible .entry surf_copy_nv12_p016( .param .u64 surf_copy_nv12_p016_param_0, .param .u64 surf_copy_nv12_p016_param_1, .param .u64 surf_copy_nv12_p016_param_2, .param .u64 surf_copy_nv12_p016_param_3, .param .u64 surf_copy_nv12_p016_param_4, .param .u32 surf_copy_nv12_p016_param_5, .param .u32 surf_copy_nv12_p016_param_6 ) { .reg .pred %p<7>; .reg .b16 %rs<7>; .reg .b32 %r<15>; .reg .b64 %rd<5>; ld.param.u64 %rd1, [surf_copy_nv12_p016_param_0]; ld.param.u64 %rd2, [surf_copy_nv12_p016_param_1]; ld.param.u64 %rd3, [surf_copy_nv12_p016_param_2]; ld.param.u64 %rd4, [surf_copy_nv12_p016_param_3]; ld.param.u32 %r4, [surf_copy_nv12_p016_param_5]; ld.param.u32 %r5, [surf_copy_nv12_p016_param_6]; mov.u32 %r6, %ctaid.x; mov.u32 %r7, %ntid.x; mov.u32 %r8, %tid.x; mad.lo.s32 %r1, %r7, %r6, %r8; mov.u32 %r9, %ntid.y; mov.u32 %r10, %ctaid.y; mov.u32 %r11, %tid.y; mad.lo.s32 %r2, %r9, %r10, %r11; setp.ge.s32 %p1, %r2, %r5; setp.ge.s32 %p2, %r1, %r4; or.pred %p3, %p1, %p2; @%p3 bra BB12_3; suld.b.2d.b8.trap {%rs1}, [%rd1, {%r1, %r2}]; shl.b16 %rs2, %rs1, 8; shl.b32 %r3, %r1, 1; sust.b.2d.b16.trap [%rd3, {%r3, %r2}], {%rs2}; shr.s32 %r12, %r5, 1; setp.ge.s32 %p4, %r2, %r12; shr.s32 %r13, %r4, 1; setp.ge.s32 %p5, %r1, %r13; or.pred %p6, %p4, %p5; @%p6 bra BB12_3; suld.b.2d.v2.b8.trap {%rs3, %rs4}, [%rd2, {%r3, %r2}]; shl.b16 %rs5, %rs3, 8; shl.b16 %rs6, %rs4, 8; shl.b32 %r14, %r1, 2; sust.b.2d.v2.b16.trap [%rd4, {%r14, %r2}], {%rs5, %rs6}; BB12_3: ret; } // .globl surf_fill_p016 .visible .entry surf_fill_p016( .param .u64 surf_fill_p016_param_0, .param .u64 surf_fill_p016_param_1, .param .u32 surf_fill_p016_param_2, .param .u32 surf_fill_p016_param_3, .param .u8 surf_fill_p016_param_4, .param .u8 surf_fill_p016_param_5, .param .u8 surf_fill_p016_param_6 ) { .reg .pred %p<7>; .reg .b16 %rs<7>; .reg .b32 %r<15>; .reg .b64 %rd<3>; ld.param.u64 %rd1, [surf_fill_p016_param_0]; ld.param.u64 %rd2, [surf_fill_p016_param_1]; ld.param.u32 %r3, [surf_fill_p016_param_2]; ld.param.u32 %r4, [surf_fill_p016_param_3]; ld.param.u8 %rs3, [surf_fill_p016_param_6]; ld.param.u8 %rs2, [surf_fill_p016_param_5]; ld.param.u8 %rs1, [surf_fill_p016_param_4]; mov.u32 %r5, %ctaid.x; mov.u32 %r6, %ntid.x; mov.u32 %r7, %tid.x; mad.lo.s32 %r1, %r6, %r5, %r7; mov.u32 %r8, %ntid.y; mov.u32 %r9, %ctaid.y; mov.u32 %r10, %tid.y; mad.lo.s32 %r2, %r8, %r9, %r10; setp.ge.s32 %p1, %r2, %r4; setp.ge.s32 %p2, %r1, %r3; or.pred %p3, %p1, %p2; @%p3 bra BB13_3; shl.b16 %rs4, %rs1, 8; shl.b32 %r11, %r1, 1; sust.b.2d.b16.trap [%rd1, {%r11, %r2}], {%rs4}; shr.s32 %r12, %r4, 1; setp.ge.s32 %p4, %r2, %r12; shr.s32 %r13, %r3, 1; setp.ge.s32 %p5, %r1, %r13; or.pred %p6, %p4, %p5; @%p6 bra BB13_3; shl.b16 %rs5, %rs2, 8; shl.b16 %rs6, %rs3, 8; shl.b32 %r14, %r1, 2; sust.b.2d.v2.b16.trap [%rd2, {%r14, %r2}], {%rs5, %rs6}; BB13_3: ret; } // .globl surf_copy_p016_nv12 .visible .entry surf_copy_p016_nv12( .param .u64 surf_copy_p016_nv12_param_0, .param .u64 surf_copy_p016_nv12_param_1, .param .u64 surf_copy_p016_nv12_param_2, .param .u64 surf_copy_p016_nv12_param_3, .param .u64 surf_copy_p016_nv12_param_4, .param .u32 surf_copy_p016_nv12_param_5, .param .u32 surf_copy_p016_nv12_param_6 ) { .reg .pred %p<7>; .reg .b16 %rs<7>; .reg .b32 %r<15>; .reg .b64 %rd<5>; ld.param.u64 %rd1, [surf_copy_p016_nv12_param_0]; ld.param.u64 %rd2, [surf_copy_p016_nv12_param_1]; ld.param.u64 %rd3, [surf_copy_p016_nv12_param_2]; ld.param.u64 %rd4, [surf_copy_p016_nv12_param_3]; ld.param.u32 %r4, [surf_copy_p016_nv12_param_5]; ld.param.u32 %r5, [surf_copy_p016_nv12_param_6]; mov.u32 %r6, %ctaid.x; mov.u32 %r7, %ntid.x; mov.u32 %r8, %tid.x; mad.lo.s32 %r1, %r7, %r6, %r8; mov.u32 %r9, %ntid.y; mov.u32 %r10, %ctaid.y; mov.u32 %r11, %tid.y; mad.lo.s32 %r2, %r9, %r10, %r11; setp.ge.s32 %p1, %r2, %r5; setp.ge.s32 %p2, %r1, %r4; or.pred %p3, %p1, %p2; @%p3 bra BB14_3; shl.b32 %r3, %r1, 1; suld.b.2d.b16.trap {%rs1}, [%rd1, {%r3, %r2}]; shr.u16 %rs2, %rs1, 8; sust.b.2d.b8.trap [%rd3, {%r1, %r2}], {%rs2}; shr.s32 %r12, %r5, 1; setp.ge.s32 %p4, %r2, %r12; shr.s32 %r13, %r4, 1; setp.ge.s32 %p5, %r1, %r13; or.pred %p6, %p4, %p5; @%p6 bra BB14_3; shl.b32 %r14, %r1, 2; suld.b.2d.v2.b16.trap {%rs3, %rs4}, [%rd2, {%r14, %r2}]; shr.u16 %rs5, %rs3, 8; shr.u16 %rs6, %rs4, 8; sust.b.2d.v2.b8.trap [%rd4, {%r3, %r2}], {%rs5, %rs6}; BB14_3: ret; } // // Generated by NVIDIA NVVM Compiler // // Compiler Build ID: CL-29663091 // Cuda compilation tools, release 10.2, V10.2.300 // Based on LLVM 3.4svn // .version 6.5 .target sm_30 .address_size 64 // .globl blend_rgba_nv12 .global .align 4 .b8 apix[4] = {16, 128, 128, 0}; .visible .entry blend_rgba_nv12( .param .u64 blend_rgba_nv12_param_0, .param .u32 blend_rgba_nv12_param_1, .param .u32 blend_rgba_nv12_param_2, .param .u32 blend_rgba_nv12_param_3, .param .u32 blend_rgba_nv12_param_4, .param .u64 blend_rgba_nv12_param_5, .param .u64 blend_rgba_nv12_param_6, .param .u64 blend_rgba_nv12_param_7, .param .u32 blend_rgba_nv12_param_8, .param .u32 blend_rgba_nv12_param_9 ) { .reg .pred %p<10>; .reg .b16 %rs<66>; .reg .b32 %r<69>; .reg .b64 %rd<4>; ld.param.u64 %rd1, [blend_rgba_nv12_param_0]; ld.param.u32 %r5, [blend_rgba_nv12_param_1]; ld.param.u32 %r6, [blend_rgba_nv12_param_2]; ld.param.u32 %r7, [blend_rgba_nv12_param_3]; ld.param.u32 %r8, [blend_rgba_nv12_param_4]; ld.param.u64 %rd2, [blend_rgba_nv12_param_5]; ld.param.u64 %rd3, [blend_rgba_nv12_param_6]; mov.u32 %r9, %ctaid.x; mov.u32 %r10, %ntid.x; mov.u32 %r11, %tid.x; mad.lo.s32 %r1, %r10, %r9, %r11; mov.u32 %r12, %ntid.y; mov.u32 %r13, %ctaid.y; mov.u32 %r14, %tid.y; mad.lo.s32 %r2, %r12, %r13, %r14; setp.ge.s32 %p1, %r1, %r7; setp.ge.s32 %p2, %r2, %r8; or.pred %p3, %p1, %p2; @%p3 bra BB0_7; and.b32 %r3, %r5, 1; setp.eq.b32 %p4, %r3, 1; setp.eq.s32 %p5, %r1, 0; and.pred %p6, %p5, %p4; @%p6 bra BB0_5; bra.uni BB0_2; BB0_5: ld.global.v4.u8 {%rs62, %rs63, %rs64, %rs65}, [apix]; bra.uni BB0_6; BB0_2: setp.ne.s32 %p7, %r3, 0; shl.b32 %r4, %r1, 2; @%p7 bra BB0_4; bra.uni BB0_3; BB0_4: add.s32 %r15, %r4, -4; suld.b.2d.v4.b8.trap {%rs62, %rs63, %rs64, %rs65}, [%rd1, {%r15, %r2}]; bra.uni BB0_6; BB0_3: suld.b.2d.v4.b8.trap {%rs62, %rs63, %rs64, %rs65}, [%rd1, {%r4, %r2}]; BB0_6: setp.ne.s32 %p8, %r3, 0; selp.b32 %r16, -1, 0, %p8; add.s32 %r17, %r1, %r5; add.s32 %r18, %r17, %r16; and.b32 %r19, %r18, -2; and.b32 %r20, %r6, 1; setp.eq.b32 %p9, %r20, 1; selp.b32 %r21, -1, 0, %p9; add.s32 %r22, %r2, %r6; add.s32 %r23, %r22, %r21; shr.s32 %r24, %r23, 1; suld.b.2d.b8.trap {%rs38}, [%rd2, {%r17, %r22}]; cvt.u32.u16 %r25, %rs38; cvt.u32.u16 %r26, %rs62; and.b32 %r27, %r26, 255; cvt.u32.u16 %r28, %rs63; and.b32 %r29, %r28, 255; and.b16 %rs39, %rs63, 255; mul.wide.u16 %r30, %rs39, 129; cvt.u32.u16 %r31, %rs64; and.b32 %r32, %r31, 255; mad.lo.s32 %r33, %r27, 66, %r30; add.s32 %r34, %r33, 128; mad.lo.s32 %r35, %r32, 25, %r34; shr.u32 %r36, %r35, 8; add.s32 %r37, %r36, 16; cvt.u32.u16 %r38, %rs65; and.b32 %r39, %r38, 255; add.s32 %r40, %r39, 1; and.b32 %r41, %r37, 255; mov.u32 %r42, 256; sub.s32 %r43, %r42, %r39; and.b32 %r44, %r25, 255; mul.lo.s32 %r45, %r44, %r43; mad.lo.s32 %r46, %r41, %r40, %r45; cvt.u16.u32 %rs40, %r46; shr.u16 %rs41, %rs40, 8; sust.b.2d.b8.trap [%rd2, {%r17, %r22}], {%rs41}; suld.b.2d.v2.b8.trap {%rs42, %rs43}, [%rd3, {%r19, %r24}]; cvt.u32.u16 %r47, %rs42; cvt.u32.u16 %r48, %rs43; mul.lo.s32 %r49, %r29, -74; mad.lo.s32 %r50, %r27, -38, %r49; add.s32 %r51, %r50, 128; mad.lo.s32 %r52, %r32, 112, %r51; shr.u32 %r53, %r52, 8; add.s32 %r54, %r53, 128; and.b32 %r55, %r54, 255; mul.lo.s32 %r56, %r55, %r40; and.b32 %r57, %r47, 255; mad.lo.s32 %r58, %r57, %r43, %r56; cvt.u16.u32 %rs44, %r58; mul.lo.s32 %r59, %r29, -94; mad.lo.s32 %r60, %r27, 112, %r59; add.s32 %r61, %r60, 128; mad.lo.s32 %r62, %r32, -18, %r61; shr.u32 %r63, %r62, 8; add.s32 %r64, %r63, 128; and.b32 %r65, %r64, 255; mul.lo.s32 %r66, %r65, %r40; and.b32 %r67, %r48, 255; mad.lo.s32 %r68, %r67, %r43, %r66; cvt.u16.u32 %rs45, %r68; shr.u16 %rs46, %rs44, 8; shr.u16 %rs47, %rs45, 8; sust.b.2d.v2.b8.trap [%rd3, {%r19, %r24}], {%rs46, %rs47}; BB0_7: ret; } // .globl blend_rgba_p016 .visible .entry blend_rgba_p016( .param .u64 blend_rgba_p016_param_0, .param .u32 blend_rgba_p016_param_1, .param .u32 blend_rgba_p016_param_2, .param .u32 blend_rgba_p016_param_3, .param .u32 blend_rgba_p016_param_4, .param .u64 blend_rgba_p016_param_5, .param .u64 blend_rgba_p016_param_6, .param .u64 blend_rgba_p016_param_7, .param .u32 blend_rgba_p016_param_8, .param .u32 blend_rgba_p016_param_9 ) { .reg .pred %p<10>; .reg .b16 %rs<75>; .reg .b32 %r<68>; .reg .b64 %rd<4>; ld.param.u64 %rd1, [blend_rgba_p016_param_0]; ld.param.u32 %r5, [blend_rgba_p016_param_1]; ld.param.u32 %r6, [blend_rgba_p016_param_2]; ld.param.u32 %r7, [blend_rgba_p016_param_3]; ld.param.u32 %r8, [blend_rgba_p016_param_4]; ld.param.u64 %rd2, [blend_rgba_p016_param_5]; ld.param.u64 %rd3, [blend_rgba_p016_param_6]; mov.u32 %r9, %ctaid.x; mov.u32 %r10, %ntid.x; mov.u32 %r11, %tid.x; mad.lo.s32 %r1, %r10, %r9, %r11; mov.u32 %r12, %ntid.y; mov.u32 %r13, %ctaid.y; mov.u32 %r14, %tid.y; mad.lo.s32 %r2, %r12, %r13, %r14; setp.ge.s32 %p1, %r1, %r7; setp.ge.s32 %p2, %r2, %r8; or.pred %p3, %p1, %p2; @%p3 bra BB1_7; and.b32 %r3, %r5, 1; setp.eq.b32 %p4, %r3, 1; setp.eq.s32 %p5, %r1, 0; and.pred %p6, %p5, %p4; @%p6 bra BB1_5; bra.uni BB1_2; BB1_5: ld.global.v4.u8 {%rs71, %rs72, %rs73, %rs74}, [apix]; bra.uni BB1_6; BB1_2: setp.ne.s32 %p7, %r3, 0; shl.b32 %r4, %r1, 2; @%p7 bra BB1_4; bra.uni BB1_3; BB1_4: add.s32 %r15, %r4, -4; suld.b.2d.v4.b8.trap {%rs71, %rs72, %rs73, %rs74}, [%rd1, {%r15, %r2}]; bra.uni BB1_6; BB1_3: suld.b.2d.v4.b8.trap {%rs71, %rs72, %rs73, %rs74}, [%rd1, {%r4, %r2}]; BB1_6: and.b32 %r16, %r6, 1; setp.eq.b32 %p8, %r16, 1; selp.b32 %r17, -1, 0, %p8; add.s32 %r18, %r2, %r6; add.s32 %r19, %r18, %r17; shr.s32 %r20, %r19, 1; add.s32 %r21, %r1, %r5; shl.b32 %r22, %r21, 1; suld.b.2d.b16.trap {%rs38}, [%rd2, {%r22, %r18}]; cvt.u32.u16 %r23, %rs71; and.b32 %r24, %r23, 255; cvt.u32.u16 %r25, %rs72; and.b32 %r26, %r25, 255; and.b16 %rs39, %rs72, 255; mul.wide.u16 %r27, %rs39, 129; cvt.u32.u16 %r28, %rs73; and.b32 %r29, %r28, 255; mad.lo.s32 %r30, %r24, 66, %r27; add.s32 %r31, %r30, 128; mad.lo.s32 %r32, %r29, 25, %r31; shr.u32 %r33, %r32, 8; add.s32 %r34, %r33, 16; shr.u16 %rs40, %rs38, 8; cvt.u32.u16 %r35, %rs40; cvt.u32.u16 %r36, %rs74; and.b32 %r37, %r36, 255; add.s32 %r38, %r37, 1; and.b32 %r39, %r34, 255; mov.u32 %r40, 256; sub.s32 %r41, %r40, %r37; mul.lo.s32 %r42, %r35, %r41; mad.lo.s32 %r43, %r39, %r38, %r42; and.b16 %rs41, %rs38, 255; cvt.u16.u32 %rs42, %r43; and.b16 %rs43, %rs42, -256; or.b16 %rs44, %rs43, %rs41; sust.b.2d.b16.trap [%rd2, {%r22, %r18}], {%rs44}; setp.ne.s32 %p9, %r3, 0; selp.b32 %r44, -1, 0, %p9; add.s32 %r45, %r21, %r44; shl.b32 %r46, %r45, 1; and.b32 %r47, %r46, -4; suld.b.2d.v2.b16.trap {%rs45, %rs46}, [%rd3, {%r47, %r20}]; mul.lo.s32 %r48, %r26, -74; mad.lo.s32 %r49, %r24, -38, %r48; add.s32 %r50, %r49, 128; mad.lo.s32 %r51, %r29, 112, %r50; shr.u32 %r52, %r51, 8; add.s32 %r53, %r52, 128; shr.u16 %rs47, %rs45, 8; cvt.u32.u16 %r54, %rs47; and.b32 %r55, %r53, 255; mul.lo.s32 %r56, %r55, %r38; mad.lo.s32 %r57, %r54, %r41, %r56; and.b16 %rs48, %rs45, 255; cvt.u16.u32 %rs49, %r57; and.b16 %rs50, %rs49, -256; or.b16 %rs51, %rs50, %rs48; mul.lo.s32 %r58, %r26, -94; mad.lo.s32 %r59, %r24, 112, %r58; add.s32 %r60, %r59, 128; mad.lo.s32 %r61, %r29, -18, %r60; shr.u32 %r62, %r61, 8; add.s32 %r63, %r62, 128; shr.u16 %rs52, %rs46, 8; cvt.u32.u16 %r64, %rs52; and.b32 %r65, %r63, 255; mul.lo.s32 %r66, %r65, %r38; mad.lo.s32 %r67, %r64, %r41, %r66; and.b16 %rs53, %rs46, 255; cvt.u16.u32 %rs54, %r67; and.b16 %rs55, %rs54, -256; or.b16 %rs56, %rs55, %rs53; sust.b.2d.v2.b16.trap [%rd3, {%r47, %r20}], {%rs51, %rs56}; BB1_7: ret; } // .globl blend_rgba_yuv420p .visible .entry blend_rgba_yuv420p( .param .u64 blend_rgba_yuv420p_param_0, .param .u32 blend_rgba_yuv420p_param_1, .param .u32 blend_rgba_yuv420p_param_2, .param .u32 blend_rgba_yuv420p_param_3, .param .u32 blend_rgba_yuv420p_param_4, .param .u64 blend_rgba_yuv420p_param_5, .param .u64 blend_rgba_yuv420p_param_6, .param .u64 blend_rgba_yuv420p_param_7, .param .u32 blend_rgba_yuv420p_param_8, .param .u32 blend_rgba_yuv420p_param_9 ) { .reg .pred %p<10>; .reg .b16 %rs<65>; .reg .b32 %r<58>; .reg .b64 %rd<5>; ld.param.u64 %rd1, [blend_rgba_yuv420p_param_0]; ld.param.u32 %r5, [blend_rgba_yuv420p_param_1]; ld.param.u32 %r6, [blend_rgba_yuv420p_param_2]; ld.param.u32 %r7, [blend_rgba_yuv420p_param_3]; ld.param.u32 %r8, [blend_rgba_yuv420p_param_4]; ld.param.u64 %rd2, [blend_rgba_yuv420p_param_5]; ld.param.u64 %rd3, [blend_rgba_yuv420p_param_6]; ld.param.u64 %rd4, [blend_rgba_yuv420p_param_7]; mov.u32 %r9, %ctaid.x; mov.u32 %r10, %ntid.x; mov.u32 %r11, %tid.x; mad.lo.s32 %r1, %r10, %r9, %r11; mov.u32 %r12, %ntid.y; mov.u32 %r13, %ctaid.y; mov.u32 %r14, %tid.y; mad.lo.s32 %r2, %r12, %r13, %r14; setp.ge.s32 %p1, %r1, %r7; setp.ge.s32 %p2, %r2, %r8; or.pred %p3, %p1, %p2; @%p3 bra BB2_7; and.b32 %r3, %r5, 1; setp.eq.b32 %p4, %r3, 1; setp.eq.s32 %p5, %r1, 0; and.pred %p6, %p5, %p4; @%p6 bra BB2_5; bra.uni BB2_2; BB2_5: ld.global.v4.u8 {%rs61, %rs62, %rs63, %rs64}, [apix]; bra.uni BB2_6; BB2_2: setp.ne.s32 %p7, %r3, 0; shl.b32 %r4, %r1, 2; @%p7 bra BB2_4; bra.uni BB2_3; BB2_4: add.s32 %r15, %r4, -4; suld.b.2d.v4.b8.trap {%rs61, %rs62, %rs63, %rs64}, [%rd1, {%r15, %r2}]; bra.uni BB2_6; BB2_3: suld.b.2d.v4.b8.trap {%rs61, %rs62, %rs63, %rs64}, [%rd1, {%r4, %r2}]; BB2_6: setp.ne.s32 %p8, %r3, 0; selp.b32 %r16, -1, 0, %p8; add.s32 %r17, %r1, %r5; add.s32 %r18, %r17, %r16; shr.s32 %r19, %r18, 1; and.b32 %r20, %r6, 1; setp.eq.b32 %p9, %r20, 1; selp.b32 %r21, -1, 0, %p9; add.s32 %r22, %r2, %r6; add.s32 %r23, %r22, %r21; shr.s32 %r24, %r23, 1; suld.b.2d.b8.trap {%rs38}, [%rd2, {%r17, %r22}]; cvt.u32.u16 %r25, %rs38; cvt.u32.u16 %r26, %rs61; and.b32 %r27, %r26, 255; cvt.u32.u16 %r28, %rs62; and.b32 %r29, %r28, 255; and.b16 %rs39, %rs62, 255; mul.wide.u16 %r30, %rs39, 129; cvt.u32.u16 %r31, %rs63; and.b32 %r32, %r31, 255; mad.lo.s32 %r33, %r27, 66, %r30; add.s32 %r34, %r33, 128; mad.lo.s32 %r35, %r32, 25, %r34; shr.u32 %r36, %r35, 8; add.s32 %r37, %r36, 16; cvt.u32.u16 %r38, %rs64; and.b32 %r39, %r38, 255; add.s32 %r40, %r39, 1; and.b32 %r41, %r37, 255; mov.u32 %r42, 256; sub.s32 %r43, %r42, %r39; and.b32 %r44, %r25, 255; mul.lo.s32 %r45, %r44, %r43; mad.lo.s32 %r46, %r41, %r40, %r45; cvt.u16.u32 %rs40, %r46; shr.u16 %rs41, %rs40, 8; sust.b.2d.b8.trap [%rd2, {%r17, %r22}], {%rs41}; suld.b.2d.b8.trap {%rs42}, [%rd3, {%r19, %r24}]; cvt.u32.u16 %r47, %rs42; suld.b.2d.b8.trap {%rs43}, [%rd4, {%r19, %r24}]; mul.lo.s32 %r48, %r29, -74; mad.lo.s32 %r49, %r27, -38, %r48; add.s32 %r50, %r49, 128; mad.lo.s32 %r51, %r32, 112, %r50; shr.u32 %r52, %r51, 8; add.s32 %r53, %r52, 128; and.b32 %r54, %r53, 255; and.b32 %r55, %r47, 255; mul.lo.s32 %r56, %r55, %r43; mad.lo.s32 %r57, %r54, %r40, %r56; cvt.u16.u32 %rs44, %r57; shr.u16 %rs45, %rs44, 8; sust.b.2d.b8.trap [%rd3, {%r19, %r24}], {%rs45}; and.b16 %rs46, %rs43, 255; sust.b.2d.b8.trap [%rd4, {%r19, %r24}], {%rs46}; BB2_7: ret; } _;(D<Pt%&)L+HX+\p6666@:=,ATC|DFG$K Q8TV`WZ<\` d4 |e\ i Tl `n p s, @sT @{ p   l( P | Ж H  ( P Ȩx  ̩   $ P x D < p  Hp8ĸx@@hDp 4`8004X@\p8d\@ xH$p 0\40tX|, lHplL$ Lltx<  4L\ <<"# x#$/P/x0122 p98:`=>ABG0,HXxIJKLN(HPTP|`TVhWX$YL\[t\H^zRx 0,@ @,HA BNdPzR| ((D H    H$D@D0H h 0H$l D`H  `H("LDH 4 H$ $$ D0H  0H/D T/ D X$,/(D H P  H$T/lDpH T pH$|43hDPH P PH$t6 D`H  `H$l:DPH  PH$ <D@H  @H$<LDpH 4 pHD?$D \$\?$D`H   `H(CDH  H$H<DpH $ pH$MD@H p @H$LODPH  PH$(QDPH  PH,P\SD H  h  H(WDH  H$L[DH  H(\,D H    H(`DH  H$,b DPH  PH$T|dDPH  PH$|fDPH  PH$h8D0H ` 0H(hDL   L(p0DL   L$$tD`H h `H(Ly,DL    L$x{PD@H x @H$<{0D@H  @H(D}D H    H(HDL  ( L(  xDH ` H$LlHD@H p @H(tD H  l  H$D`H  `H$DPH  PHHD T$LDpH  pH$0 D0H  0HXD T(pTD@H | @H(О$D H     H$ȢD`H  `H(8DL  l L(DL   L$Hd4D H \  H$ppD@H p @H$ЧD0H | 0H$<D@H  @H$D@H  @H$ DPH t PH(8  DL   L$d D0H | 0H( \DL   L$ pD0H X 0H( @DD H  ,  H$ XPD H x  H$4 D0H  0H$\ pDPH X PH$ \tD`H \ `H( DH  H, <DH $ H$ |D0H d 0H$0 |D0H d 0H$X HD0H  0H$ D0H  0H$ \D0H  0H$ Ժ D0H  0H DT$ D0H  0H$8 <D0H  0H$` D@H  @H$ X8D@H   @H( h DL   L( HDL   L 0D $ $D`H   `H$H D@H  @H$p @D H   H$ |D0H d 0H$ (D0H  0H( D0H t 0H$ ,DPH  PH$<0D0H  0H$dD H   H,D0$(D H    H$tDPH  PH(DDL   L,(D H  h  H$XD`H | `H$D0H  0H$@D0H h 0H$8D0H ` 0H$$DH  H$ `D H H  H$HPD@H x @H$p,DH  H$XDPH @ PH$@D@H ( @H4DPDH$XD0H @ 0H$D@8D0H ` 0H(lPDL   L(DL   L$XDPH @ PH( DPH  PH$(D H P  H$@4D H \  HhhD`$(D`H  `H$@D0H h 0H(D H    HtDT$xD@H  @H(@lDH  Hl@ D0 pDh(L  DH  H$D@H  @HdD0$(DH  H$0XDPH @ PH(XDpH  pH(DH  H$0hD@H P @H(p DH  H$ #DH  H(,#|DL  \ L$X &DH  H$&D0H  0H$,HD H p  H$,LD0H 4 0H$-LD0H 4 0H$ /0D0H  0H$H0D0H  0H,p0DL   L(2dDH L H$3\D H D  H( 4DL   L( 78D H     H$L9D H   H$td:D0H  0H$:D`H  `H$<hD@H P @H$=DPH p PH(@?dDH L H(@x@DH  H T@@cP_  ա?T!!X@a_ ` !!"A !AT"B\@b_{ `"|9@5 `@A=P R`"<9 @{¨_{  Ѡs^q7_(e@@ B%R)"P_^@R$RPq7_(e@@ bc%R6 P_^@ R$RPq7_(e@@ Ղ%ROr_^ @Oq7_(e@@ #&RT;O^_^ @Oq7_(e@@ b&R`6OI^ @))i@)@ 7s^ @))m@)@ 7(Rs _(e@@  c'RT;O&3/+s@q ))JJr##@a%B4-џO@O33@7_(e@@ բ(R<"}O'#@ a%B=уO @O//@7_(e@@ "c)RaO'#@a%B hOiO++@7_(e@@ բ#*REO'(q@@!XO@@*WO_@  UO@_@@ @  POq7_(e@@ B+R-O'}@?Os@@(@ (R(i*@ *(i*@ (i*@(i*@(i*@ *(i*3@@ *(i*/@@ *(i*+@@ *(i*@ *7@(i*_O@ )i)7@O7q(7@ O_(u@@7@ "-Rt6O',_aRO_@@!Oq(7@N_(u@@7@ բc.R6N'_@N@N_(y@@@@#+@7+@N/@7/@N3@73@N'@_ @{¨_{ _ @@@aRN{B_{CS/_^@cNq7_(e@@ ՂCtRLN@q7_(e@@ btR;N_@!^!Ý Nq7_(e@@ buR)#Np_@A^!Ý yNq7_(e@@ bcuRX4 NX_@ @#hNq7_(e@@ uRd0MC@@b^!Ýy)_@@CQNq7_(e@@ bcvR-M(@H7@@((y@@_@@@/Nq7_(e@@ "cwR0+M(y@@_{E_{Cѿ_^Nq7R|_] Nq7RqC]q7C]Aq臟7~Rd]}S#]}S#;]< @ !~7^ ) T*M@^(R(i*C]^*(i*^ *\+7@/;@3@(i*3@ *(i*'#_ѡ]CMq7_^ ) '@CM(y@@h7#@^ *) ' *k'(y* ^ ) T*BMȀR_{G_ =_{ _!4M(y@_!"M(}@_!p|M(@_!d*uM(@_!"nM(@_!p*gM(@_!( `M(@_!"YM(@_!*RM(@_!?KM(@_!40DM(@_!==M(@_!6M(@_!"/M(e@_!((M(@_!@!M(u@_!M(@_!*M(@_!% M(i@_!H M(m@_!\.L(@_!L(@_!,L(@_!2L(@_!#L(@_!|1L(@_!(6L(@_!$#L(@_!L(@_!PL(@_!.L(@_!`6L(@_!:L(@_!&L(@_!L(@_!L(@_!2L(@_!|L(@_!9L(A_!p>yL(A_!@rL( A_!h kL( A_!dL(A_!+]L(A_!>VL(A_!OL(A_!>HL(!A_!$?AL(%A_!:L()A_!t?3L(-A_!p,L(1A_!7%L(5A_!L0L(9A_!+L(=A_! L(AA_!4 L(EA_! =L(IA_!(K(MA_!K(QA_!6K(UA_!1K(YA_!:K(]A_!:K(aA_!:K(eA_!\#K(iA_!?K(mA_! K(qA_!1K(uA_!K(yA_!dK(}A_!4&K(A_!K(A_!*K(A_!?K(A_!4;K(A_!}K(A_!4:vK(A_!loK(A_!LhK(A_!.aK(A_!|&ZK(A_!2SK(A_!LK(A_!P=EK(A_!/>K(A_!$7K(A_!0K(A_!+)K(A_!."K(A_!hK(A_! K(A_!& K(A_!=K(A_!J(AhR(q@@7_@!4BJ(q@(A@7_@!B@J(A @J*{B_ *_ *_{C@J{A_{/+'^ q7_(e@@ \JR) J_^@(q@@ѢJq7_(e@@ YCKR6J^@))@)@ 7*^@))@)@ ))r@#_^ @VJq7_(e@@ SKRT;Iu#@ q7 @q7 ]AWJ/]AUJ/ _(e@@ բN#MRIR/@q7+@q7_(u@@/@ KMRPI;_*1J''@+@ k77]A#@'@&J/q7_(u@@/@ BGNRI_@J _ @@J'@'_(y@@@9_{F_C{^ q7_(e@@ @#4R)GI_@@(q@@cIq7_(e@@ =4R6/I@@I17_(e@@ b;4RI_@ @3Iq7_(e@@ 8C5RT;I}@q77_(e@@ բ65RT;Hl@A@R**Iq7_(u@@@ Ղ36RIS@A3fIq7_(u@@@ 06Rh I=@A*"RoIq7_(u@@@ -7RХ$:H&@ *(i*@ ) @4"BpQIq7_(e@@ )7R H(y@@_{DC_{C^q7_(e@@ բ%SR)mH_^@(q@@Hq7_(e@@ բ"SR6UH^@H17_(e@@ " CTRAH_^ @Hq7_(e@@ ՂTRT;,H_^ @sHq7_e@@ UR`6H_^@cxHq7_e@@ BcURЄ<H_^@3cHq7_e@@ բURЄ.Gl^@H17_e@@ "#VRЄ>GX@q7H@ A a(i*7(}ҨF@@ *(i*@q7R@A@@#@@lH'q7_u@@'@  XRL7G@A@#@@*CH'q7_u@@'@ XRХGy@@_{E_C{c_@@q@@Gq7_e@@ ÇRЄ6_GU@@ )@)@ 7*@@ )@)@ ))r@@ q7 @q7 @aG@ C G _e@@ BRЄ"G@q7_u@@@ RХLGy@@_{DC_{ ^q7_e@@ 1RЄ)F _ @@q@@#qGq7_e@@ 2RЄ6F@y@@_{C_{^ q7_e@@ բCRЄ)Fy_^@q@@7Gq7_e@@ բRЄ6Fa_^@Gq7_e@@ cRFL_^ @Fq7_e@@ bÌRЄT;sF7C +C]3@(R#R'(R+] !G17_u@@ CR$}RХF_y@@_]ii G@@ _{F_@@_{C^ q7_e@@ "8RЄ)F_^@q@@Fq7_e@@ 9RЄ6F_^@iFq7_e@@ b#:RE_^ @TFq7_e@@ :RЄT;E@A`F@q7_u@@@ ;RХh E@AMF@q7_u@@@ ;RХh;E~@A*"RQFq7_u@@@ ";R$:Eg@A"R*:Fq7_u@@@ BcRЄ 0Ey@@_{E_{^ q7_e@@ ՂaRЄ) E__^@q@@юEq7_e@@ ՂCbRЄ6DG^@ )@)@ 7(R_^@Eq7_e@@ bbRD&_^ @cEq7_e@@ բccRЄT;D] )iiE])iiqh7] )iiE_u@@])ii] )ii #dRDsE] )iiqh7@q7] )ii q] )iiSE@q7@@@@] )iiqh7] )ii7E_e@@ բeRЄ(6UD] )ii觟7] )iiE@@@qh7] )ii@ 7H7] )ii@ E@h7] )iiD_e@@ CgR*Dj@] *(i*@] *(i*@] *(i*@q7] *(R(i*] *(i*] )ii@@D]a ii RHr k(7] )ii@D@] *(i*] )iiD] )iiD])iiq7_u@@])ii] )ii biRDDy@@_{G_{^q7_e@@ C~R)C2_^@q@@Dq7_e@@ ~R6xC] iiq7_e@@ բR7eC_^@Cq7_e@@ R.PC^@ )@)@ *))r#] )iiD])iiqh7] )iiD_u@@])ii] )ii {R C] )ii7] )iiC@@] )iiqh7#@q7] )ii q] )iiC@q7@@@@] )iiqh7] )iiC_e@@ BqR(6Bl_^@*Cq7^ @ ^ @_@Cq7_] @@C''@q7] )iiyC_u@@'@ jダR0?B8]**(i*C]]*(i*@] *(i*#@qH7] *(R(i*] *((i*] *(i*] )iiOC] )ii:Cy@@_{F_{ ^ q7_e@@ b`kR)CBG_ @@q@@#Bq7_e@@ b]kR6+B/_ @@Cq7_e@@ ZClRЄB@1@q@@ C @Bq7}@@y@@_{C_C{ ^ q7_e@@ TnR)A_@@q@@cbBq7_e@@ QCoR6A~_@@3)Bq7_e@@ bNoRЄAi@ C @|~B  @@) @ k77@A x)}  5@@ k7@A x)}   @qH7@A x)}   @  @7_e@@ EqR%nA$ @1 @q @@|;B@ C @*Bq7_e@@ B#rR9PAy@@_{DC_C{c_@@q@@Aq7_e@@ =PR6(AD@@ )@)@ 7* @@ )@)@ ))r  @@ @ )@)@ *))r@A@@Aq7_u@@@ B5QR)&Ay@@_{DC_{O`B *@hNJF^ q7_e@@ /ZR)@_^@q@@cAAq7_e@@ ,#[R6@_^@b"c@q7_e@@ "*[RЄ@h@ iiq7_e@@ '[R7~@hJ@ q跟7_e@@ բ%c\RH6m@_h@@bB R$RXAq7_e@@ "\RV@h@@@`F17_e@@ B C]R,.B@}hJ@qǟH 7h@!hhJ@qh.h.@i ~)=*|) ?ihh@i.@ k7(7h@_i@j!yji @q7_e@@ Ղ^R82 @(Rhh@hh@@ )A)@ 7i@ (R(i*c@h@AaF@b.@@`Nh@h@qh7! h@AaF@*@`NhN@q7_u@@dN@ b`R/@y@@_OA{¨_C{^ q7_e@@ B cR)?ѿ_^@@q7_e@@ B CR.?_^@ѣ?q7_e@@ ՂRL(|?]q7C*?@_\!!1?q7_e@@ "ÑRX?Y?_\ACq?q7_e@@ b#R$*C?r_\ @C?q7_e@@ R.?]_\@C?q7_e@@ 㒀R?G_\@C!?q7_e@@ BCR?1_\@?q7_e@@ բRp<>]5@]C?17_u@@ R$}R*?y@@_{HC_C{ ѿ_^@?q7_e@@ BcRd(>]5@C?qH7_y@@_]5@f?@@o}@@_{LC_{ Ѡ^q7_e@@ "#R)y>_^@R$Rt>q7_e@@ BR6b>'#Sa%Bh"e>@g>##@7_e@@ b㬀R4C>S a%BhJ> @P>@7_e@@ R;(>{Sa%BT:0>1>''@7_e@@ բcR >`A@-!>@@* >_@a >_@@a@>q7_e@@ cR-=:@ >#@@(@@( '@@(_>@q@@4B=~>q7_e@@ ՂR =_@=@=_y@@@$'@7'@=@7@=#@7#@=@7@=@_ @{¨_C{ ^ q7_e@@ cR)o=_^@A@=q7_e@@ ôR6W=y_^@Ѹ=q7_e@@ B#RB=d_^ @=q7_e@@ բRT;-=O](>q7_e@@ BRX>=<\@=\I@q7\@=_u@@\I@\M@ ·öRƠ= \U@qH7\@=@@]\(U]\(-\ @=\@=y@@_{LC_C{_@@A@cU=q7_e@@ bùR6<`_@@S=q7_e@@ ¬#R<K@q7_e@@ բR<:@@d=@Aq跟7@yi@ kH7@ *(y*(R @h7@@@A=@q7y@@}@@_{DC_C{^q7_e@@ "!R)A<Rq@@!U<@7_e@@ "#"R<)<:@"*B<_@  @<@_@@ @  ;<q7_e@@ բ"RЄ-<@ (R(i*@hR(_@A< @.<_y@@ @_{DC_C{^q7_e@@ Ղ#@R);_@@q@@cV<q7_e@@ Ղ@R6;@@ q7_e@@ B@R.;y_@@S <q7_e@@ բCAR;d_@ @C;q7_e@@ ART;;O@ @ )@)@ 7*R*r@ @ )@)@ 7(R(˪r _e@@ BBRЄt2Z;)@qǟ7@qן7_e@@ bBR;C;@@(i*@@ (i*@@ (i*y@@_{DC_{ _ @@ARu;{B_{_^ )ii;^ )ii跟 RwH7^)iiqww@6^ )ii^ )ii;^)iiq(7^ )ii;CC]^)@ k7H7^C x)}   @qH7^C x)}  C]C^ )ii;]7^*(~R(i*^ *R(i*]@^ )ii^ )3ii:3@^ii](1^ )/ii](q^a]Z;C^ )7ii[;/@3@7@^Hi,^Hi+^iiS;C^q7C^^*(i*^ *R(i*H^ )ii9;^)iiqH7d;\^E@c;\)A@\%^  \[X;\Z;^ )ii;]1@q7(R^ ) ^4Bp0:q7?;@^*(i*^ *R(i*^^^q7^)iiq RW7^@q藟WW@6^a;CC^q7C^^*(i*^ *R(i* ]1@qh7^)iiq7^ )#ii':#@'@^(i*^A*:^ )ii:@@^(i*^ )ii::^)iiq(7[^E@:[@@Zu^)iiq7^ ) ^  @Y[Z:=[^E@:[[^)ii?:[^ )ii8:@@A[MA@YXC^)iiq7^ ) ^   @??@[C@Y:[Z:^ )ii:^ )ii:@^ +HRHi+^ii :@^ii:^ )ii:@^ii9_@{P_{_C+c*/9@^A3:CC^q7C^^*(i*^ *R(i*[@q R7^)iiq RH7^ )iiq@h7^ )iiqH7^)iiq7^CC^q7C^^*(i*^ *HR(i*^)iiq7^Ac9qh7[@q7^*}R(i*^ *(R(i*~^ C C9CC^q7C^-q7}R9C^.qh7dC^^*(i*^ *(R(i*U^ )ii>9^)iiqh7^)iiqh7a9''@+@5@9#'@^E@[9'@+@q@V9'@'@^A |ii 9'@^A|ii9@@ '@%A@@#@@@^  '@@49'@69^ )ii8d_@{P_{CCC_q RH7_)iiq@ 6_ )ii8_ )ii跟 RH7_)iiq@6_ )ii_ )ii8_)iiq7_ )ii8CC^_)@ k7H7_C x)}   @qH7_C x)}  C^C_ )iiw8^7_acѲ8_q7_-q7}R8~__*(i*_ *hR(i*_ )iiH7^@_ )ii_ )ii7@_ii^(1_ )ii^(q ^1_ )ii^(q_a^8_q7__*(i*_ *R(i*I_ )ii 8@_ +Hi+_ +Hi+_ii8_ )ii7_)iiqH7$8@_E@#8@)A@@_  @@8@8_ )ii7^1@q7(RCC_{E_{C*C7@_ C 8q7^o_Asc8_A |ii@ k7_A|ii@ k7U_As7q7^IsѨC^C_AA7q7^:_A"R*P7q7^- ^C^ k7H7_A^!RC%7q7^_ C @H7q7^ ^^_@{Q_{ __@ @@@R6{C_{_ )iih7_ )ii qh7_ )ii 7_*R(i*_ *R(i*_A*O7_ *(i*_ )ii 7__ii6_ )ii77_ )ii(7_ )ii(7_ *(i*_ )ii7_ )iiC6@_(i*_ )ii7_ ) ii,6 @_(i*_@q7_ C * 7@_a7{C_{ =Ѡ_*#ց5@^ *(i*^ )iiq(7^ )ii#6q7]^*(i*^ *hR(i*^ ) ^,>B >6@^)@ k767^A@*6]q7]^*(i*^ *R(i*^ )ii96^)iiq7^)iiq7\6??@)A@?@^5@6@@;^  ?@;@M6?@O6^ )ii 6^ )ii6^ )ii R'H7^)iiq''@6^ )ii^ )ii6^)iiq7^ )ii57^ )ii5^ )ii 7oo@^)@ k77^Ao@ )}  !o@o^)iiq7^a6q7]^*(i*^ *(R(i*^ )iiq7^ ) k@^ ) 51R*JJr^ )ii^ ) k@^ ) #5]q7]^*(i*^ *HR(i*gg@^)@ k77k@^)Ag J} ) (%g@g^ )ii@5^)iiq7k5//@^ )ii$5+/@-A@+@*'^  /@'@]5/@_5^ )ii5GG@^)@ k7H7^)ii^)AG J} ) (%G@GCC@^)@ k7h7^AC@ )}   @^)AC@J} ) )@} ^)AC@J} ) (!C@C^ )ii^(q^R(^a^4^ )ii4@^ +Hi+^ii4]q7]^*(i*^ *R(i*^1@q7(R@E^]#]q7^A*4^ )ii 4@ @^(i*^ )iiq(7^ )ii#4^ )iip4@^ +HRHi+^iik4_= @{¨_{__)iiq7_ C cђ4CC^q7C^-q7}R4C^_*(i*_ *(R(i*^1@qh7^5@C^@^1@C_ )iiCZ 77_ )iiCZ*4Y7_*pR(i*_ *R(i*uY_ *(i*CZ_ *(i*_ )iiZCZ*D3@_iiCZ*#e4_ )ii3_)iiqH7 4Y#W4CY^q@4?^@r7 A@ A@@;Y%A@C@?@;@_  YX3Y3_ )ii3_ C ^3CC^q7C^_*(i*_ *HR(i*6_ )ii337_)iiq(77@_E@337@@@3@s/_)iiq7_ ) _  @++@7@/@3=7@_E@3'7@ 7@_)ii=37@_ )ii63@ @?#7@MA@'@#@_)iiq7_ ) _  @@7@@W37@X3_ )ii3@{P_{ C_)iiq R'7^q''@h86ѿ_a73^q7^-q7 R:3^_*(i*_ *hR(i*_a]5@Y3_ )ii2_)iiq7_)iiq72])A@ ]]5@2 @@_  ]\2]2_ )ii2_ )ii2_ )ii RH7_)iiq@6_ )ii_ )ii2_)iiq7_ )iin26]_ )iif2_ )ii H7(RCC\^)@ k77^AC\ )}  !C\C_)iiq7_aѦ2q7^_*(i*_ *(R(i*_ )iiq7_ ) \_ ) 21R*JJr_ )ii__ ) \_ ) t2^q7^_*(i*_ *HR(i*[^)@ k77\^)AÛ J} ) (%[_ )ii1_)iiq7_)iiq71##@_ )ii1#@-A@@_  #@@1#@1_ )ii1//@^)@ k7H7_)ii^)A/ J} ) (%/@/++@^)@ k7h7^A+@ )}   @^)A+@J} ) )@} ^)A+@J} ) (!+@+_ )ii^(q^R(_a^M1_ )iiN1@_ +Hi+_iiI1^q7^_*(i*_ *R(i*.^{I_{ __@ @@R0{C_{ @Ѡ_S*W\0W@ցX0_I@qH"7_@1_U@q RS7_I@qSS@6_ @_@1_I@q7_@0_@0_U@1q(7_R(I_R(M@q(7_}R(I_HR(M^7@@_q@D1^(7_HR(I_HR(M^_U@_-@41Cq(7C^_(I_(R(M_@0_Q@q70KK@'eA@#K@_U@w0#@'@G_aK@G@0K@0_@v0_@o0_U_scC_@h0@q(7_pR(I_R(Md^@0Cq7_ ꛀ*(y*C^.qh7NC^_(I_ȥR(MG_@90_Q@q7f0CC@C@@ 0C@@0@@%?C@iA@C@@0C@S@N0@@@?@;_aC@;@E0C@G0_@0_@0_sc?C_@/_@/_Q@q7077@ 7@_I@/7@_M@/@ @37@MA@aA@3@/_a7@/@07@0_@/^Z0@ @{¨_{C'@Aq跟7_yiq7_yiq7_yi'@#_ (R(y*_yiqh7@'@q7@Aq跟H7  ^%@^5@(RRRT/q7'@#'@_)(y*_ (R(y*'@qH7'@^(#@@(pR_{E_C{_^ )iiH7^ ) ii'/ @^+(RHi+^ii"/^@ qH7CC^^)ij k77^iiqǟh7^ii/^*(i*C^Ce^^ )ii(7^ )ii/^ *(i*^ *(i*^@q7^Ay/^ )ii7^ )ii>.@^(i*^ ij#^*(i*#@q7#@O/^ij^*(i*@q7@@/^ij^(i*@17@1/{DC_  _{_] )iiz..]@@ @@@]E@*].@ @f@@ ) .@.@]+(RHi+]iiV.{F_{ @@!@7@I@q7@@9.@(wR(I@ @G.@@4.@!@w.@!@@7@@-@ @7@ @-@@7@@-@Aq跟h7@yiq7@yi.@ *(y*@yiqh7@{B_  _{ _@@-@(R(Q@@-{C_{ Ѡ_A^AR_.W17}RW@ T.q7-@_@17}RC]r7}R_@W(W a(i*W 1 0#*?,@?@#@W1(R(i+W@W(i*W!,WHR( W@W( W@Oq7O@q 7WR**.q7}RW1 >҈R(i*W ?*(i*_C' Pc)w,@y,@W <W |7pRmW ! (R(i*W)! (i*_C' բPc)T,@V,W<W|7pRMJW1 >(R(i*W ?HR(i*_C' b{Pc2,@4, @W <W |7pR(W ! ҈R(i*W ! hR(i*_C' vPc,@,W<W|7pRH~R_ @{¨_{CCoC= =^^C^_@AˊP-@q7,@_{E_{CC_!C*+_@q7@q7}^ R r k7 @ R)r k7 @ RHr k7 @ IRɆr k7 @ R鈦r k(7 @ IRɦr k7^C#(R39^/],}RF^ R)r k7@ɩR)Fr k7@)R)Fr k(7@)RIr kh7]  ] C^^^,q7]C^C^^#] ii39],}R}R]_@{Q_{CC_ 1 C*,+_@q(7@q 7^ IRɆr k(7 @ IRɦr k(7^C#C^^(R39^/]l,}R8^ R)r k(7@ɩR)Fr kh7]  ] C^^^Q,q7]C^C^^#] ii39]8,}R}R]_@{Q_{C@@$Rr{A_{ ^# _@aÊ,@q7+@@ @(_{C_{ @** @@@!ˊ+@q7r+@-q7l+@_{B_{_@!+q7@,_@q7@#@@)@ k7(7_@*#*q7@ _ aii@(u@_{C_{ ^ q7 @q(7_!@ @+ _ 1 @ @+(R@_{C_C{_ 1 3l+q(7@^(@@( _ @^(_ @@({DC_{c@ 1 q*Cq7C^p@@@qH 7@@ 1 *Cq7C^[@@ b Ҹ)@@@ @ | " @| * @@@@C@ @Cik@ii*1*Cq7C^4@/@*Cq7C^'/@q7(~R@@@/@ k77@/@)@(@ @ 1  @*Cq7C^_@{R_{CC_ q7 @q7_!@\* _ 1 @U*(R@{B_C{ )#!??o = =*RC( (R^ŠRr k7@i9R)r kq臟( 7@9R)r k q臟(7@LR)r kH7@iMR)r kh7@)`R)r k q臟h7@`R)r kq臟( 7@`R)r k7@ bR)r k7@ibR)r k7@bR)r kh7@bR)r kq臟7@icR)r k 7@cR)r k7@IdR)r k q臟(7@ eR)r kH7@IeR)r kh7@eR)r k7@i R)r k7^ q7HR @ !@)@!@@*x( @@YOYK)@!)^q7HR @ !@J@# !@J@' !@J @+ !@J @/ !@J@3 !@J@7^ѫI)@ъ^q7HR{ @^ѩ !@I)@wm^q7HRh @ !@*@s^ѫI)@aW^q7HRR @^ !@)@C)@ND^q7HR? @^ !@)@C)@;1^q7HR, @^ !@)@C)@(^q7HR @)@@q7H~R @^)@hR]_@{WC_{  @@@Ɋ)@q7(@@_{B_{ }y *'@@sҌ'@ @}@ 1 (Cq7C^i@}@ 1 (CqH7@@*y@H@)y@(QqH7@Q*y@HQ)y@(Q%qH7@Q*y@HQ)y@(@}@1iiq7@}@ 1 s(Cq7C^$#@rH7@/@*y@H +@)y@(@7@)y@(Qq7@R)y@( @R)y@( _ @{¨_{C@ 1 '@!'@(@}*'{A_{ @ aii7 @ aiiu(@ @(i* @@q7 @@k(q7'@ @_{B_C{C3o==)R(RdR(r#C_/_{DC_{C;o==(R (RdR(r+_(_{E_{ C^@r7^4*#~&@@CoC= =+R +R aR+r3I^) @k'_@Ȋ'Cq跟7S'@^^(_{I_{ C#*E&@Co= =*R *R(^ @k^+^/_@qh7@qH7cR(r3`R(r3H~R_@ȊCo'Cq跟7'@_{I_{  @@@( @@@aNJҁQ'@q7&@_{B_{  @@@( @@@2'@q7&@_{B_{  @@@( @@@'@q7&@U @@q7@B9 @ (i*@@ @ (i*@ @ )ij k77@q@)}  @ @) @} ) ( @q@)} ii @) @J} ) (@ @ (R(i*@@ @ (i*@@ @ҋ(i+@ @ @(i*@ @ @ (i*_{B_{  @@q7~R @@@( @@@Šҁ&@q7*&@_{B_{  @@q7~R @@@( @@@Šҁq&@q7&@_{B_{ @ @ )ij k71@q7@q7 @@ @AŠҁ@&  @@ @aŠҁ7& h~R @q7%@@ @ (i*_{B_C@ iiC_{ @ @q7 @ iiq7}R @/& @ iiq7@ q( 7@qh7  @ @) @ k77 @A @ x)}  &q7@ @  @*q7@o @ ii q7  @ @) @ k77 @A @ x)}  %q7@u @  @*q7@d>  @ @) @ k77@ @)A @ xJ} ) )@ k77 @A @ x)}  A )}  @c%@ @  @*jq7@*~R& @|7 @|$  @ @) @ k77 @A @ x)}  x*# @  @ _{B_{_|z$CC__.)ij k跟7_ Cyi%C_C_.*i)_ #@@_ #_|U$@{B_{^q臟7~RW# ^ _@_ ii_@ҁ$@q7k$@; @q(7^q7/ @q7 @_( @ @ k7H7_A@ x)}  _@_ ii_ ii_ @$@(~R_{C_{_ iiq7~R|_@@@q7_#@m_ @@(@_) @ k7 7_ ii q7 @ q7 @q7D_A@ x)}  _@"R$q7@<4_ ii qr_A@ x)}  _ii_ ii_ ii_@@j$q7@_A@ x)}  _@\$~R@_{C_C{959@ iiq(7@9@ @qH7@9@ iiq7hR *}"@  x" @@@ ! 9@ )ij 9@)@C 9@ )ij9@@!Š #@C9@|#C^q7Q#@-q7X r(7@9@|#(wRc@9@|";#@[@ 9@I @)qI 9@)AX xJ} ) -X -@I @ H)} @) -@Iq-@ /@/@)-@)@ k7h7@(!@/@Jzij)-@)A/@ J} ) (!/@/@9@ iiq7@9@@)q7@(!@ @)-@(5@5@7@(-@)5@(@9@|"_@{[C_{511@7@5@|"@1@ @q7h}R@ 1@)RI 5@|t" *!@  Ҳ! @@(!(1@@(1@@C(1@@(1@ @(1@@C(1@q@ H ʚk}  ((1@q@ ʚ( '@'@)1@)@ k77@ 1@)A'@ J} ) )!@ !@'@kzIi+5@ ii q(7@(1@A'@ J}  -@*!@'@J H(1@A'@J}   @)!@'@) (@1@@q7@(1@A'@ J}  %@)!@'@) ( '@'@5@@ џ"@C5@|!C^q7@ 1@? 5@|!&"@ @ 5@I @)I 5@|!_@{Z_  @ @* k7(~R@ @) @ k臟7~R  @A@ x)}  @(@_{C3_ ii"q7~R1_@%q7@)q7R@R~Rc#_ii_ ii(R+@/@3^_!17(}R_{E_{_|Z!_.iiq觟7_|T!~RX@_.)ij k跟h7_ yiq7@@@1H7_@ "@q7@@_* _.kiHy+_  _ i(RHy+_  (@(_ yi@(_  (R(y*_| @_{C_{C@| @@.)ij k跟7@ yi@ k(7@  *(y*@@| *{A_{  @@@( @@@Šҁ`!@q7 @_{B_{ @ iiq7 @ @q7@-q7@q7@_{B_{ @ iiq7 @ q(7@1 @@W @_{B_C{C_*^_(__(C__(@_( _ __@ҁ@_ |@q7 7_A@ )}  (%_A@)}  @@J}  = =@(_A@)}   _A@)}  !_A@)}  )_A@)}  -@@)} ii@@J} ) )@} @@J} ) ) @} _)A@J} ) (@@)} ii@@J} ) ) @} _)A@J} ) ( @*{DC_{ @@q7R< @@q7~R3@ @)@ k7(7 @A@ )}   @H7 @A@ )}  @ @A@)}  -@*  @A@ )}   @_{B_{C@@q7~R@@q7~R @5@ 17(}R_{A_@ R)r k$7 @)RI&r kH7 @ɩR)Fr k7 @)R)Fr k7 @ɊR)Fr k7 @)RIr k(7 @)HRIr k(7 @IRIr kH7 @)R)Ʀr k( 7 @FRIr kH 7 @)ˊR)r k7 @)R)˪r k(7 @HR(r k7 @*R*r kh7 @*R*r k7@hR(@@(@@(@@()@@(@@(@@(-@(R( @(@(1@hR(@@(@IR @H@ @()@@(@@(@@(-@(R( @(@(1@hR(@@(@@(@@()@@(@IR @H@ @(-@(R( @(@(1z@hR(@@(@IR @H@ @H)@@H@ @H@ @(-@(R( @(@(1X@HR(@@I@) @I@@I@) @I@)RI @(A@(R(@@I@@I@( 5@(R(@@(@@(@HR( (@(R(@@(@@(@R( @HR(@@I@) @I@@I@) @I@( @R(}R///@_{  Ѡ_@q7~R *D@  ? @@_) @_)@_)@C_@^! Cq7@*  @_@_ @^] k7 7^HRr^ŠbCC^q7@\@q7(}RR@_)A^ * k} ) (%\^ @_)A^k} ) (-\^  @_)A^k} ) ()_A^)}  -@*_A^)}  %@_A^)}  )@*bR#Rp@_A^)}  _A^)}   @7H}R^_ @{¨_C{ +^*R*r k7(R^)R)˪r k7RH}RC_@_@#@'(R+_ 17H}R_{DC_{_@ q7~RY *E@ @@@_) @_)@_)@C_@^! q7@2CC^] k7(7\C^ @_)AC^ * k} ) (-_AC^)}  )_AC^)}  -@*@_AC^)}  C^C_@{Z_{ ѨQCQ@@q7~R@59=C^Cq7~Rx *@  ҿ @@@ Q@J @ Q@J@C Q@J@ !Q@@^! q7@P]qH 7@Q@@ %q7 @)q7 R(R @R(R~R1#ѿ^^ZC[[17(}R/@/@)Q@)@ k77@])Q@)A/@ J} ) (%/@/_ @{¨_{CC^^_@_C^ _@_@ @@1__@<{E_{ _@9q7 @_@9q(7___ @@@{B_{ _ @@@aR({B_{ _ @@AR{B_{C^^ _ @/_C^w_^s_]#_^@_@@]@@__@{G_{C@17@qh7@qH7@q(7@q"7@qh"7@qH!7@q(7@q7@q7@q7@q7@q7@q7@q7@q7@q7@qh7@qH7@q(7@q7@q7@q7@q7@q7@qh7@!qH7@%q(7@)q7@-q7@1q7@5q7@9q7@=qh7@AqH7@Eq(7@Iq7@Mq7@Qq7@Uq7@Yq7@]qh7@aqH7@eq(7@iq7@mq7@qq7@uq7@yq7@.qh7@.qH7 ՈFPq0u;*=1;!510A3q,?:9 !/,}0yI/u6q/miE+eY2a]%8YU-U-QM I-EiA/== 9I15%1 -E ))5%)'! !>= A< 2@@{A_{ __@ @@@R{C_C{C^^ _ @_C^V_^R_]_]J_^]@_@@\@@@ #__#@{HC_C{ __^@@ @@R{DC_{ 3_@@q7:@q73_@@ @q7(_@ @ @q7_@ @ @!q7_@@ @1q7(R_{C_C @ R)r kH7@ɩR)Fr k7@)R)Fr k7@)RIr kH7 R R RR@C_C@)QA)@ 7ŠRr@)UA)@ 7IRȆrx@)YA)@ 7IRȦrl@)]A)@ 7 RHr`@)@)@ 7ȩR(FrT@)@)@ 7R(rH@)mA)@ 79R(r<@)qA)@ 7hMR(r0@)uA)@ 7(`R(r$@)yA)@ 7LR(r@)}A)@ 7cR(r @)A)@ 7H`R(r@)A)@ 7bR(r@)A)@ 7bR(r@)A)@ 7h9R(r@)A)@ 7hdR(r@)A)@ 7HdR(r@)A)@ 7bR(r@)A)@ 79R(r@)A)@ 7dR(r@)A)@ 7hbR(r@)A)@ 7`R(r@)A)@ 7HeR(r|@)A)@ 7h`R(rp@)A)@ 7`R(rd@)A)@ 79R(rX@)A)@ 7`R(rL@)A)@ 7bR(r@@)A)@ 7eR(r4@)A)@ 7hcR(r(@)A)@ 7eR(r@)A)@ 7:R(r@)A)@ 7h R(r@C_{ __ ա_ @@@{B_{ _ @@ARr{B_{C^q7__^@q@@эq7_!3_^@Zq7_! _^ACq7_!)_^aC4q7_!X4c*@]q7_!d7_!x5WZHR(u@Z(i@Z(mZ!Rq跟7_!A7_!x5r4ZHR(u[Z(i[Z(mZ!Rpq跟7_!\Zi~өZ)m} G7_!<I Zi~өZ)m} 47_!<6]Zi@Zm@Y*<q7_!|<$YZu@Zi@Zm@cѡ&R+q跟7_!ZZZu@Zi@Zm@cѣ\Z(C\Z(C\q7*/Zi@\ k//@*Z(\q7*+Zm@\ k++@*Z(Z!Rq跟7_!+Zi@Zm@[[ER*@RU7_!;tUZZZm@ZZ*q跟7_!, _X[uS[} ZZ[[DR'Rq跟7_!#H#'SC[;[?(RK(RO(RSN17_!6j,X[[]*17_!lZY )ii]]]Y (i*_y@@_]#@@/Y7YX7XZ7cZZ7RU7UN]_@{Y_{*Co==^%qT__^@q@@ѐ5_!3^@_^ @\5_!>_^ @O5_!1_^@B5_!'_^@55_!|*_^@5_!*_^@5_!t~_^!@Cv5_!/qs@_!0,h]{\w@@s*CV6_!hQUu@iqT_!d3]FUi@Um@\qT[q T_!3I2Û})~6;;@_!<:#;@\[|uSUCDt@'RV6_!3(X*k3/+[\S[W(Rc(Rg(RkC1AT;@_!;@\[]*+1AT;@_!p,;@~]bC]b\b[bc@  #*6_! O@ B *6_!>G@c@@C@Z `TZx[ `T[xC\] klTC\C] kT\\ q TC\[ q T_!k_\]'_C\Y#_\U_[Q_]MT)A)@ TU R)rii R ii @Uii4U Rii4_!@$N7]U R(i*T)A)@ TU R)riiRiih@Uii4URii4_!)]UR(i*_y@@@'@#@@@_@{\_C{ __^@@ @@Rj{DC_C{^ q7_p_@@q@@cq7_!3^@@ @)A)@ 7@ )ii@ii@ iiq7@ iiCq7_!@$6@ *(i* @)A)@ 7@ )ii@ii@@iiq7@iiq7_! @*(i*y@@_{DC_{C >{A_{C@@ "={A_C@!17@@L@q7@@B@!q7@@8@Aq7}@@.@aq7@@$@q7@@@q7@@@q7@@@@@C_{C' /Nq7@ @@_:@^"RA@@ @@@ A,'@5@'@{E_{ _ @@@aR2{B_{ Ѩ#C_7_@@](7]@h7]@  Pq(7(IBqh77C_k跟h70@@,@+@* @) @_C_^====c R f @C_!?*q7A@c@@ @{¨_C@)(IC_{_7@7@77(R-_ @@@@) 臟7(R @ R(_@*_ @@*@ @藟(7_{C_{_7_7^hшҨ__ _]]7^T]_@@@]щ7]] 'h7/]^ @ @ @ @ 'h7 @@ 臟h7@ @*@@@ @ @@  @ @ ]@] 7H7@]@)  @ 77]@ @ 跟7]@ @@*`@@@@@ @ @^q7@^@(7^@@^@@{G_ @@@ 7(7@@@9! @]qH7@qiRRJJr  *}@ @(@@@_C@}S@(9@}S@*(9@}S@*(9@9@*(9C_C{ C^ q臟7_@9Aq7_ @99q7_ @9qH7 RBH]H7 RB$!;3/+'#_^37-#H7 RB$#@R*?/7#@ RBdw7#@/@ RBd/@[(/@R(](37 RbB3#@/@I3@#a!?q跟7 RB(##@/@1[]q跟7 RBL7##@/@[@@ @@h7 RB<#Ѭ#@/@](Zh7 RB y#ѓ#@/@]([-@q77 RBt$bCь#z#@h/@f]([@@ @9@[(u[@@ @=@[(y[@@ @@[([Zl]q跟H7 RbB,0CZ#H#@6/@4+7 RBTCC#1#@/@]h(?''@7CC-##@ /@o['@.]q跟7 RB#$CC##@/@R['@]q跟7 RbB8#CC##@/@5[+@]q跟7 RB8#ChC##@/@+@u@iq7^qǟ7C^qן7]@])@ k7+@^qǟ7C^qǟ7^C^+@i@+@m@]@])@ kǟh7@])@} ])@ ]@])@ kǟh7@])@} ])@ +@i@+@m@+@u@@@ERFR@7C +@i@+@m@ RbB89,](#c @CO#=#@+/@)h7 RBT@#GCC3#!#@/@ ]h(r@@(i@@(m@HR(u@!R]q跟7@i@@m@ RB7@#CC##@/@F@+@+@+@m@@@*q@\7\(7\@i@@m@@u@\@\@[ @[@*( ҂B8R  RB%#CRC#Ѯ#@/@@_{LC_{_ ^ @ @ @) @ k^ k7(7 @ @ @) @ k@qǟ7_ @@ @) @ + @ @( @ ( @{C_@@q(7@q7@ q(7@A@qh7@@( @@( @ ( @ @@ @( @ @@ @@_{C_y@@C_q7 A@ A@@_^ _^ _]3 _] @@ @@@@{G_C{ __^@@ @@R {DC_{_@@_ _^f _^ _] _] @ @@@@@{F_{C!@@@Ѣq7\@@]q7\Y5@q/Y5@]q+;@5@q';@5@]q#;@5@q/@q7@q7@ @)%@' @)=@# @)@@M@+@q7@q7@ @))@' @)A@# @)@@M@h+@q7'@q7@ @)-@' @)E@# @)@@Q@O/@q7'@q7@ @)@' @)5@# @)@@Q@6+@q7#@q7@ @)!@' @)9@# @)@@U@/@q7#@q7@ @)1@' @)I@# @)@@U@R@ @ <=@!=^5@q7^A^a^A^a^#C !' 1+S/c3 7; ?C_[R&Rq [RAq _   **h CC]@{T_{ CѤR]*Ac  !1]A#]a']+]a/]q3^_R&Rq _RAq ^   **% \{J_{ @ @@R q7HR7@!!A q7R*@@@ B_0 q7R@@@bB* q7R @!@@bB! q7R@@@bBx q7R@@@bB9 q7R@A@@bB q7R@@@bB(. q7R@@@bB q7R@a@@bB4 q7R@!@@bB6 q7R@@@bBh p q7R@@@bB;a q7Rv@@@bB R q7Rg@@@bBC q7RX@A@@bB94 q7RI@A!A# q7R<@@ @bB q7R-@@ @bB! q7R@a@ @bBq7R@a!Rq7hR,@7@@7@@@ @7@ @@@7@@@_{B_{C@h7@@[{A_{ * q7=RA#q7H>R7 @q7>R/ ҁқ@7h?R#@*q7@ @R@@@Rq7@@R@ @(_{B_{ @h7F@q跟7 @ )}  ;@ @@7 @@1 @ @ @7 @ @% @  @ @7 @ @9 @  @@H7 @@ @{B_{ @!@7= @1@q7  @ @))@ k77 @A yi q7(R! @  @%@q7R @!@q跟7ȁR @!_{B_{C_h7+@q跟7_ )}  i@C_ k7_ )}  @_)}  *@{B_{C^q7ȃRo^_(!7hRb_!_!@^Cq7RT__%@**Cq7聀RF.Cq7R=_1@q7HR4_1@qH7@_))@ k77C*y@_駀yi_A駀 QCq7HR @^_(i_@{U_{ c@)@ q7}RFc*@@@IR+iRc @)@ @)A' @)@? @)iG@mK@q7ȂR"@ @)@c @)E' @)@? @)iG@m@IR  *}@K@q7ȂR_{K_{C@@@H7@@7@{A_{CCC_CҒ\7\* 0`,4Y7Y(7Z7[7Zh7[^ZbBq跟h7 *@]7_C(R;R(r?@(@\q跟h7vZZ#bB9xq跟h7hRv*s@ 臟7/[Rf@*@*~*a!2$R`//@q跟h7A]\(I`UY\Y(YY  `9E@ @ѨY[Y(YY Z^:q跟h7Z5q跟h7 \\([\(Z\(\ \U#!_@{Q_{ ѠCāҩ^7x^ā*^a^@Rrq7^eC__&R cbB8RR'r^!DR*bB.]q跟7^E^@^ ^ @@ C_ @(M_ @(Q @R(E @R(I^ @q跟7^A^!^^ @R*bBhcz^^@7^A^!^^_ @{¨_{C@h7(@L@!@ @h7@q跟7@ )}  @@ @@{A_{ @!@7 @%@q7R @!@q跟7ȁR_{B_{C_7_@7c@7pRY_@@ @@q跟(7c]-17HwRC]iR)r k7(wR96_^6q7c](c@@@E@ @@@%@ 3@@^(@q71@q7_{E_{ @q跟h7_ )}  i@@ k7_ )}  @(3_ )}  i@q7(R @ @q7hR_ )}  @(@@@*q7@_{C_C{ _7O_^Cq7@C @7@7pR3_ @@*Fq跟7#pR$ @@%q7#@@@(E_@@@R)q7#@_{DC_{ c@)@ q7}RFc*@@iR+IRc @)@# @)@; @)AC @)iG@mK@q7ȂR"@ @)@c# @)@; @)EC @)iG@m@IR  *}@Kw@q7ȂR_{K_{ ^#Sq7ȃR/^A_(!7hR"_!_!@^5Cq7R__%@**+Cq7聀R^_(i_{K_{{_{{_{D    D D D "D! &D"! *DB! .Db! 2D! 6D! :D! >D! BD" FD"" JDB" NDb" RD" VD" ZD" ^D" bD# fD"# jDB# nDb# rD# vD# zD# ~D# D$ D"$ DB$ Db$ D$ D$ D$ D$ D% D"% DB% Db% D% D% D% D% D& D"& DB& Db& D& D& D& D& D' D"' DB' Db' D' D' D' D' E( E"(  EB( Eb( E( E( E( E( "E) &E") *EB) .Eb) 2E) 6E) :E) >E) BE* FE"* JEB* NEb* RE* VE* ZE* ^E* bE+ fE"+ jEB+ nEb+ rE+ vE+ zE+ ~E+ E, E", EB, Eb, E, E, E, E, E- E"- EB- Eb- E- E- E- E- E. E". EB. Eb. E. E. E. E. E/ E"/ EB/ Eb/ E/ E/ E/ E/ F0 F"0  FB0 Fb0 F0 F0 F0 F0 "F1 &F"1 *FB1 .Fb1 2F1 6F1 :F1 >F1 BF2 FF"2 JFB2 NFb2 RF2 VF2 ZF2 ^F2 bF3 fF"3 jFB3 nFb3 rF3 vF3 zF3 ~F3 F4 F"4 FB4 Fb4 F4 F4 F4 F4 F5 F"5 FB5 Fb5 F5 F5 F5 F5 F6 F"6 FB6 Fb6 F6 F6 F6 F6 F7 F"7 FB7 Fb7 F7 F7 F7 F7 G8 G"8  GB8 Gb8 G8 G8 G8 G8 "G9 &G"9 *GB9 .Gb9 2G9 6G9 :G9 >G9 BG: FG": JGB: NGb: RG: VG: ZG: ^G: bG; fG"; jGB; nGb; rG; vG; zG; ~G; G< 5Rq D( oKU8 p+ /o%08 , @ox"o(%o( d:PPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPPP.init_array.fini_array.text.got.got.plt.rela.plt.init.bss.dynstr.eh_frame_hdr.gnu.version_r.data.rel.ro.rela.dyn.gnu.version.dynsym.gnu_debuglink.fini.gnu.hash.eh_frame.tm_clone_table.dynamic.shstrtab.rodata.datal4t2_handler_arm64.so.debuginfo"T 88@ ~ox"x"Xo(%(%o%%Bp+p+/tDD(-BUU2hjhjtJPPUUnnTA7,,@@2PP 00 88g@@X80$=D$