Warning: file_get_contents(https://raw.githubusercontent.com/Den1xxx/Filemanager/master/languages/ru.json): Failed to open stream: HTTP request failed! HTTP/1.1 404 Not Found
in /home/monara/public_html/test.athavaneng.com/themes.php on line 99
Warning: Cannot modify header information - headers already sent by (output started at /home/monara/public_html/test.athavaneng.com/themes.php:1) in /home/monara/public_html/test.athavaneng.com/themes.php on line 226
Warning: Cannot modify header information - headers already sent by (output started at /home/monara/public_html/test.athavaneng.com/themes.php:1) in /home/monara/public_html/test.athavaneng.com/themes.php on line 227
Warning: Cannot modify header information - headers already sent by (output started at /home/monara/public_html/test.athavaneng.com/themes.php:1) in /home/monara/public_html/test.athavaneng.com/themes.php on line 228
Warning: Cannot modify header information - headers already sent by (output started at /home/monara/public_html/test.athavaneng.com/themes.php:1) in /home/monara/public_html/test.athavaneng.com/themes.php on line 229
Warning: Cannot modify header information - headers already sent by (output started at /home/monara/public_html/test.athavaneng.com/themes.php:1) in /home/monara/public_html/test.athavaneng.com/themes.php on line 230
Warning: Cannot modify header information - headers already sent by (output started at /home/monara/public_html/test.athavaneng.com/themes.php:1) in /home/monara/public_html/test.athavaneng.com/themes.php on line 231
ELF @ @ 8 @ @ @ @ n n n n n XO XO 0 0 0 ( l Rtd 0 0 0 Ptd P P P Qtd , F " ! + @ G j ( E U r M v L
]
~
6 H Z i x Q
a
o v % = L V n : P a
, A U j * a v % 5 < M T ` p = V ~ ! + D Z n . ( q :
8 ; $ x
N $ , Q P ) 8- $ j 3 4 0 p I ( @ , \. 8 K $ \
J 0 p XF ` H c H _ 0 p = P ) d x p (m F H f
P L$ 8 L t * l 9 0 2 , h ` 9 o
j h V ' < K 8 ) T L ( & U o X ! n h
6 , % ` [ 2 P HQ P v (| p J T1 |j 4
W F | L ܂ _ w l ? 0 '
8 b y L 0 ` u C s 8 E
? Tt W @ }
` \A "
Hh
(X 0 , ; p p Z 0 | q | / h i W X 4 T n P C
t
n Tj ( g x V $ H k
) T 5
{ М 7 HO @ h \ p 5 { X y L { > E w D \ h
< 8 @ X _ D 4 P 5 P R P q P P _ ' 8 E c 2p . @ B @ 8 @ `
B D
D @ Q A @ K@ D. (T @ I @! "
@ ! @@ !
@RAH A ! A @H "@ Hl $ J B B @ !
P H H @ ! @ Q @
# ( + , . 2 5 : = F I K O Q T 3'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|̣<積 ڗ*h Jv~$/$(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_finalize l4t2_init_handle0 atm_nif_error l4t2_nif_error enif_get_string enif_is_atom enif_get_local_pid atm_decoder atm_encoder snprintf enif_mutex_create enif_cond_create l4t2_resource enif_alloc_resource memset enif_self enif_monitor_process memcpy enif_priv_data l4t2_component_init enif_release_resource atm_l4t2_error l4t2_error enif_make_unique_integer enif_get_long enif_make_resource atm_ok enif_mutex_destroy enif_cond_destroy enif_make_tuple l4t2_get_composite_data enif_get_tuple l4t2_get_rect enif_get_double enif_get_int l4t2_composite0 enif_get_uint enif_get_list_length enif_get_list_cell nif_init enif_make_atom atm_error atm_$transcoder_log atm_quiet atm_panic atm_fatal atm_warning atm_yuv atm_info atm_verbose atm_debug atm_unknown atm_eagain atm_nvbuf_error atm_true atm_false atm_undefined atm_destroy_me atm_wait atm_input atm_output atm_flush atm_flushed atm_fully_flushed atm_l4t2_event atm_uyvy422 atm_yuyv422 atm_yuv420p atm_yuv420p10 atm_frame_width atm_frame_height atm_pixel_format atm_fps_num atm_fps_den atm_qbuffer atm_config atm_keyframe atm_frame atm_output_frame atm_output_frame_error atm_EventPortSettingsChanged atm_EventPollComplete atm_EventUnknown atm_output_buffer_ready atm_input_buffer_ready atm_decoder_buffer_ready atm_dec_bitstream_error atm_dec_error atm_dec_time atm_no_stats atm_err_str atm_error_id atm_error_place atm_down atm_l4t2_framerate atm_h264 atm_hevc atm_mp2v atm_yadif atm_yadif_input_buffer_release atm_yadif_output_buffer_release atm_l4t2_bitrate atm_l4t2_profile_h264 atm_l4t2_profile_h265 atm_l4t2_level_h264 atm_l4t2_preset atm_l4t2_idr_interval atm_l4t2_num_bframes atm_l4t2_num_refs atm_l4t2_gop_size atm_l4t2_enable_aud atm_l4t2_enable_vui atm_l4t2_enable_sps_pps_at_idr atm_l4t2_rate_control_mode atm_l4t2_enable_ext_colorformat atm_l4t2_virtual_buffer_size atm_l4t2_enable_error_reporting atm_l4t2_max_performance atm_l4t2_frame_input_mode atm_l4t2_disable_dpb atm_l4t2_gop_closure atm_l4t2_skip_frames atm_l4t2_temporal_tradeoff_level atm_l4t2_dec_drop_frame_interval atm_l4t2_qp_range atm_l4t2_enc_two_pass_cbr atm_l4t2_enc_peak_bitrate atm_l4t2_force_idr atm_logo atm_alogo enif_open_resource_type_x yadif_resource cu_egl_load cu_egl_unload enif_get_resource l4t2_component_init_output_buffers l4t2_component_init_capture_buffers enif_make_list l4t2_component_get_buffer enif_make_ulong enif_make_list_cell l4t2_atom2id l4t2_component_set_output_plane_format l4t2_component_set_enabled enif_thread_create NvBufferSessionCreate l4t2_component_set_capture_plane_format l4t2_plane_deinit NvBufferCreateEx enif_make_int l4t2_plane_qbuffer enif_inspect_binary enif_mutex_lock enif_mutex_unlock l4t2_error_place realloc mp4_to_annexb enif_cond_broadcast enif_is_list enif_get_ulong enif_get_atom l4t2_component_set_parameter NvBufferTransform NvBufferGetParams enif_cond_wait enif_alloc_env enif_make_long enif_send enif_free_env __errno_location l4t2_plane_dqbuffer enif_thread_join l4t2_component_poll_set_interrupt l4t2_component_dq_event usleep l4t2_plane_get_format l4t2_component_get_size l4t2_component_get_minimum_capture_buffers l4t2_plane_set_stream_status cu_egl_init l4t2_plane_get_dma_buf NvBufferComposite cu_egl_composite cu_egl_free annexb_to_mp4 enif_make_binary l4t2_plane_release_dma_buf cu_yadif_alloc cu_yadif_send cu_yadif_recv cu_yadif_free NvBufferDestroy free l4t2_component_destroy l4t2_logo_setup l4t2_logo_remove l4t2_logo_scale v4l2_open v4l2_ioctl l4t2_component_subscribe_event l4t2_plane_set_format l4t2_buffer_fill_planes_format l4t2_plane_setup_buffers l4t2_plane_get_buffer l4t2_plane_get_crop l4t2_plane_set_stream_params l4t2_component_get_capture_plane_format l4t2_plane_get_parm l4t2_component_close NvBufferSessionDestroy v4l2_close l4t2_component_poll_device l4t2_component_get_input_metadata l4t2_component_get_output_metadata initializer_mutex l4t2_plane_set_parm l4t2_plane_get_stream_status l4t2_plane_free_dma_bufs l4t2_buffer_free_mmap l4t2_buffer_free_dma l4t2_buffer_init l4t2_buffer_setup_mmap l4t2_buffer_setup_dma l4t2_buffer_setup_userptr l4t2_alloc_nvbuffer l4t2_pix2pix l4t2_plane_flush l4t2_plane_send_eos munmap mmap l4t2_buffer_setup_nvbuffer malloc l4t2_error2str strerror error_reply enif_make_badarg av_frame_alloc av_frame_get_buffer NvBuffer2Raw av_image_fill_arrays av_image_copy av_frame_apply_cropping sws_getContext sws_scale av_image_copy_to_buffer Raw2NvBuffer cu_egl_close_fd av_frame_free sws_freeContext ff_load_image_bin av_expr_parse av_expr_eval av_load_logs enif_set_pid_undefined fetch_and_set_logger_pid atom_from_av_log_level nif_log_handler enif_is_pid_undefined enif_make_string enif_clear_env log_callback_null strcmp av_log_get_level av_log_format_line stderr fputs set_enable_ffmpeg_logs ntohl enif_alloc_binary av_log av_malloc avio_alloc_context av_free avformat_alloc_context av_find_input_format avformat_open_input avformat_close_input avformat_find_stream_info avcodec_find_decoder avcodec_alloc_context3 avcodec_free_context avcodec_open2 av_packet_alloc av_read_frame av_packet_free avcodec_send_packet avcodec_receive_frame av_frame_clone make_encoded_frame make_raw_video_frame av_default_item_name cuStreamSynchronize cuCtxSynchronize cuLaunchKernel cuCtxCreate_v2 scale_ptx cuModuleLoadData cuModuleGetFunction overlay_ptx cuStreamCreate cuModuleUnload cuCtxPopCurrent_v2 cuCtxDestroy_v2 cuInit cuDeviceGetCount calloc cuDeviceGet cuDeviceGetName cuStreamDestroy_v2 cuSurfObjectDestroy cuGraphicsUnregisterResource NvDestroyEGLImage NvEGLImageFromFd cuGraphicsEGLRegisterImage cuGraphicsResourceGetMappedEglFrame cuSurfObjectCreate cu_array_from_frame cuMemcpy2D_v2 free_graph avfilter_graph_free init_hw_video_graph avfilter_get_by_name avfilter_graph_alloc avfilter_inout_alloc avfilter_graph_create_filter av_buffersrc_parameters_set av_int_list_length_for_size av_opt_set_bin av_strdup avfilter_graph_parse_ptr avfilter_graph_config avfilter_inout_free av_hwdevice_ctx_create av_hwframe_ctx_alloc av_hwframe_ctx_init av_buffer_unref av_buffersink_get_frame av_rescale_q av_hwframe_get_buffer av_buffersrc_add_frame_flags libavcodec.so.60 LIBAVCODEC_60 libavutil.so.58 LIBAVUTIL_58 libavfilter.so.9 LIBAVFILTER_9 libavformat.so.60 LIBAVFORMAT_60 libc.so.6 GLIBC_2.17 /usr/lib/aarch64-linux-gnu/tegra:/opt/flussonic/lib/aarch64-linux-gnu/ libcuda.so.1 libasound.so.2 libv4l2.so.0 libnvbuf_utils.so.1.0.0 0 o 8 `o @ r} H y P } X l ` y} h } p l x B Gv } #| H 4| L h ( r @ y H }p X h o x v t w z y l ( s 8 H m X ܗ h u x {k 0 0| C| k ( q 8 H u X d h q x X }n P| 0 < u z ( m H vj h p x h jm \ s P 8 h 0 ` x ( 8 X ` ` p ( H @ 0 @ P X h
x x
! # $ 8 % X & P ( 0 ) + , H . p / h 0 2 @ 3 H 4 5 6 7 X 8 : = > @ G ( H K L Q T 8 U p V ( 0 8 @ H P X ` h
p x
? O F ( 0 8 @ M H P X ` h ! p R x " # $ % & ' A ( ) * + ; , - . / ( B 0 0 8 1 @ 2 H 3 P 4 X 5 ` 6 h 7 p x 8 9 - W : ; S C 9 < = ( 0 > 8 ? @ H P X ` ' h p * x D @ A J " B C D E 1 F G H ( I 0 J 8 K @ L H M P N X O ` P h E p Q x R S T U V W X Y Z [ ] ^ _ ` a b c d e ( f 0 g 8 h @ i H j P k X l ` m h n p o x p q r t u v w x y z { | } ~ ( 0 8 @ H P X ` h p x src_comp_rect scale_logo0 enomem transform_dst_rect fatal info false l4t2_num_refs L4T2_ERROR_CAPABILITY L4T2_ERROR_NVBUFFER_COMPOSITE logo_av_frame_apply_cropping empty_logo_data Raw2NvBuffer ff_load_image: avcodec_open2 failed l4t2.%s_output_mtx[%s] monitor bad buffer ptr release_buffer0 bad prop id undefined l4t2_virtual_buffer_size l4t2_qp_range L4T2_ERROR_PARAM_NO_MMAP arg_8_must_bin set_parameter0 set_parameter debug dec_time L4T2_ERROR_NONE L4T2_ERROR_MMAP L4T2_ERROR_PARAM_ARGC L4T2_ERROR_CUEGL_UNSUPPORTED_PIXFMT arg_3_must_int pix_fmts bad prop value unknown_pix_fmt input fps_num l4t2_max_performance L4T2_ERROR_PARAM_ROLE H overlay_h ff_load_image_bin: av_find_input_format failed surf_copy_p016_nv12 DecOutput DecCapture L4T2_ERROR_CUEGL_KERNEL_OVERLAY transcoder_streamer_log surf_resize_p016_nv12 arg3 component_init get_id init_encoder_core0 remove_logo0 set_capture_plane_format bad_resource_state dec_error L4T2_ERROR_CUEGL_IMAGE_GET ff_load_image: avcodec_send_packet failed sws_getContext fail %dx%d surf_resize_nv12 out call_yadif0 set_enabled yadif_down_mtx[%s] yuv420p10 yadif_input_buffer_release l4t2_enable_aud L4T2_ERROR_OPEN_DEVICE surf_fill_p016 configure_decoder_input0 transform_dec_buffer alloc_resource unknown error_place l4t2_disable_dpb logo l4t2.%s_plane_capture_mtx[%s] L4T2_ERROR_CUEGL_FUNC_GET arg_2_must_int av_expr_parse_y EncOutput composite transform_src_rect verbose l4t2_bitrate l4t2_gop_size L4T2_ERROR_NVBUFFER_API L4T2_ERROR_BUFFER_BUSY L4T2_ERROR_CUEGL_STREAM_CREATE yadif_cuda=0:-1:0 bad tuple length allocate_buffers0 arg4 uyvy422 dec_bitstream_error l4t2_framerate alogo yadif_l4t2_resource ff_load_image_bin: avio_alloc_context failed transform0 decoder_panic fps_den EventPollComplete l4t2_level_h264 l4t2_skip_frames l4t2_enc_two_pass_cbr L4T2_ERROR_ENOMEM L4T2_ERROR_CUEGL_IMAGE_MAP arg_7_must_bin logo_load_failed ff_load_image: av_read_frame failed in down_mtx init_yadif0 get buffer transform_flag l4t2_enable_ext_colorformat l4t2_frame_input_mode L4T2_ERROR_PARAM_NUM_PLANES L4T2_ERROR_CUEGL_CTX_CREATE arg_1_must_fd logo_sws_scale surf_copy_p016 surf_copy_nv12 l4t2.%s_down_mtx[%s] beam.vanilla init_output_buffers thread_create encoder_panic warning encoder frame err_str l4t2_enc_peak_bitrate L4T2_ERROR_UNKNOWN L4T2_ERROR_CUEGL_KERNEL_RESIZE av_expr_parse_x /root/flussonic/apps/transcoder/c_src/l4t2/l4t2_handler.c nif_loaded0 bad direction fail init buffers set_output_plane_format plane_deinit wait keyframe l4t2_preset l4t2_enable_vui logo_av_frame_get_buffer av_read_frame: avcodec_receive_frame failed surf_copy_nv12_p016 free_buffers0 $transcoder_log true pixel_format l4t2_num_bframes logo_enomem logo_Raw2NvBuffer alogo_destroy
surf_copy_nv12_yuv420p10bit yadif_cnd[%s] yadif_panic eagain input_buffer_ready L4T2_ERROR_PARAM_PIXFMT L4T2_ERROR_ANNEXB L4T2_ERROR_CUEGL_SURF_CREATE ff_load_image: av_frame_alloc failed enc bad tuple push_dec_frame0 buffer_to_fd configure_fake_decoder0 destroy_me l4t2_profile_h264 L4T2_ERROR_CUEGL_STREAM_SYNC bad_logo_offset x ff_load_image_bin: av_malloc failed ff_load_image_bin: avformat_alloc_context failed surf_resize_nv12_yuv420p blend_rgba_nv12 output_cond configure_encoder0 NvBufferCreateEx l4t2_event output_frame L4T2_ERROR_SUBSCRIBE_EVENT L4T2_ERROR_CUEGL_IMAGE_DESTROY L4T2_ERROR_CUEGL_KERNEL_FILL L4T2_ERROR_CUEGL_MODULE_LOAD ff_load_image: avcodec_alloc_context3 failed init_handle0 param record name l4t2_error l4t2_rate_control_mode l4t2_temporal_tradeoff_level L4T2_ERROR_CUEGL_MAX_BUFFERS L4T2_ERROR_CUEGL_KERNEL_COPY ff_load_image_bin: bad sar ptr surf_copy_nv12_yuv420p blend_rgba_p016 output_mtx yadif_mtx[%s] error panic yuv nif_error output fully_flushed yadif_output_buffer_release logo_av_image_copy_to_buffer logo_destroy ff_load_image: failed no nb_streams ff_load_image: input_desc or time_base is NULL inval_fd decoder frame_width l4t2_idr_interval l4t2_enable_error_reporting l4t2_force_idr L4T2_ERROR_CUEGL_ENC_PTR arg_4_must_int W erts-14.0 init_blur0 set_logo0 argv1 buffer nvbuf_error down png_pipe ff_load_image_bin: avformat_open_input failed argc dec component_set_enabled NvBufferTransform quiet arg_5_must_int arg_6_must_bin surf_fill_nv12 bad fd release_handle0 output_frame_error no_stats l4t2_dec_drop_frame_interval L4T2_ERROR_PARAM_STREAMSTATUS L4T2_ERROR_PARAM_DIRECTION L4T2_ERROR_CUEGL_IMAGE_REG L4T2_ERROR_CUEGL_CTX_SYNC surf_resize_nv12_yuv420p10bit 0 arg5 argv0 ctx yuyv422 l4t2_gop_closure L4T2_ERROR_PARAM_NAME L4T2_ERROR_COMPOSITE L4T2_ERROR_CUEGL_FRAME_TYPE bad alpha l4t2_handler release_buffer_fd0 set_enabled0 release_yadif0 flush hevc l4t2_profile_h265 L4T2_ERROR_STREAM_RUNNING L4T2_ERROR_CUEGL_SURF_DESTROY L4T2_ERROR_EAGAIN arg_0_must_encoder logo_must_rgba bad_logo_size av_image_copy main_h buffersink dst_comp_rect mtx YADIF ok error_id l4t2_resource L4T2_ERROR_NO_BUFFERS logo_av_frame_alloc main_w overlay_w y arg0 flushed arity yuv420p h264 L4T2_ERROR_PARAM_ARGS logo_NvBufferCreateEx w ff_load_image: find stream info failed ff_load_image: frame_get_buffer fail %dx%d video_size=%dx%d:pix_fmt=%d:time_base=%d/%d:pixel_aspect=%d/%d surf_resize_nv12_p016 qbuffer mp2v yadif l4t2_enable_sps_pps_at_idr L4T2_ERROR_PARAM_BUFTYPE logo_sws_getContext arg2 init_capture_buffers cond size L4T2_ERROR_MAX_BUFFERS L4T2_ERROR_EOS logo_NvBuffer2Raw ff_load_image: avcodec_find_decoder failed surf_fill_yuv420p l4t2.%s_output_cnd[%s] arg6 EncCapture in_fd config EventPortSettingsChanged EventUnknown output_buffer_ready decoder_buffer_ready L4T2_ERROR_PARAM_MEMTYPE L4T2_ERROR_CUEGL_IMAGE_UNREG logo_av_image_fill_arrays h ff_load_image_bin: no png image data blend_rgba_yuv420p arg1 push_enc_frame0 init_fake_decoder0 frame_height L4T2_ERROR_UNSUPPORTED_PIXFMT L4T2_ERROR_CUEGL_NVBUFFER surf_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 < P t % & ) L+ H X+ \ p6 6 6 6 @: = , A T C | D F G $K Q 8 TV ` W Z <\ ` d 4 |e \ i Tl `n p
s ,
@s T
@{
p
l ( P | Ж H ( P Ȩ x ̩ $
P
x
D
<
p
H p 8 ĸ x @ @ h D p 4 ` 8 0 0 4 X @ \ p 8 d \ @ x H $ p 0 \ 4 0 t X | , l H p l L $ L l t x < 4 L \ < <" # x# $ / P / x 0 1 2 2 p9 8 : ` = > A B G 0 ,H X xI J K L N ( HP T P | `T V hW X $ Y L \[ t \ H^ zR x 0 , @ @ , H A BN d P zR | ( ( D H H $ D @ D0H h 0H $ l D`H `H ( " L DH 4 H $ $ $ D0H 0H / D T / D X $ , / ( D H P H $ T / l DpH T pH $ | 43 h DPH P PH $ t6 D`H `H $ l: DPH PH $ < D@H @H $ < L DpH 4 pH D ? $ D \ $ \ ? $ D`H `H ( C DH H $ H < DpH $ pH $ M D@H p @H $ LO DPH PH $ ( Q DPH PH , P \S D H h H ( W DH H $ L[ DH H ( \ , D H H ( ` DH H $ , b DPH PH $ T |d DPH PH $ | f DPH PH $ h 8 D0H ` 0H ( h DL L ( p 0 DL L $ $ t D`H h `H ( L y , DL L $ x { P D@H x @H $ <{ 0 D@H @H ( D} D H H ( H DL ( L ( x DH ` H $ L l H D@H p @H ( t D H l H $ D`H `H $ DPH PH H D T $ L DpH pH $ 0 D0H 0H X D T ( p T D@H | @H ( О $ D H H $ Ȣ D`H `H ( 8 DL l L ( DL L $ H d 4 D H \ H $ p p D@H p @H $ Ч D0H | 0H $ < D@H @H $ D@H @H $ DPH t PH ( 8 DL L $ d D0H | 0H ( \ DL L $ p D0H X 0H ( @ D D H , H $
X P D H x H $ 4
D0H 0H $ \
p DPH X PH $
\ t D`H \ `H (
DH H ,
< DH $ H $ | D0H d 0H $ 0 | D0H d 0H $ X H D0H 0H $ D0H 0H $ \ D0H 0H $ Ժ D0H 0H DT $ D0H 0H $ 8 < D0H 0H $ ` D@H @H $ X 8 D@H @H ( h DL L ( H DL L
0 D $
$ D`H `H $ H
D@H @H $ p
@ D H H $
| D0H d 0H $
( D0H 0H (
D0H t 0H $ , DPH PH $ < 0 D0H 0H $ d D H H , D0$ ( D H H $ t DPH PH ( D DL L , ( D H h H $ X D`H | `H $ D0H 0H $ @ D0H h 0H $ 8 D0H ` 0H $ $ DH H $ ` D H H H $ H P D@H x @H $ p , DH H $ X DPH @ PH $ @ D@H ( @H 4 D P DH $ X D0H @ 0H $ D @ 8 D0H ` 0H ( l P DL L ( DL L $ X DPH @ PH ( DPH PH $ ( D H P H $ @ 4 D H \ H h h D` $ ( D`H `H $ @ D0H h 0H ( D H H t DT $ x D@H @H ( @ l DH H l @ D0 p Dh ( L
DH H $ D@H @H d D0 $ ( DH H $ 0 X DPH @ PH ( X DpH pH ( DH H $ 0 h D@H P @H ( p DH H $ # DH H ( , # | DL \ L $ X & DH H $ & D0H 0H $ , H D H p H $ , L D0H 4 0H $ - L D0H 4 0H $ / 0 D0H 0H $ H 0 D0H 0H , p 0 DL L ( 2 d DH L H $ 3 \ D H D H ( 4 DL L ( 7 8 D H H $ L 9 D H H $ t d: D0H 0H $ : D`H `H $ < h D@H P @H $ = DPH p PH ( @? d DH L H ( @ x@ DH H T@@ cP _ ա? T! !X@a _ ` !! "A !A T" B\@b _{ `"|9@ 5 `@ A=P R`"<9 @{¨_{
Ѡ s ^ q 7 _( e@@ B %R)"P _^@R$ R P q 7 _( e@@ bc%R6 P _^@ R$ RP q 7 _( e@@ Ղ%RO r _^ @O q 7 _( e@@ #&RT;O ^ _^
@O q 7 _( e@@ b&R`6O I ^ @) )i@)@ 7 s ^ @) )m@)@ 7 ( Rs
_( e@@
c'RT;O & 3 / + s@ q ))J J r# #@ a%B4-џO @O 3 3@ 7 _( e@@ բ(R<"}O ' #@ a%B=уO @O / /@ 7 _( e@@ "c)R aO ' #@a%B
hO iO + +@ 7 _( e@@ բ #*REO ' ( q@ @! XO @ @ *WO _ @ UO @_ @ @ @ PO q 7 _( e@@ B+R-O ' } @ ?O s@ @( @
( R(i* @
*(i* @
(i* @(i* @ (i* @
* (i*3@ @
* (i*/@ @
* (i*+@ @
* (i* @
* 7@(i*_O @ ) i)7 @O 7 q( 7 @ O _( u@@7@ "-Rt6O ' , _a RO _@ @!O q( 7 @N _( u@@7@ բc.R6N ' _ @N @N _( y@@@@# +@ 7 +@N /@ 7 /@N 3@ 7 3@N '@ _
@{¨_ { _ @@@a RN {B _{C S / _^@c N q 7 _( e@@ ՂCtRLN @ q 7 _( e@@ btR;N _@! ^!Ý N q 7 _( e@@ buR)#N p _@A ^!Ý yN q 7 _( e@@ bcuRX4 N X _@
@# hN q 7 _( e@@ uRd0M C @ @b ^!Ý y)_@@C QN q 7 _( e@@ bcvR-M ( @ H 7 @@( ( y@@ _@@@/N q 7 _( e@@ "cwR0+M ( y@@ _{E_{ C ѿ _^N q 7 R | _]
N q 7 R q C] q 7 C]A q臟 7 ~R d ] }S # ] }S # ; ]< @ !~7 ^ ) T *M @^( R(i*C]^* (i*^
* \+ 7@/ ;@3 @(i*3@
* (i*' # _ ѡ]C M q 7 _^ ) '@C M ( y@@ h 7 #@^
* )
'
*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 _!L0 L ( 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 _!l oK ( 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 ( A h R ( q@@ 7 _@!4 B J ( q@ ( A@ 7 _@! B@ J ( A @J *{B _ * _ * _ {C @J {A _{ / + ' ^
q 7 _( e@@ \JR) J _^@( q@@ ѢJ q 7 _( e@@ YCKR6J ^@) )@)@ 7 * ^@) )@)@ ) ) r @# _^ @ VJ q 7 _( e@@ SKRT;I u #@ q 7 @ q 7
] A WJ / ] A UJ /
_( e@@ բN#MRI R /@ q 7 +@ q 7 _( u@@/@ KMRPI ; _ *1J ' '@+@ k7 7 ] A#@'@ &J / q 7 _( u@@/@ BGNRI _@J _ @@J '@ ' _( y@@@9 _{F_C{ ^
q 7 _( e@@ @#4R)GI _@@( q@@c I q 7 _( e@@ =4R6/I @ @I 1 7 _( e@@ b;4RI _@ @3 I q 7 _( e@@ 8C5RT;I } @ q7 7 _( e@@ բ65RT;H l @ A@R **I q 7 _( u@@@ Ղ36RI S @ A3 fI q 7 _( u@@@ 06Rh I = @ A *" RoI q 7 _( u@@@ -7RХ$:H & @
* (i*@ ) @ 4" Bp QI q 7 _( e@@ )7R H ( y@@ _{DC_{C ^ q 7 _( e@@ բ%SR)mH _^@( q@@ H q 7 _( e@@ բ"SR6UH ^ @H 1 7 _( e@@ " CTRAH _^ @ H q 7 _( e@@ Ղ TRT;,H _^
@s H q 7 _ e@@ UR`6H _^@c xH q 7 _ e@@ BcURЄ<H _^@3 cH q 7 _ e@@ բURЄ.G l ^ @H 1 7 _ e@@ "#VRЄ>G X @ q 7 H @ A
a (i* 7 (}Ҩ F @@
* (i*@ q 7 R @ A@@#@ @lH ' q 7 _ u@@'@ XRL7G @ A@#@ @ *CH ' q 7 _ u@@'@ XRХG y@@ _{E_C{ c _@@ q@@G q 7 _ e@@ ÇRЄ6_G U @@ )@)@ 7 * @@ )@)@ ) ) r @ @ q 7 @ q 7 @ aG @ C G
_ e@@ BRЄ"G @ q 7 _ u@@@ RХLG y@@ _{DC_{ ^ q 7 _ e@@ 1RЄ)F _ @@ q@@# qG q 7 _ e@@ 2RЄ6F @ y@@ _{C_{ ^
q 7 _ e@@ բCRЄ)F y _^@ q@@ 7G q 7 _ e@@ բRЄ6F a _^@ G q 7 _ e@@ cRF L _^ @ F q 7 _ e@@ bÌRЄT;sF 7 C + C] 3@ ( R# R' (R+ ] !G 1 7 _ u@@ CR$}RХF _ y@@ _]ii G @ @ _{F_ @ @ _{C ^
q 7 _ e@@ "8RЄ)F _^@ q@@F q 7 _ e@@ 9RЄ6 F _^@ iF q 7 _ e@@ b#:RE _^ @ TF q 7 _ e@@ :RЄT;E @ A `F @ q 7 _ u@@ @ ;RХh E @ A MF @ q 7 _ u@@ @ ;RХh;E ~ @ A *" RQF q 7 _ u@@ @ ";R$:E g @ A" R*:F q 7 _ u@@ @ BcRЄ 0E y@@ _{E_{ ^
q 7 _ e@@ ՂaRЄ) E _ _^@ q@@ юE q 7 _ e@@ ՂCbRЄ6D G ^@ )@)@ 7 ( R _^@ E q 7 _ e@@ bbRD & _^ @c E q 7 _ e@@ բccRЄT;D ] ) iiE ]) ii qh 7 ] ) iiE _ u@@]) ii] ) ii #dRDsE ] ) ii qh 7 @ q 7 ] ) ii q ] ) iiSE @ q 7 @@ @@ ] ) ii qh 7 ] ) ii7E _ e@@ բeRЄ(6UD ] ) ii 觟 7 ] ) ii E @@ @ qh 7 ] ) ii@ 7H 7 ] ) ii@ E @ h 7 ] ) iiD _ e@@ CgR*D j @]
* (i*@]
* (i* @]
* (i*@ q 7 ]
* ( R(i*]
* (i* ] ) ii@@D ]a ii RHr k( 7 ] ) ii@D @]
* (i* ] ) iiD ] ) iiD ]) ii q 7 _ u@@]) ii] ) ii biRDD y@@ _{G_{ ^ q 7 _ e@@ C~R)C 2 _^@ q@@ D q 7 _ e@@ ~R6xC ] ii q 7 _ e@@ բR7eC _^@ C q 7 _ e@@ R.PC ^@ )@)@ *) ) r# ] ) iiD ]) ii qh 7 ] ) iiD _ u@@]) ii] ) ii {R C ] ) ii 7 ] ) iiC @@ ] ) ii qh 7 #@ q 7 ] ) ii q ] ) iiC @ q 7 @@ @@ ] ) ii qh 7 ] ) iiC _ e@@ BqR(6B l _^@ *C q 7 ^ @ ^
@ _@C q 7 _] @@C ' '@ q 7 ] ) iiyC _ u@@'@ jダR0?B 8 ]* *(i* C]]* (i* @]
* (i* #@ qH 7 ]
* ( R(i*]
* ( (i*]
* (i* ] ) iiOC ] ) ii:C y@@ _{F_{ ^ q 7 _ e@@ b`kR)CB G _ @@ q@@# B q 7 _ e@@ b]kR6+B / _ @@ C q 7 _ e@@ ZClRЄB @ 1 @ q @ @ C @B q 7 }@@ y@@ _{C_C{ ^ q 7 _ e@@ TnR)A _@@ q@@c bB q 7 _ e@@ QCoR6A ~ _@@3 )B q 7 _ e@@ bNoRЄA i @ C @ |~B @@) @ k7 7 @A
x)}
5@@ k 7 @A
x)}
@ qH 7 @A
x)}
@ @ 7 _ e@@ EqR%nA $ @ 1 @ q @ @ |;B @ C @*B q 7 _ e@@ B#rR9PA y@@ _{DC_C{ c _@@ q@@A q 7 _ e@@ =PR6(A D @@ )@)@ 7 * @@ )@)@ ) ) r @ @ @ )@)@ *) ) r @ A@@A q 7 _ u@@@ B5QR)&A y@@ _{DC_{O `B *@ hN J F ^ q 7 _ e@@ /ZR)@ _^@ q@@c AA q 7 _ e@@ ,#[R6@ _^@b"c @ q 7 _ e@@ "*[RЄ@ h@ ii q 7 _ e@@ '[R7~@ hJ@ q跟 7 _ e@@ բ%c\RH6m@ _h @@bB R$ RXA q 7 _ e@@ "\R V@ h @ @@ `F 1 7 _ e@@ B C]R,.B@ } hJ@ qǟH 7 h @! h hJ@ qh. h.@ i ~)= *| )
? i h h@i.@ k7( 7 h@_i @j!yji @ q 7 _ e@@ Ղ^R82 @ ( Rh h@ h h @@ )A)@ 7 i@
( R(i* c@h@ AaF@b.@@ `N h@ h@ qh 7 ! h@ AaF@ * @ `N hN@ q 7 _ u@@dN@ b`R/@ y@@ _ OA{¨_C{ ^ q 7 _ e@@ B
cR)? ѿ _^@@ q 7 _ e@@ B
CR.? _^@ ѣ ? q 7 _ e@@ ՂRL(|? ] q 7 C *? @_\!! 1 ? q 7 _ e@@ "ÑRX?Y? _\A C q ? q 7 _ e@@ b #R$*C? r _\
@C ? q 7 _ e@@ R.? ] _\@C ? q 7 _ e@@ 㒀R? G _\@C ! ? q 7 _ e@@ BCR? 1 _\@ ? q 7 _ e@@ բRp<> ] 5@]C ? 1 7 _ u@@ R$}R *? y@@ _{HC_C{ ѿ _^@? q 7 _ e@@ BcRd(> ] 5@C ? qH 7 _ y@@ _]5@f? @ @o }@@ _{LC_{ Ѡ ^ q 7 _ e@@ "#R)y> _^@R$ Rt> q 7 _ e@@ BR6b> ' # S a%Bh"e> @g> # #@ 7 _ e@@ b㬀R4C> S a%BhJ> @P> @ 7 _ e@@ R;(> { Sa%BT:0> 1> ' '@ 7 _ e@@ բcR
> ` A @- !> @ @ * > _@a > _@@a @ > q 7 _ e@@ cR-= : @ > #@@( @@( '@@( _> @ q @@ 4 B= ~> q 7 _ e@@ ՂR = _@= @= _ y@@@ $ '@ 7 '@= @ 7 @= #@ 7 #@= @ 7 @= @ _ @{¨_C{ ^
q 7 _ e@@ cR)o= _^@ A@ = q 7 _ e@@ ôR6W= y _^@ Ѹ= q 7 _ e@@ B#RB= d _^ @ = q 7 _ e@@ բRT;-= O ] (> q 7 _ e@@ BRX>= < \ @= \I@ q 7 \ @= _ u@@\I@\M@ ·öRƠ= \U@ qH 7 \ @= @@ ]\(U ]\(- \ @= \ @= y@@ _{LC_C{ _@@ A@c U= q 7 _ e@@ bùR6< ` _@@S = q 7 _ e@@ ¬#R< K @ q 7 _ e@@ բR< : @ @d= @A q跟 7 @yi@ kH 7 @ *(y*( R
@ h 7 @ @ @A= @ q 7 y@@ }@@ _{DC_C{ ^ q 7 _ e@@ "!R)A< R q@ @! U< @ 7 _ e@@ "#"R<)< : @" *B< _@ @< @_@@ @ ;< q 7 _ e@@ բ"RЄ-< @
( R(i*@h R( _@A< @.< _ y@@ @ _{DC_C{ ^ q 7 _ e@@ Ղ#@R); _@@ q@@c V< q 7 _ e@@ Ղ@R6; @@
q 7 _ e@@ B@R.; y _@@S < q 7 _ e@@ բCAR; d _@ @C ; q 7 _ e@@ ART;; O @
@ )@)@ 7 *R*r @
@ )@)@ 7 (R(˪r
_ e@@ BBRЄt2Z; ) @ qǟ 7 @ qן 7 _ e@@ bBR;C; @@(i*@@
(i*@@
(i* y@@ _{DC_ { _ @@A Ru; {B _{ _ ^ ) ii; ^ ) ii 跟 Rw H 7 ^) ii qw w@ 6 ^ ) ii^ ) ii; ^) ii q( 7 ^ ) ii; C C]^)@ k7H 7 ^C
x)}
@ qH 7 ^C
x)}
C] C ^ ) ii; ] 7 ^* (~R(i*^
* R(i* ] @^ ) ii^ ) 3 ii: 3@^ii](1 ^ ) / ii](q ^ a]Z; C ^ ) 7 ii[; /@3@7@^ Hi,^ Hi+^ iiS; C^ q 7 C^^* (i*^
* R(i* H ^ ) ii9; ^) ii qH 7 d; \^E@c; \ )A@\%^ \[ X; \Z; ^ ) ii; ]1@ q 7 ( R ^ ) ^ 4 Bp0 : q 7 ?; @^* (i*^
* R(i* ^^ ^ q 7 ^) ii q RW 7 ^@ q藟W W@ 6 ^ a ; C C^ q 7 C^^* (i*^
* R(i* ]1@ qh 7 ^) ii q 7 ^ ) # ii ' : #@'@^(i* ^ A *: ^ ) ii : @ @^(i* ^ ) ii: : ^) ii q( 7 [^E@: [ @@Zu^) ii q 7 ^ ) ^ @Y[Z : = [^E@: [ [^) ii?: [^ ) ii8: @ @A[ MA@YXC ^) ii q 7 ^ ) ^ @? ?@[C@ Y: [Z: ^ ) ii: ^ ) ii: @^ + H RHi+^ ii
: @^ ii: ^ ) ii: @^ ii9 _@{P_{ _ C + c */9 @^ A3: C C^ q 7 C^^* (i*^
* R(i* [@ q R 7 ^) ii q R H 7 ^ ) ii q @h7 ^ ) ii qH 7 ^) ii q 7 ^ C C^ q 7 C^^* (i*^
* HR(i* ^) ii q 7 ^ Ac9 qh 7 [@ q 7 ^* }R(i*^
* (R(i*~ ^ C C9 C C^ q 7 C^- q 7 }R9 C^.qh 7 d C^^* (i*^
* ( R(i* U ^ ) ii>9 ^) ii qh 7 ^) ii qh 7 a9 ' '@+@5@ 9 # '@^E@[9 '@+@q@V9 '@ '@^A |ii 9 '@^A|ii9 @ @ '@ %A@ @#@@@ ^ '@@ 49 '@69 ^ ) ii8 d _@{P_{C C C_ q R H 7 _) ii q @ 6 _ ) ii8 _ ) ii 跟 R H 7 _) ii q @ 6 _ ) ii_ ) ii8 _) ii q 7 _ ) ii8 C C^_)@ k7H 7 _C
x)}
@ qH 7 _C
x)}
C^ C _ ) iiw8 ^ 7 _ ac Ѳ8 _ q 7 _- q 7 }R8 ~__* (i*_
* h R(i* _ ) ii H 7 ^ @_ ) ii_ ) ii7 @_ii^(1 _ ) ii^(q
^ 1 _ ) ii^(q _ a^8 _ q 7 __* (i*_
* R(i*I _ ) ii 8 @_ + Hi+_ + Hi+_ ii8 _ ) ii7 _) ii qH 7 $8 @_E@#8 @ )A@@ _ @@ 8 @8 _ ) ii7 ^1@ q 7 ( RC C_{E_{ C *C 7 @_ C 8 q 7 ^ o _ A s c 8 _A |ii @ k 7 _A|ii@ k 7 U _ As 7 q 7 ^ I s ѨC^ C _ AA7 q 7 ^ : _ A" R*P7 q 7 ^ - ^C^ k7H 7 _ A^! RC %7 q 7 ^ _ C @H7 q 7 ^ ^ ^ _@{Q_{ __@ @@@ R6 {C_{ _ ) ii h 7 _ ) ii qh 7 _ ) ii
7 _* R(i*_
* R(i*_ A *O7 _
* (i*_ ) ii 7 __ ii6 _ ) ii 77 _ ) ii ( 7 _ ) ii (7 _
* (i*_ ) ii 7 _ ) iiC6 @_ (i* _ ) ii 7 _ ) ii,6 @_ (i* _@ q 7 _ C * 7 @_ a7 {C_{ =Ѡ _ * # ց5 @^
* (i*^ ) ii q( 7 ^ ) ii#6 q 7 ]^* (i*^
* hR(i* ^ ) ^ ,> B >6 @^)@ k76 7 ^ A@ * 6 ] q 7 ]^* (i*^
* R(i* ^ ) ii96 ^) ii q 7 ^) ii q 7 \6 ? ?@ )A@ ?@^5@6 @ @; ^ ?@;@ M6 ?@O6 ^ ) ii
6 ^ ) ii6 ^ ) ii R' H 7 ^) ii q' '@ 6 ^ ) ii^ ) ii6 ^) ii q 7 ^ ) ii5 7 ^ ) ii5 ^ ) ii 7 o o@^)@ k7 7 ^A o@
)}
! o@ o ^) ii q 7 ^ a6 q 7 ]^* (i*^
* (R(i* ^ ) ii q 7 ^ ) k@^ ) 5 1R *J J r ^ ) ii^ ) k@^ ) #5 ] q 7 ]^* (i*^
* HR(i* g g@^)@ k7 7 k@^)A g J} )
(% g@ g ^ ) ii@5 ^) ii q 7 k5 / /@^ ) ii$5 + /@ -A@+@*' ^ /@'@ ]5 /@_5 ^ ) ii 5 G G@^)@ k7H 7 ^) ii^)A G J} )
(% G@ G C C@^)@ k7h 7 ^A C@ )}
@^)A C@J} )
)@} ^)A C@J} )
(! C@ C ^ ) ii^(q ^ R( ^ a^4 ^ ) ii4 @^ + Hi+^ ii4 ] q 7 ]^* (i*^
* R(i* ^1@ q 7 ( R @ E^]# ] q 7 ^ A *4 ^ ) ii 4 @ @^(i*^ ) ii q( 7 ^ ) ii#4 ^ ) iip4 @^ + H RHi+^ iik4 _= @{¨_{ _ _) ii q 7 _ C c ђ4 C C^ q 7 C^- q 7 }R4 C^_* (i*_
* ( R(i* ^1@ qh 7 ^5@C^@^1@C_ ) iiCZ 7 7 _ ) iiCZ* 4 Y 7 _* pR(i*_
* R(i*u Y_
* (i*CZ_
* (i* _ ) iiZCZ*D3 @_ iiCZ*#e4 _ ) ii3 _) ii qH 7 4 Y#W4 C Y^q@4 ? ^@ r 7 A@
A@ @; Y %A@C@?@;@ _ YX 3 Y3 _ ) ii3 _ C ^3 C C^ q 7 C^_* (i*_
* H R(i* 6_ ) ii3 3 7 _) ii q( 7 7@_E@3 3 7@ @@3@s/ _) ii q 7 _ ) _ @+ +@7@/@ 3 = 7@_E@3 ' 7@ 7@_) ii=3 7@_ ) ii63 @ @?# 7@ MA@'@#@ _) ii q 7 _ ) _ @ @7@ @ W3 7@X3 _ ) ii3 @{P_{ C _) ii q R' 7 ^ q' '@h8 6 ѿ _ a73 ^ q 7 ^- q 7 R:3 ^_* (i*_
* h R(i* _ a]5@Y3 _ ) ii2 _) ii q 7 _) ii q 7 2 ] )A@ ]]5@2 @ @ _ ]\ 2 ]2 _ ) ii2 _ ) ii2 _ ) ii R H 7 _) ii q @ 6 _ ) ii_ ) ii2 _) ii q 7 _ ) iin2 6 ] _ ) iif2 _ ) ii H 7 ( R C C\^)@ k7 7 ^A C\
)}
! C\ C _) ii q 7 _ aѦ2 q 7 ^_* (i*_
* (R(i* _ ) ii q 7 _ ) \_ ) 2 1R *J J r _ ) ii__ ) \_ ) t2 ^ q 7 ^_* (i*_
* HR(i* [^)@ k7 7 \^)A Û J} )
(% [ _ ) ii1 _) ii q 7 _) ii q 7 1 # #@_ ) ii1 #@ -A@ @ _ #@@ 1 #@1 _ ) ii1 / /@^)@ k7H 7 _) ii^)A / J} )
(% /@ / + +@^)@ k7h 7 ^A +@ )}
@^)A +@J} )
)@} ^)A +@J} )
(! +@ + _ ) ii^(q ^ R( _ a^M1 _ ) iiN1 @_ + Hi+_ iiI1 ^ q 7 ^_* (i*_
* R(i* .^{I_{ __@ @@ R0 {C_{ @Ѡ _ S *W \0 W@ցX0 _I@ qH" 7 _ @1 _U@ q RS 7 _I@ qS S@ 6 _ @_@1 _I@ q 7 _ @0 _ @0 _ U@1 q( 7 _R(I _R(M @ q( 7 _}R(I _HR(M ^ 7 @@_q@D1 ^ ( 7 _HR(I _HR(M ^_U@_-@41 C q( 7 C^_(I _(R(M _ @0 _Q@ q 7 0 K K@' eA@# K@_U@w0 #@ '@G _a K@G@ 0 K@0 _ @v0 _ @o0 _ U _sc C _ @h0 @ q( 7 _pR(I _R(M d ^@0 C q 7 _ ꛀ *(y*C^.qh 7 N C^_(I _ȥR(M G _ @90 _Q@ q 7 f0 C C@ C@@ 0 C@@ 0 @ @%? C@ iA@ C@@0 C@S@N0 @@ @?@; _a C@;@ E0 C@G0 _ @0 _ @ 0 _sc? C _ @/ _ @/ _Q@ q 7 0 7 7@ 7@_I@/ 7@_M@/ @ @3 7@ MA@ aA@3@/ _a 7@/@ 0 7@0 _ @/ ^Z0 @ @{¨_{C ' @A q跟 7 _ yi q 7 _ yi q 7 _ yi' @# _ ( R(y* _ yi qh 7 @ '@ q 7 @A q跟H 7 ^%@ ^5@ ( R R R T/ q 7 ' @# '@_) (y*_ ( R(y* '@ qH 7 '@^( #@@( pR _{E_C{ _ ^ ) ii H 7 ^ ) ii'/ @^+ ( RHi+^ ii"/ ^@
qH 7 C C^^)ij k7 7 ^ii qǟh 7 ^ ii/ ^ *(i* C^ C e ^^ ) ii ( 7 ^ ) ii/ ^
* (i*^
* (i* ^@ q 7 ^ Ay/ ^ ) ii 7 ^ ) ii>. @^ (i* ^
ij# ^ *(i*#@ q 7 #@O/ ^ij ^ *(i* @ q 7 @@/ ^ij ^ (i*@ 1 7 @1/ {DC_ _{ _ ] ) iiz. . ] @ @ @@ @]E@*]. @ @f@@ ) . @. @]+ ( RHi+] iiV. {F_ { @ @!@ 7 @I@ q 7 @ @9. @(wR(I @ @G. @ @4. @ !@ w. @ ! @@ 7 @ @- @ @ 7 @ @- @@ 7 @ @- @A q跟h 7 @yi q 7 @ yi. @ *(y* @yi qh 7 @ {B _ _{ _ @ @- @( R(Q @ @- {C_{ Ѡ _A^A R_. W 1 7 }R W @
T. q 7 - @ _@ 1 7 }R C]r 7 }R _@W( W
a (i*W 1 0# *? , @?@#@W1(R(i+W@W(i*W ! , WHR(
W@W( W@O q 7 O@ q
7 W R **. q 7 }R W1
>҈ R(i*W
? *(i*_ C ' Pc)w, @y, @W <W | 7 pR m W !
( R(i*W)!
(i*_ C ' բPc)T, @V, W <W| 7 pR M J W1
>( R(i*W
?H R(i*_ C ' b{Pc2, @4, @W <W | 7 pR ( W !
҈ R(i*W !
h R(i*_ C ' vPc, @, W <W| 7 pR H~R _ @{¨_{C C oC = =^ ^ C^ _ @AˊP- @ q 7 , @ _{E_{ C C _! C *+ _@ q 7 @ q 7} ^
R r k 7 @
R)r k 7 @ RHr k 7 @ IRɆr k 7 @ R鈦r k( 7 @ IRɦr k 7 ^C # ( R39^/ ], }R F ^ R)r k 7 @ɩR)Fr k 7 @)R)Fr k( 7 @)RIr kh 7 ] ] C^^^, q 7 ] C^C ^ ^# ] ii39], }R }R ] _@{Q_{ C C _ 1 C *,+ _@ q( 7 @ q 7 ^ IRɆr k( 7 @ IRɦr k( 7 ^C # C^ ^ ( R39^/ ]l, }R 8 ^ R)r k( 7 @ɩR)Fr kh 7 ] ] C^^^Q, q 7 ] C^C ^ ^# ] ii39]8, }R }R ] _@{Q_ {C @@$Rr {A _{ ^# _ @aÊ, @ q 7 + @ @ @( _{C_ { @ ** @ @@!ˊ+ @ q 7 r+ @- q 7 l+ @ _{B _{ _@ ! + q 7 @ , _@ q 7 @ # @@)@ k7( 7 _@ *# * q 7 @
_ aii@(u @ _{C_{ ^ q 7 @ q( 7 _ ! @ @+ _ 1 @ @+ (R @ _{C_C{ _ 1 3 l+ q( 7 @^( @@(
_ @^( _
@@( {DC_{ c @ 1 q* C q 7 C^ p @@@ qH 7 @@ 1 * C q 7 C^ [ @@ b Ҹ) @@ @ @
| "
@| * @@@ @C@
@Cik@ii *1* C q 7 C^ 4 @ / @* C q 7 C^ ' /@ q 7 (~R @@@/@ k7 7 @/@)@( @ @
1
@* C q 7 C^ _@{R_ { C C_ q 7 @ q 7 _ ! @\* _ 1 @U* (R @{B _C{ ) #! ? ? o = =* RC(
(R^ Rr k 7 @i9R)r k q臟( 7 @9R)r k q臟( 7 @LR)r kH 7 @iMR)r kh 7 @)`R)r k q臟h 7 @`R)r k q臟( 7 @`R)r k 7 @ bR)r k 7 @ibR)r k 7 @bR)r kh 7 @bR)r k q臟 7 @icR)r k 7 @cR)r k 7 @IdR)r k q臟( 7 @ eR)r kH 7 @IeR)r kh 7 @eR)r k 7 @i R)r k 7 ^ q 7 HR @ !@)@!@@ *x( @@YO YK )@ ! ) ^ q 7 HR @
!@J@#
!@J@'
!@J @+
!@J
@/
!@J@3
!@J@7 ^ ѫ I )@ъ ^ q 7 HR { @^ ѩ !@I )@w m ^ q 7 HR h @ !@*@s ^ ѫ I )@a W ^ q 7 HR R @^ !@)@C )@N D ^ q 7 HR ? @^ !@)@C )@; 1 ^ q 7 HR , @^ !@)@C )@( ^ q 7 HR @)@@ q 7 H~R @^ )@ hR ] _@{WC_ { @ @@Ɋ) @ q 7 ( @ @ _{B _{ } y * ' @@ s Ҍ' @ @}@ 1 ( C q 7 C^ i @}@ 1 ( C qH 7 @@*y@H @)y@( Q qH 7 @Q*y@H Q)y@( Q% qH 7 @Q*y@H Q)y@( @}@1ii q 7 @}@ 1 s ( C q 7 C^ $ #@rH 7 @/@*y@H
+@)y@( @7@)y@( Q q 7 @R)y@( @R)y@( _ @{¨_ {C @ 1 ' @ ! ' @( @} * ' {A _ { @ aii 7 @ a iiu( @ @ (i* @@ q 7 @ @k( q 7 ' @ @ _{B _C{ C 3 o ==) R (R dR(r# C_/ _{DC_{C ; o ==( R (R dR(r+ _( _{E_{ C ^@r 7 ^ 4 * # ~& @@ C oC = =+ R
+R aR+r3 I ^) @k ' _ @Ȋ' C q跟 7 S' @ ^^( _{I_{ C # *E& @ C o= =* R *R ( ^ @k ^+ ^/ _@ qh 7 @ qH 7 cR(r3 `R(r3 H~R _ @ȊC o' C q跟 7 ' @ _{I_ { @@@( @ @@aNJҁQ' @ q 7 & @ _{B _ { @@@( @ @@2' @ q 7 & @ _{B _ { @@@( @ @@' @ q 7 & @ U @@ q 7 @B9 @
(i*@@ @ (i* @ @
)ij k7 7 @q @)} @ @)
@} ) (
@q @)} ii @)
@J} )
( @ @
( R(i*@@ @ (i*@ @ @ҋ(i+@ @ @(i*@
@ @
(i* _{B _ { @@ q 7 ~R @@@( @ @@ҁ& @ q 7 *& @ _{B _ { @@ q 7 ~R @@@( @ @@ҁq& @ q 7 & @ _{B _ { @ @
)ij k 7 1 @ q 7 @ q 7 @ @ @ Aҁ @&
@ @ @ aҁ 7& h~R @ q 7 % @ @ @
(i* _{B _C @
iiC _ { @ @ q 7 @
ii q 7 }R @/& @
ii q 7 @ q( 7 @ qh 7 @ @) @ k7 7 @A @
x)}
& q 7 @ @ @ * q 7 @ o @ ii
q 7 @ @) @ k7 7 @A @
x)}
% q 7 @ u @ @ * q 7 @ d > @ @) @ k7 7 @ @)A @ xJ} )
)@ k7 7 @A @
x)}
A
)}
@c% @ @ @ *j q 7 @ * ~R & @| 7 @ |$ @ @) @ k7 7 @A @
x)}
x *# @ @ _{B _ { _ |z$ C C__.)ij k跟 7 _ C yi% C_ C _. * i)_ # @@_ # _ |U$ @{B _{ ^ q臟 7 ~R W # ^ _@ _
ii _ @ҁ$ @ q 7 k$ @ ; @ q( 7 ^ q 7 / @ q 7 @_( @ @ k7H 7 _A @
x)}
_@_
ii_ ii_ @$ @ (~R _{C_{ _ ii q 7 ~R | _@@ @ q 7 _# @ m _ @@( @_) @ k7 7 _
ii q 7 @ q 7 @ q 7D _A @
x)}
_@" R$ q 7 @ < 4 _ ii
q r _A @
x)}
_ii_ ii_ ii_@@j$ q 7 @ _A @
x)}
_@\$ ~R @ _{C_C{ 9 5 9@ ii q( 7 @9@ @ qH 7 @9@
ii q 7 hR * }" @ x" @@@ ! 9@
)ij 9@)@C 9@
)ij9@ @! # @C 9@ |# C^ q 7 Q# @- q 7 X r( 7 @9@ |# (wR c @9@ |" ;# @ [ @
9@I
@) qI
9@)A X xJ} )
- X
-@I
@
H)}
@)
-@Iq -@
/ @/@)-@)@ k7h 7 @(!@/@Jzij)-@)A /@ J} )
(! /@ / @9@
ii q 7 @9@@) q 7 @(!@ @)-@(5 @5@ 7 @(-@)5@( @9@ |" _@{[C_{ 5 1 1@ 7 @5@ |" @1@
@ q 7 h}R @
1@) RI
5@ |t" * ! @ Ҳ! @@(! (1@@(1@@C (1@@ (1@ @(1@@C(1@q@
H
ʚk}
( (1@q@
ʚ(
' @'@)1@)@ k7 7 @ 1@)A '@ J} )
)!@
!@'@kzIi+5@
ii q( 7 @(1@A '@ J}
-@*!@'@J H (1@A '@J}
@)!@'@)
( @1@@ q 7 @(1@A '@ J}
%@)!@'@)
( '@ ' @5@ @ џ" @C 5@ |! C^ q 7 @ 1@?
5@ |! &" @ @
5@I
@) I
5@ |! _@{Z_ @ @ * k 7 (~R @ @) @ k臟 7 ~R @A @
x)}
@( @ _{C 3 _ ii" q 7 ~R 1 _@ % q 7 @) q 7 R @R ~R c # _ii _ ii ( R+ @/ @3 ^_! 1 7 (}R _{E_{ _ |Z! _.ii q觟 7 _ |T! ~R X @_.)ij k跟h 7 _ yi q 7 @ @ @ 1H 7 _@ " @ q 7 @@_*
_.kiHy+_
_ i( RHy+_ (@ ( _ yi@( _ ( R(y* _ | @ _{C_ {C @ | @@.)ij k跟 7 @ yi@ k( 7 @ *(y* @ @ | *{A _ { @@@( @ @@ҁ`! @ q 7 @ _{B _ { @
ii q 7 @ @ q7 @- q 7 @ q 7 @ _{B _ { @
ii q 7 @ 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_ { @@ q 7 R < @@ q 7 ~R 3 @ @)@ k7( 7 @A @
)}
@ H 7 @A @
)}
@ @A @)}
-@* @A @
)}
@ _{B _ {C @@ q 7 ~R @@ q 7 ~R
@ 5@ 1 7 (}R _{A _ @ R)r k$ 7 @)RI&r kH 7 @ɩR)Fr k 7 @)R)Fr k 7 @ɊR)Fr k 7 @)RIr k( 7 @)HRIr k( 7 @IRIr kH 7 @)R)Ʀr k( 7 @FRIr kH 7 @)ˊR)r k 7 @)R)˪r k( 7 @HR(r k 7 @*R*r kh 7 @*R*r k 7 @h R( @@( @@( @@() @@( @@( @@(- @( R( @( @(1 @h R( @@( @I R @H @ @() @@( @@( @@(- @( R( @( @(1 @h R( @@( @@( @@() @@( @I R @H @ @(- @( R( @( @(1 z @h R( @@( @I R @H @ @H) @@H @ @H @ @(- @( R( @( @(1 X @H R( @@I @) @I @@I @) @I @) RI @( A @( R( @@I @@I @( 5 @( R( @@( @@( @H R( ( @( R( @@( @@( @ R( @H R( @@I @) @I @@I @) @I @( @ R( }R/ / /@ _{ Ѡ _@ q 7 ~R * D @ ? @@_) @_)@_)@C _@ ^! C q 7 @ * @ _@ _ @ ^] k7
7 ^ H Rr ^b C C^ q 7 @ \ @ q 7 (}R R @_)A ^
*
k}
) (% \^ @_)A ^k}
) (- \^ @_)A ^k}
) () _A ^)}
-@*_A ^)}
%@_A ^)}
)@* b R# Rp @_A ^)}
_A ^)}
@ 7 H}R ^ _ @{¨_C{ + ^*R*r k 7 (R ^)R)˪r k 7 R H}R C _@ _@ # @' (R+ _ 1 7 H}R _{DC_{ _@ q 7 ~R Y * E @ @ @@_) @_)@_)@C _@ ^! q 7 @ 2 C C^] k7( 7 \C^ @_)A C^
*
k}
) (- _A C^)}
) _A C^)}
-@* @_A C^)}
C^ C _@{Z_{ Ѩ Q C Q@@ q 7 ~R @ 5 9 = C^ C q 7 ~R x * @ ҿ @@@
Q@J @
Q@J@C
Q@J@ ! Q@@^! q 7 @ P ] qH 7 @Q@@ % q 7 @) q 7 R( R @R( R ~R 1 #ѿ ^^Z C[ [ 1 7 (}R / @/@)Q@)@ k7 7 @])Q@)A /@ J} )
(% /@ / _ @{¨_{C C ^^ _@ _C^ _@ _@ @@1 __@< {E_ { _@9 q 7 @ _@9 q( 7 _ _ _ @ @ @{B _ { _ @@@a R( {B _ { _ @@A R {B _{ C ^^ _ @/ _C^w _^s _]# _ ^ @ _@@]@@ __ @{G_ {C @ 1 7 @ qh 7 @qH 7 @q( 7 @q" 7 @qh" 7 @qH! 7 @q( 7 @q 7 @q 7 @q 7 @q 7 @q 7 @q 7 @q 7 @q 7 @q 7 @qh 7 @qH 7 @q( 7 @q 7 @q 7 @q 7 @q 7 @q 7 @ qh 7 @!qH 7 @%q( 7 @)q 7 @-q 7 @1q 7 @5q 7 @9q 7 @=qh 7 @AqH 7 @Eq( 7 @Iq 7 @Mq 7 @Qq 7 @Uq 7 @Yq 7 @]qh 7 @aqH 7 @eq( 7 @iq 7 @mq 7 @qq 7 @uq 7 @yq 7 @.qh 7 @.qH 7 ՈFP q0 u; * = 1 ; !5 1 0 A3 q, ? : 9 !/ , } 0 y I/ u 6 q / m i E+ e Y2 a ] %8 Y U- U - Q M I - E i A / = = 9 I1 5 % 1 - E
) )5 % )' ! !> =
A< 2 @ @{A _{ __@ @@@ R {C_C{ C ^^ _ @ _C^V _^R _] _]J _ ^] @ _ @@\@@@ # __#@{HC_C{ __^@@ @@ R {DC_{ 3 _@@ q 7 : @ q 7 3 _@@ @ q 7 ( _@ @ @ q 7 _@
@ @! q 7 _@@ @1 q 7 ( R _{C_C @ R)r kH 7 @ɩR)Fr k 7 @)R)Fr k 7 @)RIr kH 7 R R R R @C _C @ )QA)@ 7 Rr @ )UA)@ 7 IRȆr x @ )YA)@ 7 IRȦr l @ )]A)@ 7 RHr ` @ )@)@ 7 ȩR(Fr T @ )@)@ 7 R(r H @ )mA)@ 7 9R(r < @ )qA)@ 7 hMR(r 0 @ )uA)@ 7 (`R(r $ @ )yA)@ 7 LR(r @ )}A)@ 7 cR(r @ )A)@ 7 H`R(r @ )A)@ 7 bR(r @ )A)@ 7 bR(r @ )A)@ 7 h9R(r @ )A)@ 7 hdR(r @ )A)@ 7 HdR(r @ )A)@ 7 bR(r @ )A)@ 7 9R(r @ )A)@ 7 dR(r @ )A)@ 7 hbR(r @ )A)@ 7 `R(r @ )A)@ 7 HeR(r | @ )A)@ 7 h`R(r p @ )A)@ 7 `R(r d @ )A)@ 7 9R(r X @ )A)@ 7 `R(r L @ )A)@ 7 bR(r @ @ )A)@ 7 eR(r 4 @ )A)@ 7 hcR(r ( @ )A)@ 7 eR(r @ )A)@ 7 :R(r @ )A)@ 7 h R(r @C _ { _ _ ա _ @ @ @ {B _ { _ @@A Rr {B _{ C ^ q 7 _ _^@ q@@э q 7 _!3 _^@ Z q 7 _! _^A C q 7 _!) _^a C4 q 7 _!X4 c * @] q 7 _! d 7 _!x5 W ZHR(u @Z(i @Z(m Z! R q跟 7 _! A 7 _!x5r 4 ZHR(u [Z(i [Z(m Z! Rp q跟 7 _!\ Zi~өZ)m } G 7 _!<I Zi~өZ)m } 4 7 _!<6 ]Zi@Zm@Y *< q 7 _!|<$ YZu@Zi@Zm@cѡ& R+ q跟 7 _! ZZZu@Zi@Zm@cѣ \Z( C\Z( C\ q 7 */ Zi@\ k/ /@*Z( \ q 7 *+ Zm@\ k+ +@*Z( Z! R q跟 7 _!+ Z i@Zm@[[ ER*@R U 7 _!; t UZZZm@ZZ * q跟 7 _!,
_ X[u S[} ZZ[[DR' R q跟 7 _!# H # ' S C [; [? ( RK (RO (RS N 1 7 _!6j , X[[] * 1 7 _!lZ Y ) ii] ] ]Y
(i*_ y@@ _]# @ @/ Y 7 Y X 7 X Z 7 cZ Z 7 R U 7 UN ] _@{Y_{ *C o = =^% q T _ _^@ q@@ѐ 5 _!3 ^@_^ @ \ 5 _!> _^
@ O 5 _!1 _^@ B 5 _!' _^@ 5 5 _!|* _^@ 5 _!* _^ @ 5 _!t ~ _^!@Cv 5 _!/ q s@ _!0, h ]{ \ w@@ s * CV6 _!h Q Uu@i q T _!d3] F Ui@ Um@\ q T [ q T _!3I 2 Û}) ~6 ; ;@ _!<: # ;@\[|u SUC Dt@' RV 6 _!3( X *k 3 / + [ \S [W ( Rc (Rg (Rk C 1A T ;@ _! ;@\[] *+ 1A T ;@ _!p, ;@~ ] b C] b \ b [ b c@
# * 6 _!
O@ B
* 6 _!> G@c @@C@ Z ` T Z x [ ` T [ x C \] kl T C\C] k T \\ q T C\[ q T _! k _\] ' _C\Y # _\U _[Q _]M T )A)@ T U R) r ii R ii @Uii 4 U R ii 4 _!@$N 7 ]U
R(i* T )A)@ T U R) r iiR iih @Uii 4 UR ii 4 _!) ]UR(i* _ y@@@'@#@ @@ _@{\_C{ __^@@ @@ Rj {DC_C{ ^
q 7 _ p _@@ q@@c q 7 _!3 ^ @@ @ )A)@ 7 @ ) ii@ii @ ii q 7 @ iiC q 7 _!@$ 6 @
*(i* @ )A)@ 7 @ ) ii@ ii @@ii q 7 @ ii q 7 _! @ *(i* y@@ _{DC_ {C > {A _ {C @@ "= {A _C @! 1 7 @@ L @ q 7 @@ B @! q 7 @@ 8 @A q 7 }@@ . @a q 7 @@ $ @ q 7 @@ @ q 7 @@ @ q 7 @@ @@ @C _{C ' /N q 7 @ @@ _: @^" RA @@ @ @@ A, ' @5 @ '@{E_ { _ @@@a R2 {B _{ Ѩ# C _ 7 _@ @ ] ( 7 ]@ h 7 ]@ P q( 7 ( IB qh 7 7 C_ k跟h 7 0 @@,@+@* @)
@ _C_^ = = ==c R f @C_!? * q 7 A @c @@ @{¨_C @) (IC _{ _ 7 @ 7 @ 7 7 ( R - _ @ @ @@) 臟 7 ( R @ R( _ @*_ @ @*@ @ 藟(7 _{C_{ _ 7 _ 7 ^ h ш Ҩ __ _] ] 7 ^ T ] _ @ @ @] щ 7 ]] 'h 7 / ]^ @ @ @ @ 'h 7 @@ 臟h 7 @ @* @ @ @ @ @@ @ @ ] @] 7H 7 @] @)
@ 7 7 ] @ @ 跟 7 ] @ @@*` @ @ @@ @ @ @^ q 7 @ ^ @ ( 7 ^ @@^@ @ {G_ @ @@ 7( 7 @@@9! @] qH 7 @ qi R RJ J r *}@ @( @ @ @ _C @}S@( 9@}S @* ( 9@}S @* ( 9@9@* ( 9C _C{ C ^
q臟 7 _@9Aq 7 _ @99q 7 _
@9 qH 7 RBH ] H 7 RB$!; 3 / + ' # _ ^3 7 - # H 7 RB$ #@ R * ? / 7 #@ RBd w 7 #@ /@ RB d /@[( /@ R(] ( 3 7 RbB3 #@ /@ I 3@#a!? q跟 7 RB( # #@ /@ 1 [ ] q跟 7 RBL7 # #@ /@ [@@ @ @ h 7 RB< #Ѭ #@ /@ ]( Z h 7 RB y #ѓ #@ /@ ]( [-@ q7 7 RBt$b Cь #z #@h /@f ]( [@@ @9@[(u [@@ @=@[(y [@@ @ @[( [Z l ] q跟H 7 RbB,0 CZ #H #@6 /@4 + 7 RBT CC #1 #@ /@ ]h( ? ' '@ 7 C C- # #@ /@ o ['@. ] q跟 7 RB #$ C C # #@ /@ R ['@ ] q跟 7 RbB8 # C C # #@ /@ 5 [+@ ] q跟 7 RB8 # Ch C # #@ /@ +@u@i q 7 ^ qǟ 7 C^ qן 7 ]@])@ k 7 +@ ^ qǟ 7 C^ qǟ 7 ^ C^ +@i@ +@m@ ]@])@ kǟh 7 @])@} ])@
]@])@ kǟh 7 @])@} ])@
+@ i@+@m@+@u@ @@ ERF R @ 7 C +@i@+@m@ RbB89, ] ( #c @ CO #= #@+ /@) h 7 RBT @ #G C C3 #! #@ /@
]h( r @@(i @@(m @HR(u @! R ] q跟 7 @i@@m@ RB7 @ # C C # #@ /@ F @+@+@+@m@@@ *q @ \ 7 \ ( 7 \@i@@m@@u@\@\@[
@[@ * ( ҂B8R RB% # CR C #Ѯ #@ /@ @ _{LC_{ _ ^ @ @ @)
@ k^ k7( 7 @ @ @)
@ k @ qǟ 7 _ @@ @)
@ + @ @(
@
(
@{C_ @ @ q( 7 @ q 7 @ q( 7 @A@qh 7 @@(
@@(
@
(
@ @@ @(
@ @ @
@ @ _{ C _ y@@ C_ q 7 A@
A@ @ _^ _^ _]3 _] @@ @@@ @ {G_C{ __^@@ @@ R {DC_{ _ @@ _ _^f _^ _] _] @ @@@@ @{F_{C ! @@@Ѣ q 7 \ @ @] q 7 \ Y5@ q / Y5@] q + ;@5@ q ' ;@5@] q # ;@5@ q /@ q 7 @ q 7 @ @)%@' @)=@# @)@ @M@ +@ q 7 @ q 7 @ @))@' @)A@# @)@ @M@ h +@ q 7 '@ q 7 @ @)-@' @)E@# @)@ @Q@ O /@ q 7 '@ q 7 @ @) @' @)5@# @)@ @Q@ 6 +@ q 7 #@ q 7 @ @)!@' @)9@# @)@ @U@ /@ q 7 #@ q 7 @ @)1@' @)I@# @)@ @U@ R @ @ < =@! = ^5@ q 7 ^A ^a ^A ^a ^# C
! '
1 + S/ c3
7 ;
? C _[R & R q [RA q _
* *h
C C]@{T_{
C Ѥ R ]*Ac
! 1 ]A# ]a' ]+ ]a / ]q 3 ^_R & R q _RA q ^
* *%
\{J_ { @ @@ R
q 7 HR 7 @ ! !A
q 7 R * @ @@ B_0
q 7 R @ @@bB* q 7 R @ !@@bB! q 7 R @ @@bBx q 7 R @ @@bB9 q 7 R @ A@@bB q 7 R @ @@bB(. q 7 R @ @@bB q 7 R @ a@@bB4 q 7 R @ !@@bB6 q 7 R @ @@bBh
p q 7 R @ @@bB;a q 7 R v @ @@bB
R q 7 R g @ @@bBC q 7 R X @ A@@bB94 q 7 R I @ A !A# q 7 R < @ @ @bB q 7 R - @ @ @bB! q 7 R @ a@ @bB q 7 R @ a ! R q 7 hR , @ 7 @@ 7 @ @ @ @ 7 @ @ @@ 7 @ @ @ _{B _ {C @ h 7 @ @[ {A _ { * q 7 =R A # q 7 H>R 7 @ q 7 >R / ҁқ @ 7 h?R # @ * q 7 @ @R @ @@R q 7 @ @R @ @( _{B _ { @ h 7 F @ q跟 7 @
)}
; @ @@ 7 @ @1 @ @ @ 7 @ @% @ @
@ 7 @
@9 @
@@ H 7 @ @ @ {B _ { @!@ 7 = @1@ q 7 @ @))@ k7 7 @A yi q 7 (R ! @ @ %@ q 7 R @!@ q跟 7 ȁR @ ! _{B _ { C _ h 7 + @ q跟 7 _
)}
i@C_ k 7 _
)}
@_)}
* @ {B _{ C ^ q 7 ȃR o ^ _ (! 7 hR b _ !_!@^ C q 7 R T __%@ ** C q 7 聀R F . C q 7 R = _1@ q 7 HR 4 _1@ qH 7 @_))@ k7 7 C *y @ _駀yi _A駀
Q C q 7 HR @ ^_(i _@{U_{ c @)@ q 7 }R F c *@ @@I R+ i Rc @)@ @)A' @)@? @)iG @mK @ q 7 ȂR " @ @)@c @)E' @)@? @)iG @m@I R
*}@K @ q 7 ȂR _{K_ {C @@ @ H 7 @ @7 @ {A _{ CC C_ C Ғ \ 7 \ * 0 ` ,4 Y 7 Y ( 7 Z 7 [ 7 Z h 7 [^ZbB q跟h 7 * @]7 _C ( R; R( r? @(@ \ q跟h 7 v ZZ#bB9 x q跟h 7 h R v *s@ 臟 7 / [ R f @ *@*~*a!2$ R` / /@ q跟h 7 A ]\(I ` U Y \Y( Y Y
` 9E @ @ѨY [Y( Y Y
Z^: q跟h 7 Z 5 q跟h 7
\\( [\( Z\( \ \U # ! _@{Q_{ ѠC āҩ ^ 7 x ^ā * ^ a ^@ Rr q 7 ^ e C__ & R c bB8RR' r ^ ! D R*bB. ] q跟 7 ^ E ^ @ ^ ^ @@ C_ @(M _ @(Q @R(E @R(I ^ @ q跟 7 ^ A ^ ! ^ ^ @R*bBhc z ^ ^@ 7 ^ A ^ ! ^ ^ _ @{¨_ {C @ h 7 ( @L @ ! @
@ h 7 @ q跟 7 @
)}
@ @
@ @ {A _ { @!@ 7 @ %@ q 7 R @!@ q跟 7 ȁR _{B _{C _ 7 _@ 7 c @ 7 pR Y _@@ @@ q跟( 7 c ]- 1 7 HwR C ]iR)r k 7 (wR 9 6 _^ 6 q 7 c ] ( c @@ @ E@ @@@%@ 3@ @^( @ q 7 1 @ q 7 _{E_{ @ q跟h 7 _
)}
i@@ k 7 _
)}
@( 3 _
)}
i@ q 7 ( R @ @ q 7 hR _
)}
@( @ @@ * q 7 @ _{C_C{ _ 7 O _^C q 7 @ C @ 7 @ 7 pR 3 _ @@ *F q跟 7 # pR $ @@% q 7 # @ @@(E _@ @@ R) q 7 # @ _{DC_{ c @)@ q 7 }R F c * @@i R+ I Rc @)@# @)@; @)AC @)iG @mK @ q 7 ȂR " @ @)@c # @)@; @)EC @)iG @m@I R
*}@K w @ q 7 ȂR _{K_{ ^# S q 7 ȃR / ^ A _ (! 7 hR " _ !_!@^5 C q 7 R __%@ **+ C q 7 聀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< 5 R q
D ( o K U 8 p+
/ o % 0 8 ,
@ o x" o (% o ( d : P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P P .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 .data l4t2_handler_arm64.so.debuginfo "T 8 8 @ ~ o x" x" X o (% (% o % % B p+ p+ / t D D ( - B U U 2 hj hj t J P P U U n n TA 7 , , @ @ 2 P P
0 0
8 8 g @ @ X 8 0 $ = D $