
    ^j~^                        d dl mZmZmZ d dlmZmZmZmZ d dl	m
Z
 d dlmZ d dlmZmZmZ d dlmZ d dlZd dlZd dlZd dlZd dlZd dlZd dlmZ d	efd
Zd Zd Zd Z ed       G d d             Z G d de      Z y)    )BaseBackend	GPUTargetLanguage)irpassesllvmamd)knobs)	dataclass)AnyDictTuple)
ModuleTypeN)Pathtargetc                     d S )Nc                      y)N)   r   r    )lhs_typerhs_types     g/var/www/ramen.bs-engineer-server.com/venv/lib/python3.12/site-packages/triton/backends/amd/compiler.py<lambda>z"get_min_dot_size.<locals>.<lambda>   s        r   r   s    r   get_min_dot_sizer      s
     0/r   c                     t         j                  j                  | dk(  xs | dk(  xr |du S t         j                  j                  S )Ngfx942gfx950T)r
   r	   use_block_pingpong)archuse_async_copys     r   is_pingpong_schedule_enabledr#      sI    99''/ HM!1!Ln6L R5:YY5Q5QRr   c                 t    t         j                  j                  | dk(  S t         j                  j                  S )Nr   )r
   r	   use_in_thread_transposer!   s    r   is_in_thread_transpose_enabledr'      s.    !&!B!B!JDHqPUPYPYPqPqqr   c                 6    t         j                  j                  S N)r
   r	   r"   r&   s    r   is_async_copy_enabledr*      s    99###r   T)frozenc                   R   e Zd ZU dZeed<   dZeed<   dZeed<   dZeed<   d	Z	e
ed
<   dZeed<   dZeed<   d	Zeed<   dZee   ed<   dZee   ed<   dZeed<   dZee   ed<   dZeed<   dZeed<   dZeed<   dZeed<   dZeed<   dZeed<   dZeed<   d Zeed!<   d"Zeed#<   d$ Zd% Zy	)&
HIPOptions   	num_warpsr   waves_per_eu   
num_stagesr   num_ctasNextern_libsFdebugTsanitize_overflowr!   )fp8e4nvfp8e5fp8e5b16fp8e4b8supported_fp8_dtypesr   !deprecated_fp8_dot_operand_dtypesieeedefault_dot_input_precision)r=   bf16x3bf16x6allowed_dot_input_precisionsenable_fp_fusionlaunch_cooperative_gridmatrix_instr_nonkdimkpackallow_flush_denormmax_num_imprecise_acc_defaulthipbackend_name instrumentation_modenoneschedule_hintc                    t        | j                  dd       }|dk\  rdnd}t        j                  | d|       | j                  dkD  r| j                  | j                  dz
  z  dk(  sJ d	       | j                  d
k(  rI| j
                  dk7  r:t        j                  d| j
                   d       t        j                  | dd       t        t              j                  dz  }| j                  i nt        | j                        }dD ]  }t        || dz        ||<    t        j                  | dt        |j                                      y )N   
       @   	warp_sizer   r   znum_warps must be a power of 2r   zckpack is deprecated starting from gfx950 and will be removed in later releases. So for now kpack = z7 will be overwritten to 1 to make transitioning easier.rE   lib)ocmlocklz.bcr4   )intr!   object__setattr__r/   rE   warningswarnr   __file__parentr4   dictstrtupleitems)self	gfx_majorrT   default_libdirr4   rU   s         r   __post_init__zHIPOptions.__post_init__S   sG   		!B(	#r/Br	4i8~~!t~~!9K'LQR&R 	-,	-R II!

aMMuvz  wA  wA  vB  By  z tWa0h..6 ,,4b$t?O?O:P# 	AC">se3K#?@K	A4k6G6G6I0JKr   c           	          dj                  | j                  j                         D cg c]  \  }}| d|  c}}      }t        j                  |j                  d            j                         S c c}}w )N_-utf-8)join__dict__rb   hashlibsha256encode	hexdigest)rc   namevalkeys       r   hashzHIPOptions.hashf   s]    hh9L9L9NOID#4&#OP~~cjj12<<>> Ps   A4
) __name__
__module____qualname__r/   rX   __annotations__r0   r2   r3   r4   r_   r5   boolr6   r!   r`   r;   r   r<   r>   rA   rB   rC   rD   rE   rF   rG   rI   rK   rM   rf   rt   r   r   r   r-   r-   #   s    IsL#JHcKE4"t"D#
 (S%*R46%uSz6'--/K %*K!d!$)T) !#!E3N$$)*!3*L# "#"*  M3L&?r   r-   c                   J    e Zd ZdZdZedefd       Zdeddf fdZde	fdZ
defdZd	 Zd
 Zdee	ef   fdZd Zed        Zed        Zed        Zed        Zed        Zed        Zed        Zed        Zed        Zd Z ej<                         d        Z xZ S )
HIPBackendNFr   c                      | j                   dk(  S )NrH   )backendr   s    r   supports_targetzHIPBackend.supports_targeto   s    ~~&&r   returnc                 j    t         |   |       t        |j                  t              sJ d| _        y )Nhsaco)super__init__
isinstancer!   r`   
binary_ext)rc   r   	__class__s     r   r   zHIPBackend.__init__s   s+     &++s+++!r   c                      d|j                    S )Nhip:r&   rc   optionss     r   get_target_namezHIPBackend.get_target_namex   s    gll^$$r   c                    dt         j                  j                  xs | j                  j                  i}|j                  dd      dkD  rKt        j                  | j                  j                        s"t        d| j                  j                         | j                  j                  dk(  rBt        t        j                        }|j                  dh       t        t        |            |d<   d|vr%t        t        t        j                              |d<   | j                  j                  d	k(  rCt        t        j                         }|j                  d
dh       t        t        |            |d<   d|vrt         j"                  j$                  |d<   |j                  t        j&                  j)                         D ci c]  }||v s||   |||    c}       t        di |S c c}w )Nr!   r3   r   znum_ctas > 1 not supported on r   tf32rA   r;   r   r9   r:   r<   rB   r   )r
   runtimeoverride_archr   r!   getr	   supports_multi_cta_launch
ValueErrorsetr-   rA   updatera   sortedr;   r<   languagedefault_fp_fusion__dataclass_fields__keys)rc   optsargsrA   r<   ks         r   parse_optionszHIPBackend.parse_options{   s   33Gt{{7G7GH88J"Q&s/L/LT[[M]M]/^=dkk>N>N=OPQQ ;;x'+.z/V/V+W((//938@\9]3^D/0!-+0
8W8W1X+YD'(;;x'03J4`4`0a--44j)5LM8=fEf>g8hD45T)',~~'G'GD#$)H)H)M)M)OuASTX\S\aefgahatQQZuv!D!! vs   	G+G+G+c                 H    |j                   |j                  |j                  fS r)   )r/   r3   shared)rc   metadatas     r   pack_metadatazHIPBackend.pack_metadata   s%    OO
 	
r   c                 0    dt        | j                        iS )Nmin_dot_size)r   r   r   s     r   get_codegen_implementationz%HIPBackend.get_codegen_implementation   s     0 =>>r   c                     ddl m} d|iS )Nr   )	libdeviceztriton.language.extra.libdevice)triton.language.extra.hipr   )rc   r   s     r   get_module_mapzHIPBackend.get_module_map   s    719==r   c                     t        j                  |       t        j                  r t        j                  j                  |       y y r)   )r	   load_dialectsr{   instrumentation)rc   ctxs     r   r   zHIPBackend.load_dialects   s2    #%%&&44S9 &r   c                     dd l }d}t        | d      r| j                         |k  S t        | |j                        r-t        | d      r!| j                         j                         |k  S y)Nr   i	ptr_rangeuntyped_storageF)torchhasattrr   r   Tensorr   size)argr   
MAX_INT_32s      r   is_within_2gbzHIPBackend.is_within_2gb   s]    
3$==?j00c5<<(WS:K-L&&(--/:==r   c                 H    t        j                  |       }d| v r|ddggz  }|S )NSztt.pointer_rangerR   )r   
parse_attr)descrets     r   r   zHIPBackend.parse_attr   s1    $$T*$;',--C
r   c                     t        j                  | fi |}t        j                  j                  rt
        j                  |       r|dz  }|S )Nr   )r   get_tensor_specializationr
   r	   use_buffer_opsr{   r   )r   kwargsr   s      r   r   z$HIPBackend.get_tensor_specialization   s?    33CB6B99##
(@(@(E3JC
r   c                 2   t        j                  | j                        }|j                          t        j
                  j                  |       t        j                  j                  |       t        j                  |j                        st        j                  j                  |       t        j
                  j                  |       t        j                  j                  |       t        j                  j                  |       t        j
                  j!                  |       t        j                  j#                  |       t        j
                  j%                  |       t        j                  j'                  |       |j)                  | d       | S )N	make_ttir)r   pass_managercontextenable_debugr   commonadd_inlinerttiradd_rewrite_tensor_pointerr	   supports_tdmr!   (add_rewrite_tensor_descriptor_to_pointeradd_canonicalizeradd_combineadd_reorder_broadcastadd_cseadd_triton_licmadd_symbol_dceadd_loop_unrollrun)modr   r   pms       r   r   zHIPBackend.make_ttir   s    __S[[)
!!"%..r2-KK@@D''+#))"-b!##B'$$R(##B'
sK 
r   c                    t        j                  | j                        }|j                          t        j
                  j                  |d|j                   |j                  |j                  |j                         |j                  | d       t        j                  | j                        }|j                          d}t        j                  j                  |       t        j                  j                  ||       t        j                  j                  |       t        j                  j!                  |       t"        j                  j                  j%                  ||j                  |j&                  |j(                         t        j                  j                  |       t"        j                  j                  j+                  |       t"        j                  j                  j-                  ||j                         t"        j                  j                  j/                  |       t"        j                  j                  j1                  |       t        j                  j3                  |       t        j4                  j7                  |       t        j
                  j9                  |       t        j4                  j7                  |       t;        |j                        }t=        |j                  |      }t"        j                  j                  j?                  ||j@                         t"        j                  j                  jC                  |||       |r4t"        j                  j                  jE                  ||j                         t"        j                  j                  jG                  |       t        j4                  j7                  |       |jH                  jK                         dk7  rJ|jH                  jM                  d      D ],  }t"        j                  j                  jO                  ||       . t        j                  j                  |       t        j                  jQ                  |       tS        |j                        rHt"        j                  j                  jU                  |       t        j                  j                  |       t"        j                  j                  jW                  |       |rC|j@                  dkD  r4t"        j                  j                  jY                  ||j@                         tZ        j"                  j\                  rt"        j                  j                  j_                  |       t        j4                  j7                  |       t"        j                  j                  ja                  ||j                  tZ        j"                  jb                  tZ        j"                  jd                         t"        j                  j                  jg                  |       t"        j                  j                  ji                  |       t        j4                  j7                  |       t        j4                  jk                  |       t        j4                  jm                  |       |j                  | d       | jo                         |d<   | S )	Nr   make_ttgir_earlyFrL   ,r   
make_ttgirtensordesc_meta)8r   r   r   r   r   r   add_convert_to_ttgpuirr!   r/   rT   r3   r   ttgpuiradd_coalesceadd_f32_dot_tcadd_remove_layout_conversionsadd_optimize_thread_localityr	   add_accelerate_matmulrD   rE   add_optimize_epilogueadd_optimize_dot_operandsadd_hoist_layout_conversionsadd_sink_layout_conversionsadd_fuse_nested_loopsr   r   r   r*   r#   add_schedule_loopsr2   add_pipelineadd_coalesce_async_copyadd_convert_to_tensor_opsrM   lowersplitinsert_instruction_sched_hintsadd_reduce_data_duplicationr'   add_in_thread_transposeadd_move_up_prologue_loadsadd_block_pingpongr
   r   add_canonicalize_pointersadd_convert_to_buffer_opsuse_buffer_atomics%buffer_ops_analyze_small_tensor_rangeadd_fold_true_cmpiadd_prepare_if_combiningr   r   get_tensordesc_metadata)r   r   r   r   emuTF32r"   r    hints           r   r   zHIPBackend.make_ttgir   s`   __S[[)
**2gll^/DgFWFWY`YjYj+2+;+;	=
s&'__S[[)
##B'%%b'244R833B7

00W\\7C_C_ahanano44R8

004

44RF

77;

66r:,,R0''+##B'''+.w||<9',,W

--b'2D2DE

''N<NOJJ66r7<<H

44R8''+  &&(F2--33C8 L

""AA"dKL44R82226)',,7JJ66r:NN88<

55b9'"4"4q"8JJ11"g6H6HI99##JJ88<MM++B/JJ88		,,		??	 	

--b1

33B7''+b!$$R(
sL!&)&A&A&C"#
r   c                    | }t        j                  |j                        }|j                          t        j
                  j                  |       t        j
                  j                  |       t        j                  j                  |       t        j                  j                  |       t        j
                  j                  |       t        j                  j                  |       t        j                  j                  j!                  |       t        j                  j#                  |       |j%                  |d       |j'                         |d<   |S )Ngluon_to_ttgirr   )r   r   r   r   r   gluonr   add_resolve_auto_encodingsr   add_sccpr   add_loop_aware_cser   r    add_combine_tensor_select_and_ifr	   add_warp_pipelineadd_allocate_warp_groupsr   r   )srcr   r   r   r   s        r   r   zHIPBackend.gluon_to_ttgir  s    __S[[)
  $//3r"&&r*&&r*77;

,,R0//3
s$%&)&A&A&C"#
r   c                 4   | }t        j                  |j                        }|j                          t        j
                  j                  j                  ||j                         t        j
                  j                  j                  ||j                         t
        j                  j                  |       t
        j                  j                  |       t
        j                  j                  |       t        j
                  j                  j                  |       t         j"                  r+t         j"                  j%                  d||j                         d}t        j
                  j                  j'                  ||j                  |       t        j
                  j                  j)                  ||j                         t
        j*                  j-                  |       t
        j*                  j/                  |       t
        j                  j1                  |       t
        j                  j3                  |       t
        j*                  j-                  |       t
        j*                  j/                  |       t
        j*                  j5                  |       |j6                  j9                         dk7  r?t        j
                  j                  j;                  ||j                  |j<                         t         j"                  r+t         j"                  j%                  d||j                         t>        j@                  jB                  s9t>        j@                  jD                  st
        jF                  jI                  |       t        j
                  j                  jK                  ||j                  |       |jM                  |d       t>        j@                  jD                  rt>        j@                  jB                  s`t        j                  |j                        }|j                          t
        jF                  jI                  |       |jM                  |d       t        j                  |j                        }|j                          t
        jF                  jO                  |       |jM                  |d       tQ        jR                          tQ        j                         }tQ        jT                  ||      }t	        jV                  |       d}t>        j@                  jX                  rd	}tQ        jZ                  |t        j\                  |j                  |       t	        j^                  ||j                         t	        j`                  |d
       t	        jb                  |dd       t	        jb                  |dd       t	        jb                  |dd       t	        jb                  |d|jd                  dk(         |jg                         D 	cg c]  }	|	ji                         r|	 }
}	|
d   jk                  t        jl                         |d   }|
d   jo                  d| d       |jp                  }| js                  d      }||}|
d   jo                  dd||jd                  z          d|j6                  ju                  d      v r|
d   jo                  dd       |
d   jo                  dd       |
d   jo                  d|jv                   d|jv                          |jx                  rd nd!}|
d   jo                  d"|       t>        j@                  jX                  r'|
d   j{                  d	       |
d   j}                          |j                  d#k7  rLt>        j                  j~                  s2t>        j                  j                  st	        j                  |
d          t>        j@                  jX                  r\t        t              j                  d$z  }t        |d%z        t        |d&z        t        |d'z        g}tQ        j                  ||       ne|j                  rY|j                  D cg c]  \  }}t	        j                  ||      s|  }}}t        |      dkD  rtQ        j                  ||       tQ        j                  |tP        j                  |j                  dg |j                         t	        j                  |j                        r<|
d   j                  d(       |
d   j                  d)       |
d   j                  d*       t>        j                  j                  rt	        j                  |
d          ||d+<   | js                  d,      |d-<   | js                  d.      xs d|d/<   | js                  d0      xs d1|d2<   t	        j                  |       t	        j                  |       t        |      S c c}	w c c}}w )3Nttgpuir_to_llvmirTrL   llvmir_to_llvm	make_llirzmake_llir.disable_line_infoz,make_llir.dump_ir_extract_di_local_variablesrJ   +xnacki  __oclc_finite_only_optF__oclc_correctly_rounded_sqrt32__oclc_unsafe_math_opt__oclc_wavefrontsize64rS   r   r3   zamdgpu-cluster-dimsz,1,1zttg.total-num-warpszamdgpu-flat-work-group-sizez1,zmemory-bound-attentionr   zamdgpu-sched-strategyziterative-ilpzuniform-work-group-sizetruezamdgpu-waves-per-euz, zpreserve-signr=   zdenormal-fp-math-f32gfx1250rU   z
asanrtl.bczocml.bczockl.bczamdgpu-no-workgroup-id-xzamdgpu-no-workgroup-id-yzamdgpu-no-workgroup-id-zr/   z
ttg.sharedr   zttg.profile_scratch_memory_sizeprofile_scratch_sizez$ttg.profile_scratch_memory_alignmentr   profile_scratch_align)Sr   r   r   r   r	   r   r   add_update_async_wait_countr!   add_warp_pipeline_conversionconvertadd_scf_to_cfr   r   add_index_to_llvmiradd_allocate_shared_memoryr{   r   patchadd_to_llvmiradd_warp_specialize_to_llvmr   r   r   add_cf_to_llvmiradd_arith_to_llvmirr   rM   r   lower_instruction_sched_hintsr2   r
   compilationdisable_line_info"dump_ir_extract_di_local_variablesllvmiradd_di_scopeadd_builtin_func_to_llvmirr   add_di_local_variabler   init_targets	to_moduleattach_target_tripleenable_asanattach_datalayoutTARGET_TRIPLEset_isa_versionset_abi_versionset_bool_control_constantrT   get_functionsis_declarationset_calling_convCALLING_CONV_AMDGPU_KERNELadd_fn_attrr/   get_int_attrr   r0   rF   add_fn_target_featureadd_fn_asan_attrswap_mirdump_mirset_all_fn_arg_inregr   r]   r^   r`   link_extern_libsr4   need_extern_liblenoptimize_moduleOPTIMIZE_O3rB   has_architected_sgprsremove_fn_attrscalarize_packed_fops#add_scalarize_packed_fops_llvm_passcleanup_bitcode_metadatadisable_print_inline)r  r   r   r   r   _HIPBackend__HIP_FTZr   llvm_modtarget_featuresfnfnscluster_dimtotal_warps_numtotal_num_warpsdenormal_modere   pathsrq   paths                      r   r  zHIPBackend.make_llir(  so   __S[[)


66r7<<H

77GLLI$$R(  $**2.

55b9%%&&,,-@"ckkR 	

((W\\9E

66r7<<H''+b!''+**2.''+b!$$R(  &&(F2JJ<<RwOaOab %%&&,,-=r3;;O  225;L;L;o;oMM&&r*

55b',,	R
sK ??$$66__S[[1!**2.s9: -BOOMM//3FF3FG 	,,.>>#w/  *((&Ox):):GLL/Z 	Hgll3Hc*%%h0H%P%%h0QSWX%%h0H%P%%h0H'J[J[_aJab %224PbB<M<M<OrPPA > >?z*A0[M2FG!++**+@A&-OA8BwO`O`?`>a:bc#w'<'<'B'B3'GGF6HA4f= 	A0W5I5I4J"WMaMaLb2cd+2+E+E6A1=A((F((2F##% <<9$eii.@.@EIIDVDV$$SV,((!(^22U:NN\12NY./NY./E
 !!(E2  .5.A.AiltTSEXEXYacgEhTiEi5zA~%%h6Xt'7'7r2wOgOgh $$W\\2F!!"<=F!!"<=F!!"<=99**33CF; !0 --l;+.+;+;<]+^+cbc'(,/,<,<=c,d,ihi()$$X. 	  *8}W Q^ js   h,h>hhc           	         t        j                  d|       }t        |      dk(  sJ |d   |d<   g }d|j                  v rdnd}t	        j
                  | j                  d            j                         }|d   d	z   |z   }t        j                  | t        j                  |j                  |||j                  |      }t        j                  | t        j                  |j                  |||j                  |       t        j                  j                  rwt        j                   t"        j$                  j'                  t        j                  j                  |d
z         t        j                  |j                  |||j                  d      }	n=t        j(                  | t        j                  |j                  |||j                  d      }	t        j                  j*                  rt-        d       t-        |	       |	S )Nz3define amdgpu_kernel void @([a-zA-Z_][a-zA-Z0-9_]*)r   r   rq   gfx11z-real-true16rJ   rj   rh   z.txtFz!// -----// AMDGCN Dump //----- //)refindallr:  r!   rm   rn   ro   rp   r   translate_to_mirr	   r)  rB   dump_sched_dagr
   r5  translate_mir_to_asmosrM  rk   translate_to_asmdump_amdgcnprint)
r  r   r   namesflagsfeaturesir_hashdump_file_idrh   amdgcns
             r   make_amdgcnzHIPBackend.make_amdgcn  s   
 

QSVW5zQ 8%,%<>"..G!45??AQx#~/!!#s'8'8',,RWY`YqYq".0C!2!2GLL(ESZSkSk(	*99..rww||EII<N<N<H6<Q0STWTeTegngsgs/7@X@XZ_aF **30A0A7<<QY[`+2+C+CULF99  56&Mr   c                    d}t         j                  j                  rd}d|j                  v r|dz  }t	        j
                  | |j                  |      }t        j                         5 }t        j                         5 }t        |j                  d      5 }|j                  |       d d d        t	        j                  |j                  |j                         d d d        t        |j                  d      5 }|j                         }	d d d        d d d        	S # 1 sw Y   txY w# 1 sw Y   NxY w# 1 sw Y   +xY w# 1 sw Y   	S xY w)NrJ   r  rO  z,-real-true16wbrb)r
   r  r'  r!   r	   assemble_amdgcntempfileNamedTemporaryFileopenrq   write
link_hsacoread)
r  r   r   rE  r   tmp_outtmp_infd_infd_outr   s
             r   
make_hsacozHIPBackend.make_hsaco  s   ((&Ogll".O##CG((* 	$g,,. :&&++t, 'KK&'v{{GLL9: gllD) $Vkkm$	$ 
' ': :$ $	$ 
sT   'D<<D$D%2D$D<5D0D<D!D$$D-	)D<0D9	5D<<Ec                 >    |t         j                  k(  r fd|d<    fd|d<   n|t         j                  k(  r	 fd|d<    fd|d<    fd|d	<    fd
|d<   t        j                  j
                  $t        j                  j                   ||d        y y )Nc                 *    j                  | |      S r)   )r   r  r   r   rc   s     r   r   z'HIPBackend.add_stages.<locals>.<lambda>  s    4>>#xQX3Y r   r   c                 *    j                  | |      S r)   )r   rq  s     r   r   z'HIPBackend.add_stages.<locals>.<lambda>  s    DOOCSZ4[ r   ttgirc                 *    j                  | |      S r)   )r   rq  s     r   r   z'HIPBackend.add_stages.<locals>.<lambda>  s    D4G4GXW^4_ r   c                 *    j                  | |      S r)   )r  rq  s     r   r   z'HIPBackend.add_stages.<locals>.<lambda>  s    t~~c8W/U r   llirc                 *    j                  | |      S r)   )r_  rq  s     r   r   z'HIPBackend.add_stages.<locals>.<lambda>   s    1A1A#xQX1Y r   r^  c                 *    j                  | |      S r)   )rn  rq  s     r   r   z'HIPBackend.add_stages.<locals>.<lambda>  s    Xw0W r   r   )r   TRITONGLUONr
   r   add_stages_inspection_hook)rc   stagesr   r   s   ` ` r   
add_stageszHIPBackend.add_stages  s    x&YF6N[F7O'_F7OUvYxWw==33?MM44T67HVZ[ @r   c                     | j                    S r)   r   )rc   s    r   rt   zHIPBackend.hash  s    ++r   )!ru   rv   rw   r   %supports_native_tensor_specializationstaticmethodr   r~   r   r`   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r   r  r_  rn  r}  	functools	lru_cachert   __classcell__)r   s   @r   r{   r{   k   sS   O,1)'	 ' '"y "T "
%# %"S "2
?>S*_ 5 >
:
        " ? ?B  $ ` `D  8   
\ Y   r   r{   )!triton.backends.compilerr   r   r   triton._C.libtritonr   r   r   r	   tritonr
   dataclassesr   typingr   r   r   typesr   rU  rm   rd  rP  r  r[   pathlibr   r   r#   r'   r*   r-   r{   r   r   r   <module>r     s    E E 5 5  ! # #  	   	   0Y 0R
r$ $D? D? D?N\  \ r   