ELF(4a4 (444hLhLVff  |X|h|h0044488QtdpRtdVff AndroidGNUA~p K&5E4}/<Am   . ? <B c!(+0,8,?,E,J,Q,g,z,,<====GJJJ%J9JzSSSTTTT UUUUUV3VMVcVXXXYJYPYYYYڣM5GܦM (M, l )H @\ e5 #CD ei  p s , 5 T ZT J}2  )/ X/  Vي m  - rf,A/6 }aH +q8  ԓH Ya  c- gQ [o  !C\ }} uL d g% y ?|e Da Y] QL = ` t3e$ a *5T \٧ o q TL b}L ԄŲ\ ,Ð %G qE Q a FZ u1 R+ [ !vɰ o FeT < U@ T(M, rt ʔ f L bž hx jn }  3sE p6h #z} %AL  8 7)! -7  fi gH ~5 ^ Tz V A%L /u '= =8 F. B uݰ s5 ^]T ,+X a yU; AC  s~ % #p4 L q93`  sL U+g,iH RvH  pe8 e: eH ŕE & Z ߛI i  *uM@ B | /h JS 8d q ~ # $ Q}9" ~8 >.A. ( L @3 z} و Z L ] FT u\ nkU 1 5 ̜)( b@i [ 4 \ /E[ Ep FjL V4 '!L <qL QT ey fT >Y< I Pf,N  [%" |^ k9 {uT ͺL T -[[ Gt ] T d] T G~ YH EU  +PH  A) |!T , ` J8, _ ! og,I{u G[ 1d$ n  )M }jY i)#< P [|f,9L MI d ] 3a0 95_ \ e#r rI   4"-D s  ,՗b ' vL @hd [ jM   Ym " &HP 7x Z q- eT qgL z\ -/  "lq (! !L yH fU =kH Bra MH )#< nL r] t  A<T h H Q  z{ Qu8 ; M} >_˪6 y (QW  Q= Jy 0 LW = + QT D]j  ӀL  p7,  Y^ li Ʌ 2   Z Ɏ S* v kH > L oT 2e " ?AH a)H k98 >z F^ YP" { 6QD f,WE'  &M )D xT U AT LA$  L T 05 6 | H (  O' | _ 910D u7 ۋ / )BL &1Jt qյ j\ iL L0 Z K9$ e x R!MFH ` ]h % iB4 bqyH z 1 Ag= m60 5-a IpVX1" ^] Jp y :If0 M mT ^o 6q  m gc U<f jH w $Lp ym T ×i <P TfH K ?] ?+ H5" .y Rm Y (Z 1%G ^Um cML  p)P' Bi g9 t ΍ '   PӐ ]t ['HP (-Un uH BD mW `y `H h Q g [- ӡT /!  ^LO ) 3e$ vH P&L `9yf$ "CD _]T 2h \ e !Q ԌH < X?p8 V1e$ % 2ae$ A~L  hMT \ d >$L t %I չ N& t:f0 ud- n X @J+ 9X ] @ xT  `q Ha  ,IH y we YY { (p[\ t 5 U?. W%< c a tU R `K!  z)T {p *o f}L Y\ f[ \ v\ kA  $ A mu ( )^ #2e p5e : ! bIH c"" + p 3 # ^( X B YT wQU :i ֧%pqJA c W 5t $t 7ceL Ke1H zPG 2B aYI  qѱH M1  Njm gZ l}  C ] e a)H 'EM ke: LI ^ (g,z_ip l]T U& hپL (I8 D EF &M  lT YL Gy`t g*uM@ p @f{bɽ ā u ! 1=e$ ru u ^>p I}o" W / E H \9 k  h v jɶL SM" pʧ$pL=g8" x ~qL Z k|Q  R= raq n E w 1) sL 7 J x ; ;R ZD> 'f! #7d \ f_q <f" ` d __cxa_finalizeLIBClibc.solibRSDriverArm.so__cxa_atexit__register_atfork_Z19isAllocationCpuOnlyPKN7android12renderscript10AllocationE_Z19rsdGpuIntrinsic_LUTPKN7android12renderscript7ContextEPNS0_6ScriptEPKNS0_7ElementE19RsScriptIntrinsicID_Z20rsdGpuIntrinsic_BlurPKN7android12renderscript7ContextEPNS0_6ScriptEPKNS0_7ElementE19RsScriptIntrinsicID_Z21rsdGpuIntrinsic_3DLUTPKN7android12renderscript7ContextEPNS0_6ScriptEPKNS0_7ElementE19RsScriptIntrinsicID_Z25rsdGpuIntrinsic_HistogramPKN7android12renderscript7ContextEPNS0_6ScriptEPKNS0_7ElementE19RsScriptIntrinsicIDdladdrlibdl.so_Z27rsdGpuIntrinsic_ColorMatrixPKN7android12renderscript7ContextEPNS0_6ScriptEPKNS0_7ElementE19RsScriptIntrinsicIDdl_iterate_phdr_Z27rsdGpuIntrinsic_Convolve3x3PKN7android12renderscript7ContextEPNS0_6ScriptEPKNS0_7ElementE19RsScriptIntrinsicID_Z29rsdClAllocationReleaseHostPtrPKN7android12renderscript7ContextEPKNS0_10AllocationE_ZN7android12renderscript21rsdGpuScriptIntrinsic11getNumSlotsEv_ZN7android12renderscript21rsdGpuScriptIntrinsic12launchKernelEP10_cl_kernel_ZN7android12renderscript21rsdGpuScriptIntrinsic12setGlobalObjEjPNS0_10ObjectBaseE_ZN7android12renderscript21rsdGpuScriptIntrinsic12setGlobalVarEjPKvj_ZN7android12renderscript21rsdGpuScriptIntrinsic13prepareLaunchEPK12RsScriptCallPPKNS0_10AllocationEj_ZN7android12renderscript21rsdGpuScriptIntrinsic13scaleWorkSizeEi_ZN7android12renderscript21rsdGpuScriptIntrinsic13setGlobalBindEjPNS0_10AllocationE_ZN7android12renderscript21rsdGpuScriptIntrinsic15createIntrinsicEPKNS0_7ContextEPNS0_6ScriptEPKNS0_7ElementE19RsScriptIntrinsicID_ZN7android12renderscript21rsdGpuScriptIntrinsic24setGlobalVarWithElemDimsEjPKvjPKNS0_7ElementEPKjj_ZN7android12renderscript21rsdGpuScriptIntrinsic7gpuRootEPKNS0_7ContextEPNS0_6ScriptEjPKNS0_10AllocationEPS7_PKvjPK12RsScriptCall_ZN7android12renderscript21rsdGpuScriptIntrinsic7initGpuEP9RsdHalRecPNS0_6ScriptE_ZN7android12renderscript21rsdGpuScriptIntrinsic8getClMemEPKNS0_10AllocationE_ZN7android12renderscript21rsdGpuScriptIntrinsicC1EPKNS0_7ContextEPNS0_6ScriptEPKNS0_7ElementE19RsScriptIntrinsicID_ZN7android12renderscript21rsdGpuScriptIntrinsicC2EPKNS0_7ContextEPNS0_6ScriptEPKNS0_7ElementE19RsScriptIntrinsicID_ZN7android12renderscript27rsdGpuScriptIntrinsicBuffer3mapEv_ZN7android12renderscript27rsdGpuScriptIntrinsicBuffer5unmapEv_ZN7android12renderscript27rsdGpuScriptIntrinsicBufferC1EP8RsdCLRecjyPv_ZN7android12renderscript27rsdGpuScriptIntrinsicBufferC2EP8RsdCLRecjyPv_ZN7android12renderscript27rsdGpuScriptIntrinsicBufferD0Ev_ZN7android12renderscript27rsdGpuScriptIntrinsicBufferD1Ev_ZN7android12renderscript27rsdGpuScriptIntrinsicBufferD2Ev_ZNK7android12renderscript7Context8setErrorE7RsErrorPKc_ZTVN7android12renderscript21rsdGpuScriptIntrinsicE_ZTVN7android12renderscript27rsdGpuScriptIntrinsicBufferE_ZdlPv__android_log_print__stack_chk_fail__stack_chk_guard_Z26rsdClSetAllocationArgumentPKN7android12renderscript7ContextEPKNS0_6ScriptEP13maliKernelRecjPKNS0_10AllocationE_ZN7android12renderscript26rsdGpuScriptIntrinsic3DLUT11getNumSlotsEv_ZN7android12renderscript26rsdGpuScriptIntrinsic3DLUT12gpuRunHelperEP13maliKernelRecPKNS0_10AllocationEPS4__ZN7android12renderscript26rsdGpuScriptIntrinsic3DLUT12setGlobalObjEjPNS0_10ObjectBaseE_ZN7android12renderscript26rsdGpuScriptIntrinsic3DLUT7gpuRootEPKNS0_7ContextEPNS0_6ScriptEjPKNS0_10AllocationEPS7_PKvjPK12RsScriptCall_ZN7android12renderscript26rsdGpuScriptIntrinsic3DLUT7initGpuEP9RsdHalRecPNS0_6ScriptE_ZN7android12renderscript26rsdGpuScriptIntrinsic3DLUTC1EPKNS0_7ContextEPNS0_6ScriptEPKNS0_7ElementE19RsScriptIntrinsicID_ZN7android12renderscript26rsdGpuScriptIntrinsic3DLUTC2EPKNS0_7ContextEPNS0_6ScriptEPKNS0_7ElementE19RsScriptIntrinsicID_ZNK7android12renderscript10ObjectBase9decSysRefEv_ZNK7android12renderscript10ObjectBase9incSysRefEv_ZTVN7android12renderscript26rsdGpuScriptIntrinsic3DLUTE_Znwjcallocsrc_kernels_lut3d_Z21rsdGpuIntrinsic_BlendPKN7android12renderscript7ContextEPNS0_6ScriptEPKNS0_7ElementE19RsScriptIntrinsicID_ZN7android12renderscript26rsdGpuScriptIntrinsicBlend11getNumSlotsEv__vsnprintf_chk_ZN7android12renderscript26rsdGpuScriptIntrinsicBlend12gpuRunHelperEiPKNS0_10AllocationEPS2_PK12RsScriptCalli_ZN7android12renderscript26rsdGpuScriptIntrinsicBlend7gpuRootEPKNS0_7ContextEPNS0_6ScriptEjPKNS0_10AllocationEPS7_PKvjPK12RsScriptCall_ZN7android12renderscript26rsdGpuScriptIntrinsicBlend7initGpuEP9RsdHalRecPNS0_6ScriptE_ZN7android12renderscript26rsdGpuScriptIntrinsicBlendC1EPKNS0_7ContextEPNS0_6ScriptEPKNS0_7ElementE19RsScriptIntrinsicID_ZN7android12renderscript26rsdGpuScriptIntrinsicBlendC2EPKNS0_7ContextEPNS0_6ScriptEPKNS0_7ElementE19RsScriptIntrinsicID_ZTVN7android12renderscript26rsdGpuScriptIntrinsicBlendEsrc_kernels_blend_ZN7android12renderscript25rsdGpuScriptIntrinsicBlur11getNumSlotsEv_ZN7android12renderscript25rsdGpuScriptIntrinsicBlur11gpuFirstRunEP13maliKernelRecPKNS0_10AllocationE_ZN7android12renderscript25rsdGpuScriptIntrinsicBlur12gpuSecondRunEP13maliKernelRecPKNS0_10AllocationE_ZN7android12renderscript25rsdGpuScriptIntrinsicBlur12setGlobalObjEjPNS0_10ObjectBaseE_ZN7android12renderscript25rsdGpuScriptIntrinsicBlur12setGlobalVarEjPKvj_ZN7android12renderscript25rsdGpuScriptIntrinsicBlur22ComputeGaussianWeightsEv_ZN7android12renderscript25rsdGpuScriptIntrinsicBlur7gpuRootEPKNS0_7ContextEPNS0_6ScriptEjPKNS0_10AllocationEPS7_PKvjPK12RsScriptCall_ZN7android12renderscript25rsdGpuScriptIntrinsicBlur7initGpuEP9RsdHalRecPNS0_6ScriptE_ZN7android12renderscript25rsdGpuScriptIntrinsicBlurC1EPKNS0_7ContextEPNS0_6ScriptEPKNS0_7ElementE19RsScriptIntrinsicID_ZN7android12renderscript25rsdGpuScriptIntrinsicBlurC2EPKNS0_7ContextEPNS0_6ScriptEPKNS0_7ElementE19RsScriptIntrinsicID_ZTVN7android12renderscript25rsdGpuScriptIntrinsicBlurE__aeabi_memclr8LIBC_N__aeabi_memcpypowflibm.sosrc_kernels_blur_ZN7android12renderscript32rsdGpuScriptIntrinsicColorMatrix11getNumSlotsEv_ZN7android12renderscript32rsdGpuScriptIntrinsicColorMatrix12gpuRunHelperEP13maliKernelRecPKNS0_10AllocationEPS4_i_ZN7android12renderscript32rsdGpuScriptIntrinsicColorMatrix12setGlobalVarEjPKvj_ZN7android12renderscript32rsdGpuScriptIntrinsicColorMatrix14float2half_rtzEf_ZN7android12renderscript32rsdGpuScriptIntrinsicColorMatrix18gpuRunHelperSimpleEP13maliKernelRecPKNS0_10AllocationEPS4_i_ZN7android12renderscript32rsdGpuScriptIntrinsicColorMatrix22gpuRunHelperGreyU84_U8EP13maliKernelRecPKNS0_10AllocationEPS4_i_ZN7android12renderscript32rsdGpuScriptIntrinsicColorMatrix7gpuRootEPKNS0_7ContextEPNS0_6ScriptEjPKNS0_10AllocationEPS7_PKvjPK12RsScriptCall_ZN7android12renderscript32rsdGpuScriptIntrinsicColorMatrix7initGpuEP9RsdHalRecPNS0_6ScriptE_ZN7android12renderscript32rsdGpuScriptIntrinsicColorMatrix8checkVarEv_ZN7android12renderscript32rsdGpuScriptIntrinsicColorMatrixC1EPKNS0_7ContextEPNS0_6ScriptEPKNS0_7ElementE19RsScriptIntrinsicID_ZN7android12renderscript32rsdGpuScriptIntrinsicColorMatrixC2EPKNS0_7ContextEPNS0_6ScriptEPKNS0_7ElementE19RsScriptIntrinsicID_ZTVN7android12renderscript32rsdGpuScriptIntrinsicColorMatrixEsrc_kernels_colormatrix_ZN7android12renderscript32rsdGpuScriptIntrinsicConvolve3x311findPatternEv_ZN7android12renderscript32rsdGpuScriptIntrinsicConvolve3x311getNumSlotsEv_ZN7android12renderscript32rsdGpuScriptIntrinsicConvolve3x312gpuRunHelperENS1_11c3x3_kernelEPNS0_10AllocationE_ZN7android12renderscript32rsdGpuScriptIntrinsicConvolve3x312setGlobalObjEjPNS0_10ObjectBaseE_ZN7android12renderscript32rsdGpuScriptIntrinsicConvolve3x312setGlobalVarEjPKvj_ZN7android12renderscript32rsdGpuScriptIntrinsicConvolve3x37gpuRootEPKNS0_7ContextEPNS0_6ScriptEjPKNS0_10AllocationEPS7_PKvjPK12RsScriptCall_ZN7android12renderscript32rsdGpuScriptIntrinsicConvolve3x37initGpuEP9RsdHalRecPNS0_6ScriptEfprintf_ZN7android12renderscript32rsdGpuScriptIntrinsicConvolve3x3C1EPKNS0_7ContextEPNS0_6ScriptEPKNS0_7ElementE19RsScriptIntrinsicIDfflush_ZN7android12renderscript32rsdGpuScriptIntrinsicConvolve3x3C2EPKNS0_7ContextEPNS0_6ScriptEPKNS0_7ElementE19RsScriptIntrinsicID_ZTVN7android12renderscript32rsdGpuScriptIntrinsicConvolve3x3Esrc_kernels_convolve3x3_Z27rsdGpuIntrinsic_Convolve5x5PKN7android12renderscript7ContextEPNS0_6ScriptEPKNS0_7ElementE19RsScriptIntrinsicID_ZN7android12renderscript32rsdGpuScriptIntrinsicConvolve5x511getNumSlotsEv_ZN7android12renderscript32rsdGpuScriptIntrinsicConvolve5x512gpuRunHelperEiPNS0_10AllocationEPK12RsScriptCalli_ZN7android12renderscript32rsdGpuScriptIntrinsicConvolve5x512setGlobalObjEjPNS0_10ObjectBaseE_ZN7android12renderscript32rsdGpuScriptIntrinsicConvolve5x512setGlobalVarEjPKvj__sF_ZN7android12renderscript32rsdGpuScriptIntrinsicConvolve5x57gpuRootEPKNS0_7ContextEPNS0_6ScriptEjPKNS0_10AllocationEPS7_PKvjPK12RsScriptCall_ZN7android12renderscript32rsdGpuScriptIntrinsicConvolve5x57initGpuEP9RsdHalRecPNS0_6ScriptE_ZN7android12renderscript32rsdGpuScriptIntrinsicConvolve5x5C1EPKNS0_7ContextEPNS0_6ScriptEPKNS0_7ElementE19RsScriptIntrinsicID_ZN7android12renderscript32rsdGpuScriptIntrinsicConvolve5x5C2EPKNS0_7ContextEPNS0_6ScriptEPKNS0_7ElementE19RsScriptIntrinsicID_ZTVN7android12renderscript32rsdGpuScriptIntrinsicConvolve5x5Esrc_kernels_convolve5x5_Z29rsdClAllocationAcquireHostPtrPKN7android12renderscript7ContextEPKNS0_10AllocationEb_ZN7android12renderscript30rsdGpuScriptIntrinsicHistogram11getNumSlotsEv_ZN7android12renderscript30rsdGpuScriptIntrinsicHistogram12setGlobalObjEjPNS0_10ObjectBaseE_ZN7android12renderscript30rsdGpuScriptIntrinsicHistogram12setGlobalVarEjPKvj_ZN7android12renderscript30rsdGpuScriptIntrinsicHistogram17gpuRunHelper_fullEPKNS0_10AllocationE_ZN7android12renderscript30rsdGpuScriptIntrinsicHistogram22gpuRunHelper_remainderEPKNS0_10AllocationE_ZN7android12renderscript30rsdGpuScriptIntrinsicHistogram7gpuRootEPKNS0_7ContextEPNS0_6ScriptEjPKNS0_10AllocationEPS7_PKvjPK12RsScriptCall_ZN7android12renderscript30rsdGpuScriptIntrinsicHistogram7initGpuEP9RsdHalRecPNS0_6ScriptE_ZN7android12renderscript30rsdGpuScriptIntrinsicHistogramC1EPKNS0_7ContextEPNS0_6ScriptEPKNS0_7ElementE19RsScriptIntrinsicID_ZN7android12renderscript30rsdGpuScriptIntrinsicHistogramC2EPKNS0_7ContextEPNS0_6ScriptEPKNS0_7ElementE19RsScriptIntrinsicID_ZN7android12renderscript30rsdGpuScriptIntrinsicHistogramD0Ev_ZN7android12renderscript30rsdGpuScriptIntrinsicHistogramD1Ev_ZN7android12renderscript30rsdGpuScriptIntrinsicHistogramD2Ev_ZTVN7android12renderscript30rsdGpuScriptIntrinsicHistogramE__aeabi_memclrsrc_kernels_histogram_ZN7android12renderscript24rsdGpuScriptIntrinsicLUT11getNumSlotsEv_ZN7android12renderscript24rsdGpuScriptIntrinsicLUT12setGlobalObjEjPNS0_10ObjectBaseE_ZN7android12renderscript24rsdGpuScriptIntrinsicLUT7gpuRootEPKNS0_7ContextEPNS0_6ScriptEjPKNS0_10AllocationEPS7_PKvjPK12RsScriptCall_ZN7android12renderscript24rsdGpuScriptIntrinsicLUT7initGpuEP9RsdHalRecPNS0_6ScriptE_ZN7android12renderscript24rsdGpuScriptIntrinsicLUTC1EPKNS0_7ContextEPNS0_6ScriptEPKNS0_7ElementE19RsScriptIntrinsicID_ZN7android12renderscript24rsdGpuScriptIntrinsicLUTC2EPKNS0_7ContextEPNS0_6ScriptEPKNS0_7ElementE19RsScriptIntrinsicID_ZTVN7android12renderscript24rsdGpuScriptIntrinsicLUTEsrc_kernels_lut_Z13rsdCLShutdownPKN7android12renderscript7ContextE_Z9rsdCLInitPKN7android12renderscript7ContextE_ZN7android12renderscript12property_getEPKcPcS2_dlerrordlopendlsymfreemallocpthread_mutex_destroypthread_mutex_initstrcmpstrncmp_Z17rsdAllocationInitPKN7android12renderscript7ContextEPNS0_10AllocationEb_Z19rsdAllocationData1DPKN7android12renderscript7ContextEPKNS0_10AllocationEjjjPKvj_Z19rsdAllocationData2DPKN7android12renderscript7ContextEPKNS0_10AllocationEjjj23RsAllocationCubemapFacejjPKvjj_Z19rsdAllocationData3DPKN7android12renderscript7ContextEPKNS0_10AllocationEjjjjjjjPKvjj_Z19rsdAllocationLock1DPKN7android12renderscript7ContextEPKNS0_10AllocationE_Z19rsdAllocationRead1DPKN7android12renderscript7ContextEPKNS0_10AllocationEjjjPvj_Z19rsdAllocationRead2DPKN7android12renderscript7ContextEPKNS0_10AllocationEjjj23RsAllocationCubemapFacejjPvjj_Z19rsdAllocationRead3DPKN7android12renderscript7ContextEPKNS0_10AllocationEjjjjjjjPvjj_Z19rsdAllocationResizePKN7android12renderscript7ContextEPKNS0_10AllocationEPKNS0_4TypeEb_Z19rsdClAllocationInitPKN7android12renderscript7ContextEPNS0_10AllocationEb_Z20rsdAllocationDestroyPKN7android12renderscript7ContextEPNS0_10AllocationE_Z20rsdAllocationSyncAllPKN7android12renderscript7ContextEPKNS0_10AllocationE21RsAllocationUsageType_Z21rsdAllocationUnlock1DPKN7android12renderscript7ContextEPKNS0_10AllocationE_Z21rsdClAllocationData1DPKN7android12renderscript7ContextEPKNS0_10AllocationEjjjPKvj_Z21rsdClAllocationData2DPKN7android12renderscript7ContextEPKNS0_10AllocationEjjj23RsAllocationCubemapFacejjPKvjj_Z21rsdClAllocationData3DPKN7android12renderscript7ContextEPKNS0_10AllocationEjjjjjjjPKvjj_Z21rsdClAllocationLock1DPKN7android12renderscript7ContextEPKNS0_10AllocationE_Z21rsdClAllocationMapAllPKN7android12renderscript7ContextE_Z21rsdClAllocationRead1DPKN7android12renderscript7ContextEPKNS0_10AllocationEjjjPvj_Z21rsdClAllocationRead2DPKN7android12renderscript7ContextEPKNS0_10AllocationEjjj23RsAllocationCubemapFacejjPvjj_Z21rsdClAllocationRead3DPKN7android12renderscript7ContextEPKNS0_10AllocationEjjjjjjjPvjj_Z21rsdClAllocationResizePKN7android12renderscript7ContextEPKNS0_10AllocationEPKNS0_4TypeEb_Z22rsdAllocationMarkDirtyPKN7android12renderscript7ContextEPKNS0_10AllocationE_Z22rsdClAllocationDestroyPKN7android12renderscript7ContextEPNS0_10AllocationE_Z22rsdClAllocationSyncAllPKN7android12renderscript7ContextEPKNS0_10AllocationE21RsAllocationUsageType_Z23rsdAllocationSetSurfacePKN7android12renderscript7ContextEPNS0_10AllocationEP13ANativeWindow_Z23rsdClAllocationShutdownPKN7android12renderscript7ContextE_Z23rsdClAllocationUnlock1DPKN7android12renderscript7ContextEPKNS0_10AllocationE_Z24rsdAllocationAdapterInitPKN7android12renderscript7ContextEPNS0_10AllocationE_Z24rsdAllocationElementDataPKN7android12renderscript7ContextEPKNS0_10AllocationEjjjPKvjj_Z24rsdAllocationElementReadPKN7android12renderscript7ContextEPKNS0_10AllocationEjjjPvjj_Z24rsdClAllocationMarkDirtyPKN7android12renderscript7ContextEPKNS0_10AllocationE_Z25rsdAllocationData1D_allocPKN7android12renderscript7ContextEPKNS0_10AllocationEjjjS6_jj_Z25rsdAllocationData2D_allocPKN7android12renderscript7ContextEPKNS0_10AllocationEjjj23RsAllocationCubemapFacejjS6_jjjS7__Z25rsdAllocationData3D_allocPKN7android12renderscript7ContextEPKNS0_10AllocationEjjjjjjjS6_jjjj_Z25rsdClAllocationSetSurfacePKN7android12renderscript7ContextEPNS0_10AllocationEP13ANativeWindow_Z26rsdClAllocationAdapterInitPKN7android12renderscript7ContextEPNS0_10AllocationE_Z26rsdClAllocationElementDataPKN7android12renderscript7ContextEPKNS0_10AllocationEjjjPKvjj_Z26rsdClAllocationElementReadPKN7android12renderscript7ContextEPKNS0_10AllocationEjjjPvjj_Z27rsdClAllocationData1D_allocPKN7android12renderscript7ContextEPKNS0_10AllocationEjjjS6_jj_Z27rsdClAllocationData2D_allocPKN7android12renderscript7ContextEPKNS0_10AllocationEjjj23RsAllocationCubemapFacejjS6_jjjS7__Z27rsdClAllocationData3D_allocPKN7android12renderscript7ContextEPKNS0_10AllocationEjjjjjjjS6_jjjj_Z27rsdClAllocationSwitchToHostPKN7android12renderscript7ContextEPKNS0_10AllocationE_Z28rsdAllocationGenerateMipmapsPKN7android12renderscript7ContextEPKNS0_10AllocationE_Z29rsdClAllocationSwitchToDevicePKN7android12renderscript7ContextEPKNS0_10AllocationE_Z30rsdClAllocationGenerateMipmapsPKN7android12renderscript7ContextEPKNS0_10AllocationE_ZNK7android12renderscript7Element11getSizeBitsEv_ZNSt3__16vectorIPKN7android12renderscript10AllocationENS_9allocatorIS5_EEE21__push_back_slow_pathIS5_EEvOT__ZNSt3__16vectorIPKN7android12renderscript10AllocationENS_9allocatorIS5_EEE26__swap_out_circular_bufferERNS_14__split_bufferIS5_RS7_EE__aeabi_memmoveabortpthread_mutex_lockpthread_mutex_unlock_Z13rsdScriptInitPKN7android12renderscript7ContextEPNS0_7ScriptCEPKcS7_PKhjj_Z15rsdClScriptInitPKN7android12renderscript7ContextEPNS0_7ScriptCEPKcS7_PKhjj_Z16rsdInitIntrinsicPKN7android12renderscript7ContextEPNS0_6ScriptE19RsScriptIntrinsicIDPNS0_7ElementE_Z16rsdScriptDestroyPKN7android12renderscript7ContextEPNS0_6ScriptE_Z18rsdClInitIntrinsicPKN7android12renderscript7ContextEPNS0_6ScriptE19RsScriptIntrinsicIDPNS0_7ElementE_Z18rsdClScriptDestroyPKN7android12renderscript7ContextEPNS0_6ScriptE_Z19rsdClSetPtrArgumentPKN7android12renderscript7ContextEPKNS0_6ScriptEP13maliKernelRecjPv_Z19rsdClSetPtrArgumentPKN7android12renderscript7ContextEPKNS0_6ScriptEP13maliKernelRecjPvj_Z19rsdScriptInvokeInitPKN7android12renderscript7ContextEPNS0_6ScriptE_Z19rsdScriptInvokeRootPKN7android12renderscript7ContextEPNS0_6ScriptE_Z21rsdClScriptInvokeInitPKN7android12renderscript7ContextEPNS0_6ScriptE_Z21rsdClScriptInvokeRootPKN7android12renderscript7ContextEPNS0_6ScriptE_Z21rsdScriptGroupExecutePKN7android12renderscript7ContextEPKNS0_15ScriptGroupBaseE_Z21rsdScriptSetGlobalObjPKN7android12renderscript7ContextEPKNS0_6ScriptEjPNS0_10ObjectBaseE_Z21rsdScriptSetGlobalVarPKN7android12renderscript7ContextEPKNS0_6ScriptEjPvj_Z22rsdClSetAllocationAddrPKN7android12renderscript7ContextEPKNS0_10AllocationEPj_Z22rsdScriptInvokeForEachPKN7android12renderscript7ContextEPNS0_6ScriptEjPKNS0_10AllocationEPS6_PKvjPK12RsScriptCall_Z23rsdClScriptGroupExecutePKN7android12renderscript7ContextEPKNS0_15ScriptGroupBaseE_Z23rsdClScriptSetGlobalObjPKN7android12renderscript7ContextEPKNS0_6ScriptEjPNS0_10ObjectBaseE_Z23rsdClScriptSetGlobalVarPKN7android12renderscript7ContextEPKNS0_6ScriptEjPvj_Z23rsdScriptInvokeFunctionPKN7android12renderscript7ContextEPNS0_6ScriptEjPKvj_Z24rsdClRegisterTranslationPKN7android12renderscript7ContextEPKNS0_6ScriptEPv_Z24rsdClScriptInvokeForEachPKN7android12renderscript7ContextEPNS0_6ScriptEjPKNS0_10AllocationEPS6_PKvjPK12RsScriptCall_Z25rsdClScriptInvokeFunctionPKN7android12renderscript7ContextEPNS0_6ScriptEjPKvj_Z26rsdClDeclareUsedAllocationPKN7android12renderscript7ContextEPKNS0_6ScriptEP13maliKernelRecPNS0_10AllocationE_Z27rsdScriptInvokeForEachMultiPKN7android12renderscript7ContextEPNS0_6ScriptEjPPKNS0_10AllocationEjPS6_PKvjPK12RsScriptCall_Z29rsdClScriptInvokeForEachMultiPKN7android12renderscript7ContextEPNS0_6ScriptEjPPKNS0_10AllocationEjPS6_PKvjPK12RsScriptCall_ZN7android12renderscript15RsdCpuReference15getThreadTLSKeyEv_ZNSt3__112basic_stringIcNS_11char_traitsIcEENS_9allocatorIcEEE21__grow_by_and_replaceEjjjjjjPKc_ZNSt3__112basic_stringIcNS_11char_traitsIcEENS_9allocatorIcEEE6appendEPKcj_ZNSt3__113__vector_baseIPN7android12renderscript10AllocationENS_9allocatorIS4_EEED2Ev_ZNSt3__16vectorIPN7android12renderscript10AllocationENS_9allocatorIS4_EEE21__push_back_slow_pathIRKS4_EEvOT__ZNSt3__16vectorIPN7android12renderscript10AllocationENS_9allocatorIS4_EEE26__swap_out_circular_bufferERNS_14__split_bufferIS4_RS6_EE_ZNSt3__1plIcNS_11char_traitsIcEENS_9allocatorIcEEEENS_12basic_stringIT_T0_T1_EERKS9_PKS6_fclosefopenfwritepthread_getspecificstrlen_Z11rsdShutdownPN7android12renderscript7ContextE_Z11rsdTypeInitPKN7android12renderscript7ContextEPKNS0_4TypeE_Z12rsdClHalInitPvjj_Z13rsdClHalAbortPv_Z14rsdElementInitPKN7android12renderscript7ContextEPKNS0_7ElementE_Z14rsdSamplerInitPKN7android12renderscript7ContextEPKNS0_7SamplerE_Z14rsdSetPriorityPKN7android12renderscript7ContextEi_Z14rsdTypeDestroyPKN7android12renderscript7ContextEPKNS0_4TypeE_Z16rsdClHalQueryHalN7android12renderscript14RsHalInitEnumsEPPv_Z17rsdElementDestroyPKN7android12renderscript7ContextEPKNS0_7ElementE_Z17rsdFreeRuntimeMemPv_Z17rsdSamplerDestroyPKN7android12renderscript7ContextEPKNS0_7SamplerE_Z18rsdAllocRuntimeMemjj_Z18rsdScriptGroupInitPKN7android12renderscript7ContextEPNS0_15ScriptGroupBaseE_Z19rsdAllocationIoSendPKN7android12renderscript7ContextEPNS0_10AllocationE_Z21rsdScriptGetGlobalVarPKN7android12renderscript7ContextEPKNS0_6ScriptEjPvj_Z21rsdScriptGroupDestroyPKN7android12renderscript7ContextEPKNS0_15ScriptGroupBaseE_Z21rsdScriptInvokeReducePKN7android12renderscript7ContextEPNS0_6ScriptEjPPKNS0_10AllocationEjPS6_PK12RsScriptCall_Z22rsdAllocationIoReceivePKN7android12renderscript7ContextEPNS0_10AllocationE_Z22rsdScriptGroupSetInputPKN7android12renderscript7ContextEPKNS0_11ScriptGroupEPKNS0_14ScriptKernelIDEPNS0_10AllocationE_Z22rsdScriptSetGlobalBindPKN7android12renderscript7ContextEPKNS0_6ScriptEjPNS0_10AllocationE_Z23rsdScriptGroupSetOutputPKN7android12renderscript7ContextEPKNS0_11ScriptGroupEPKNS0_14ScriptKernelIDEPNS0_10AllocationE_Z24rsdAllocationGrallocBitsPKN7android12renderscript7ContextEPNS0_10AllocationE_Z25rsdTypeUpdateCachedObjectPKN7android12renderscript7ContextEPKNS0_4TypeEPNS0_7rs_typeE_Z26rsdAllocationAdapterOffsetPKN7android12renderscript7ContextEPKNS0_10AllocationE_Z27rsdScriptInvokeFreeChildrenPKN7android12renderscript7ContextEPNS0_6ScriptE_Z27rsdScriptUpdateCachedObjectPKN7android12renderscript7ContextEPKNS0_6ScriptEPNS0_9rs_scriptE_Z28rsdElementUpdateCachedObjectPKN7android12renderscript7ContextEPKNS0_7ElementEPNS0_10rs_elementE_Z28rsdSamplerUpdateCachedObjectPKN7android12renderscript7ContextEPKNS0_7SamplerEPNS0_10rs_samplerE_Z31rsdAllocationUpdateCachedObjectPKN7android12renderscript7ContextEPKNS0_10AllocationEPNS0_13rs_allocationE_Z33rsdScriptSetGlobalVarWithElemDimsPKN7android12renderscript7ContextEPKNS0_6ScriptEjPvjPKNS0_7ElementEPKjj_ZN6bcinfo17MetadataExtractor7extractEvlibbcinfolibbcinfo.so_ZN6bcinfo17MetadataExtractorC1EPKcj_ZN6bcinfo17MetadataExtractorD1Ev_ZN7android12renderscript15RsdCpuReference6createEPNS0_7ContextEjjPFPKNS1_9CpuSymbolES3_PKcEPFPNS1_9CpuScriptES3_PKNS0_6ScriptEEPFS8_S8_jES8__ZN7android12renderscript16RsdCpuScriptImpl12BCC_EXE_PATHEatoi_ZN7android12renderscript15RsdCpuReference12getTlsScriptEv_ZN7android12renderscript15RsdCpuReference13getTlsContextEvrsdDeclareUsedAllocationrsdRegisterTranslatedrsdSetAllocationAddrrsdSetAllocationArgrsdSetPtrArgAHardwareBuffer_lockAHardwareBuffer_unlockANativeWindowBuffer_getHardwareBufferANativeWindow_cancelBufferANativeWindow_dequeueBufferANativeWindow_queueBufferANativeWindow_releaseANativeWindow_setBuffersGeometry_Z12GetOffsetPtrPKN7android12renderscript10AllocationEjjjj23RsAllocationCubemapFace_Z15rsdTypeToGLType10RsDataType_Z17rsdKindToGLFormat10RsDataKind_Z24rsdAllocationInitStridedPKN7android12renderscript7ContextEPNS0_10AllocationEbj_Z32rsdAllocationData2D_alloc_scriptPKN7android12renderscript7ContextEPKNS0_10AllocationEjjj23RsAllocationCubemapFacejjS6_jjjS7__Z32rsdAllocationData3D_alloc_scriptPKN7android12renderscript7ContextEPKNS0_10AllocationEjjjjjjjS6_jjjj_ZN7android12renderscript7rsRoundIjEET_S2_j_ZNK7android12renderscript10Allocation7decRefsEPKvjj_ZNK7android12renderscript10Allocation7incRefsEPKvjj_ZNK7android12renderscript7Element7decRefsEPKv_ZNK7android12renderscript7Element7incRefsEPKvclosememalignrealloc_Z32rsdScriptGetAllocationForPointerPKN7android12renderscript7ContextEPKNS0_6ScriptEPKvrsdHalAbortrsdHalInitrsdHalQueryHalrsdHalQueryVersion_Z10rsIsObject10rs_element_Z10rsIsObject10rs_sampler_Z10rsIsObject13rs_allocation_Z10rsIsObject16rs_program_store_Z10rsIsObject17rs_program_raster_Z10rsIsObject17rs_program_vertex_Z10rsIsObject19rs_program_fragment_Z10rsIsObject7rs_font_Z10rsIsObject7rs_mesh_Z10rsIsObject7rs_type_Z10rsIsObject9rs_script_Z11rsLocaltimeP5rs_tmPKi_Z11rsSetObjectP10rs_elementS__Z11rsSetObjectP10rs_samplerS__Z11rsSetObjectP13rs_allocationS__Z11rsSetObjectP16rs_program_storeS__Z11rsSetObjectP17rs_program_rasterS__Z11rsSetObjectP17rs_program_vertexS__Z11rsSetObjectP19rs_program_fragmentS__Z11rsSetObjectP7rs_fontS__Z11rsSetObjectP7rs_meshS__Z11rsSetObjectP7rs_typeS__Z11rsSetObjectP9rs_scriptS__Z12rsCreateType10rs_elementjjjbb13rs_yuv_format_Z13rsClearObjectP10rs_element_Z13rsClearObjectP10rs_sampler_Z13rsClearObjectP13rs_allocation_Z13rsClearObjectP16rs_program_store_Z13rsClearObjectP17rs_program_raster_Z13rsClearObjectP17rs_program_vertex_Z13rsClearObjectP19rs_program_fragment_Z13rsClearObjectP7rs_font_Z13rsClearObjectP7rs_mesh_Z13rsClearObjectP7rs_type_Z13rsClearObjectP9rs_script_Z13rsUptimeNanosv_Z14rsGetElementAt13rs_allocationj_Z14rsGetElementAt13rs_allocationjj_Z14rsGetElementAt13rs_allocationjjj_Z14rsSendToClienti_Z14rsSendToClientiPKvj_Z14rsSetElementAt13rs_allocationPKvj_Z14rsSetElementAt13rs_allocationPKvjj_Z14rsSetElementAt13rs_allocationPKvjjj_Z14rsUptimeMillisv_Z15rsCreateElementiibj_Z15rsGetAllocationPKv_Z17rsForEachInternaliP14rs_script_calliiP13rs_allocation_Z18rsAllocationIoSend13rs_allocation_Z18rsCreateAllocation7rs_type28rs_allocation_mipmap_controljPv_Z18rsGetElementAt_int13rs_allocationPij_Z18rsGetElementAt_int13rs_allocationPijj_Z18rsGetElementAt_int13rs_allocationPijjj_Z18rsSetElementAt_int13rs_allocationPKij_Z18rsSetElementAt_int13rs_allocationPKijj_Z18rsSetElementAt_int13rs_allocationPKijjj_Z19rsGetElementAt_char13rs_allocationPcj_Z19rsGetElementAt_char13rs_allocationPcjj_Z19rsGetElementAt_char13rs_allocationPcjjj_Z19rsGetElementAt_half13rs_allocationPDhj_Z19rsGetElementAt_half13rs_allocationPDhjj_Z19rsGetElementAt_half13rs_allocationPDhjjj_Z19rsGetElementAt_int213rs_allocationPDv2_ij_Z19rsGetElementAt_int213rs_allocationPDv2_ijj_Z19rsGetElementAt_int213rs_allocationPDv2_ijjj_Z19rsGetElementAt_int313rs_allocationPDv3_ij_Z19rsGetElementAt_int313rs_allocationPDv3_ijj_Z19rsGetElementAt_int313rs_allocationPDv3_ijjj_Z19rsGetElementAt_int413rs_allocationPDv4_ij_Z19rsGetElementAt_int413rs_allocationPDv4_ijj_Z19rsGetElementAt_int413rs_allocationPDv4_ijjj_Z19rsGetElementAt_long13rs_allocationPlj_Z19rsGetElementAt_long13rs_allocationPljj_Z19rsGetElementAt_long13rs_allocationPljjj_Z19rsGetElementAt_long13rs_allocationPxj_Z19rsGetElementAt_long13rs_allocationPxjj_Z19rsGetElementAt_long13rs_allocationPxjjj_Z19rsGetElementAt_uint13rs_allocationPjj_Z19rsGetElementAt_uint13rs_allocationPjjj_Z19rsGetElementAt_uint13rs_allocationPjjjj_Z19rsSetElementAt_char13rs_allocationPKcj_Z19rsSetElementAt_char13rs_allocationPKcjj_Z19rsSetElementAt_char13rs_allocationPKcjjj_Z19rsSetElementAt_half13rs_allocationPKDhj_Z19rsSetElementAt_half13rs_allocationPKDhjj_Z19rsSetElementAt_half13rs_allocationPKDhjjj_Z19rsSetElementAt_int213rs_allocationPKDv2_ij_Z19rsSetElementAt_int213rs_allocationPKDv2_ijj_Z19rsSetElementAt_int213rs_allocationPKDv2_ijjj_Z19rsSetElementAt_int313rs_allocationPKDv3_ij_Z19rsSetElementAt_int313rs_allocationPKDv3_ijj_Z19rsSetElementAt_int313rs_allocationPKDv3_ijjj_Z19rsSetElementAt_int413rs_allocationPKDv4_ij_Z19rsSetElementAt_int413rs_allocationPKDv4_ijj_Z19rsSetElementAt_int413rs_allocationPKDv4_ijjj_Z19rsSetElementAt_long13rs_allocationPKlj_Z19rsSetElementAt_long13rs_allocationPKljj_Z19rsSetElementAt_long13rs_allocationPKljjj_Z19rsSetElementAt_long13rs_allocationPKxj_Z19rsSetElementAt_long13rs_allocationPKxjj_Z19rsSetElementAt_long13rs_allocationPKxjjj_Z19rsSetElementAt_uint13rs_allocationPKjj_Z19rsSetElementAt_uint13rs_allocationPKjjj_Z19rsSetElementAt_uint13rs_allocationPKjjjj_Z20rsGetElementAt_char213rs_allocationPDv2_cj_Z20rsGetElementAt_char213rs_allocationPDv2_cjj_Z20rsGetElementAt_char213rs_allocationPDv2_cjjj_Z20rsGetElementAt_char313rs_allocationPDv3_cj_Z20rsGetElementAt_char313rs_allocationPDv3_cjj_Z20rsGetElementAt_char313rs_allocationPDv3_cjjj_Z20rsGetElementAt_char413rs_allocationPDv4_cj_Z20rsGetElementAt_char413rs_allocationPDv4_cjj_Z20rsGetElementAt_char413rs_allocationPDv4_cjjj_Z20rsGetElementAt_float13rs_allocationPfj_Z20rsGetElementAt_float13rs_allocationPfjj_Z20rsGetElementAt_float13rs_allocationPfjjj_Z20rsGetElementAt_half213rs_allocationPDv2_Dhj_Z20rsGetElementAt_half213rs_allocationPDv2_Dhjj_Z20rsGetElementAt_half213rs_allocationPDv2_Dhjjj_Z20rsGetElementAt_half313rs_allocationPDv3_Dhj_Z20rsGetElementAt_half313rs_allocationPDv3_Dhjj_Z20rsGetElementAt_half313rs_allocationPDv3_Dhjjj_Z20rsGetElementAt_half413rs_allocationPDv4_Dhj_Z20rsGetElementAt_half413rs_allocationPDv4_Dhjj_Z20rsGetElementAt_half413rs_allocationPDv4_Dhjjj_Z20rsGetElementAt_long213rs_allocationPDv2_lj_Z20rsGetElementAt_long213rs_allocationPDv2_ljj_Z20rsGetElementAt_long213rs_allocationPDv2_ljjj_Z20rsGetElementAt_long213rs_allocationPDv2_xj_Z20rsGetElementAt_long213rs_allocationPDv2_xjj_Z20rsGetElementAt_long213rs_allocationPDv2_xjjj_Z20rsGetElementAt_long313rs_allocationPDv3_lj_Z20rsGetElementAt_long313rs_allocationPDv3_ljj_Z20rsGetElementAt_long313rs_allocationPDv3_ljjj_Z20rsGetElementAt_long313rs_allocationPDv3_xj_Z20rsGetElementAt_long313rs_allocationPDv3_xjj_Z20rsGetElementAt_long313rs_allocationPDv3_xjjj_Z20rsGetElementAt_long413rs_allocationPDv4_lj_Z20rsGetElementAt_long413rs_allocationPDv4_ljj_Z20rsGetElementAt_long413rs_allocationPDv4_ljjj_Z20rsGetElementAt_long413rs_allocationPDv4_xj_Z20rsGetElementAt_long413rs_allocationPDv4_xjj_Z20rsGetElementAt_long413rs_allocationPDv4_xjjj_Z20rsGetElementAt_short13rs_allocationPsj_Z20rsGetElementAt_short13rs_allocationPsjj_Z20rsGetElementAt_short13rs_allocationPsjjj_Z20rsGetElementAt_uchar13rs_allocationPhj_Z20rsGetElementAt_uchar13rs_allocationPhjj_Z20rsGetElementAt_uchar13rs_allocationPhjjj_Z20rsGetElementAt_uint213rs_allocationPDv2_jj_Z20rsGetElementAt_uint213rs_allocationPDv2_jjj_Z20rsGetElementAt_uint213rs_allocationPDv2_jjjj_Z20rsGetElementAt_uint313rs_allocationPDv3_jj_Z20rsGetElementAt_uint313rs_allocationPDv3_jjj_Z20rsGetElementAt_uint313rs_allocationPDv3_jjjj_Z20rsGetElementAt_uint413rs_allocationPDv4_jj_Z20rsGetElementAt_uint413rs_allocationPDv4_jjj_Z20rsGetElementAt_uint413rs_allocationPDv4_jjjj_Z20rsGetElementAt_ulong13rs_allocationPmj_Z20rsGetElementAt_ulong13rs_allocationPmjj_Z20rsGetElementAt_ulong13rs_allocationPmjjj_Z20rsGetElementAt_ulong13rs_allocationPyj_Z20rsGetElementAt_ulong13rs_allocationPyjj_Z20rsGetElementAt_ulong13rs_allocationPyjjj_Z20rsSetElementAt_char213rs_allocationPKDv2_cj_Z20rsSetElementAt_char213rs_allocationPKDv2_cjj_Z20rsSetElementAt_char213rs_allocationPKDv2_cjjj_Z20rsSetElementAt_char313rs_allocationPKDv3_cj_Z20rsSetElementAt_char313rs_allocationPKDv3_cjj_Z20rsSetElementAt_char313rs_allocationPKDv3_cjjj_Z20rsSetElementAt_char413rs_allocationPKDv4_cj_Z20rsSetElementAt_char413rs_allocationPKDv4_cjj_Z20rsSetElementAt_char413rs_allocationPKDv4_cjjj_Z20rsSetElementAt_float13rs_allocationPKfj_Z20rsSetElementAt_float13rs_allocationPKfjj_Z20rsSetElementAt_float13rs_allocationPKfjjj_Z20rsSetElementAt_half213rs_allocationPKDv2_Dhj_Z20rsSetElementAt_half213rs_allocationPKDv2_Dhjj_Z20rsSetElementAt_half213rs_allocationPKDv2_Dhjjj_Z20rsSetElementAt_half313rs_allocationPKDv3_Dhj_Z20rsSetElementAt_half313rs_allocationPKDv3_Dhjj_Z20rsSetElementAt_half313rs_allocationPKDv3_Dhjjj_Z20rsSetElementAt_half413rs_allocationPKDv4_Dhj_Z20rsSetElementAt_half413rs_allocationPKDv4_Dhjj_Z20rsSetElementAt_half413rs_allocationPKDv4_Dhjjj_Z20rsSetElementAt_long213rs_allocationPKDv2_lj_Z20rsSetElementAt_long213rs_allocationPKDv2_ljj_Z20rsSetElementAt_long213rs_allocationPKDv2_ljjj_Z20rsSetElementAt_long213rs_allocationPKDv2_xj_Z20rsSetElementAt_long213rs_allocationPKDv2_xjj_Z20rsSetElementAt_long213rs_allocationPKDv2_xjjj_Z20rsSetElementAt_long313rs_allocationPKDv3_lj_Z20rsSetElementAt_long313rs_allocationPKDv3_ljj_Z20rsSetElementAt_long313rs_allocationPKDv3_ljjj_Z20rsSetElementAt_long313rs_allocationPKDv3_xj_Z20rsSetElementAt_long313rs_allocationPKDv3_xjj_Z20rsSetElementAt_long313rs_allocationPKDv3_xjjj_Z20rsSetElementAt_long413rs_allocationPKDv4_lj_Z20rsSetElementAt_long413rs_allocationPKDv4_ljj_Z20rsSetElementAt_long413rs_allocationPKDv4_ljjj_Z20rsSetElementAt_long413rs_allocationPKDv4_xj_Z20rsSetElementAt_long413rs_allocationPKDv4_xjj_Z20rsSetElementAt_long413rs_allocationPKDv4_xjjj_Z20rsSetElementAt_short13rs_allocationPKsj_Z20rsSetElementAt_short13rs_allocationPKsjj_Z20rsSetElementAt_short13rs_allocationPKsjjj_Z20rsSetElementAt_uchar13rs_allocationPKhj_Z20rsSetElementAt_uchar13rs_allocationPKhjj_Z20rsSetElementAt_uchar13rs_allocationPKhjjj_Z20rsSetElementAt_uint213rs_allocationPKDv2_jj_Z20rsSetElementAt_uint213rs_allocationPKDv2_jjj_Z20rsSetElementAt_uint213rs_allocationPKDv2_jjjj_Z20rsSetElementAt_uint313rs_allocationPKDv3_jj_Z20rsSetElementAt_uint313rs_allocationPKDv3_jjj_Z20rsSetElementAt_uint313rs_allocationPKDv3_jjjj_Z20rsSetElementAt_uint413rs_allocationPKDv4_jj_Z20rsSetElementAt_uint413rs_allocationPKDv4_jjj_Z20rsSetElementAt_uint413rs_allocationPKDv4_jjjj_Z20rsSetElementAt_ulong13rs_allocationPKmj_Z20rsSetElementAt_ulong13rs_allocationPKmjj_Z20rsSetElementAt_ulong13rs_allocationPKmjjj_Z20rsSetElementAt_ulong13rs_allocationPKyj_Z20rsSetElementAt_ulong13rs_allocationPKyjj_Z20rsSetElementAt_ulong13rs_allocationPKyjjj_Z20rsdLookupRuntimeStubPN7android12renderscript7ContextEPKc_Z21rsAllocationIoReceive13rs_allocation_Z21rsGetElementAt_double13rs_allocationPdj_Z21rsGetElementAt_double13rs_allocationPdjj_Z21rsGetElementAt_double13rs_allocationPdjjj_Z21rsGetElementAt_float213rs_allocationPDv2_fj_Z21rsGetElementAt_float213rs_allocationPDv2_fjj_Z21rsGetElementAt_float213rs_allocationPDv2_fjjj_Z21rsGetElementAt_float313rs_allocationPDv3_fj_Z21rsGetElementAt_float313rs_allocationPDv3_fjj_Z21rsGetElementAt_float313rs_allocationPDv3_fjjj_Z21rsGetElementAt_float413rs_allocationPDv4_fj_Z21rsGetElementAt_float413rs_allocationPDv4_fjj_Z21rsGetElementAt_float413rs_allocationPDv4_fjjj_Z21rsGetElementAt_short213rs_allocationPDv2_sj_Z21rsGetElementAt_short213rs_allocationPDv2_sjj_Z21rsGetElementAt_short213rs_allocationPDv2_sjjj_Z21rsGetElementAt_short313rs_allocationPDv3_sj_Z21rsGetElementAt_short313rs_allocationPDv3_sjj_Z21rsGetElementAt_short313rs_allocationPDv3_sjjj_Z21rsGetElementAt_short413rs_allocationPDv4_sj_Z21rsGetElementAt_short413rs_allocationPDv4_sjj_Z21rsGetElementAt_short413rs_allocationPDv4_sjjj_Z21rsGetElementAt_uchar213rs_allocationPDv2_hj_Z21rsGetElementAt_uchar213rs_allocationPDv2_hjj_Z21rsGetElementAt_uchar213rs_allocationPDv2_hjjj_Z21rsGetElementAt_uchar313rs_allocationPDv3_hj_Z21rsGetElementAt_uchar313rs_allocationPDv3_hjj_Z21rsGetElementAt_uchar313rs_allocationPDv3_hjjj_Z21rsGetElementAt_uchar413rs_allocationPDv4_hj_Z21rsGetElementAt_uchar413rs_allocationPDv4_hjj_Z21rsGetElementAt_uchar413rs_allocationPDv4_hjjj_Z21rsGetElementAt_ulong213rs_allocationPDv2_mj_Z21rsGetElementAt_ulong213rs_allocationPDv2_mjj_Z21rsGetElementAt_ulong213rs_allocationPDv2_mjjj_Z21rsGetElementAt_ulong213rs_allocationPDv2_yj_Z21rsGetElementAt_ulong213rs_allocationPDv2_yjj_Z21rsGetElementAt_ulong213rs_allocationPDv2_yjjj_Z21rsGetElementAt_ulong313rs_allocationPDv3_mj_Z21rsGetElementAt_ulong313rs_allocationPDv3_mjj_Z21rsGetElementAt_ulong313rs_allocationPDv3_mjjj_Z21rsGetElementAt_ulong313rs_allocationPDv3_yj_Z21rsGetElementAt_ulong313rs_allocationPDv3_yjj_Z21rsGetElementAt_ulong313rs_allocationPDv3_yjjj_Z21rsGetElementAt_ulong413rs_allocationPDv4_mj_Z21rsGetElementAt_ulong413rs_allocationPDv4_mjj_Z21rsGetElementAt_ulong413rs_allocationPDv4_mjjj_Z21rsGetElementAt_ulong413rs_allocationPDv4_yj_Z21rsGetElementAt_ulong413rs_allocationPDv4_yjj_Z21rsGetElementAt_ulong413rs_allocationPDv4_yjjj_Z21rsGetElementAt_ushort13rs_allocationPtj_Z21rsGetElementAt_ushort13rs_allocationPtjj_Z21rsGetElementAt_ushort13rs_allocationPtjjj_Z21rsSetElementAt_double13rs_allocationPKdj_Z21rsSetElementAt_double13rs_allocationPKdjj_Z21rsSetElementAt_double13rs_allocationPKdjjj_Z21rsSetElementAt_float213rs_allocationPKDv2_fj_Z21rsSetElementAt_float213rs_allocationPKDv2_fjj_Z21rsSetElementAt_float213rs_allocationPKDv2_fjjj_Z21rsSetElementAt_float313rs_allocationPKDv3_fj_Z21rsSetElementAt_float313rs_allocationPKDv3_fjj_Z21rsSetElementAt_float313rs_allocationPKDv3_fjjj_Z21rsSetElementAt_float413rs_allocationPKDv4_fj_Z21rsSetElementAt_float413rs_allocationPKDv4_fjj_Z21rsSetElementAt_float413rs_allocationPKDv4_fjjj_Z21rsSetElementAt_short213rs_allocationPKDv2_sj_Z21rsSetElementAt_short213rs_allocationPKDv2_sjj_Z21rsSetElementAt_short213rs_allocationPKDv2_sjjj_Z21rsSetElementAt_short313rs_allocationPKDv3_sj_Z21rsSetElementAt_short313rs_allocationPKDv3_sjj_Z21rsSetElementAt_short313rs_allocationPKDv3_sjjj_Z21rsSetElementAt_short413rs_allocationPKDv4_sj_Z21rsSetElementAt_short413rs_allocationPKDv4_sjj_Z21rsSetElementAt_short413rs_allocationPKDv4_sjjj_Z21rsSetElementAt_uchar213rs_allocationPKDv2_hj_Z21rsSetElementAt_uchar213rs_allocationPKDv2_hjj_Z21rsSetElementAt_uchar213rs_allocationPKDv2_hjjj_Z21rsSetElementAt_uchar313rs_allocationPKDv3_hj_Z21rsSetElementAt_uchar313rs_allocationPKDv3_hjj_Z21rsSetElementAt_uchar313rs_allocationPKDv3_hjjj_Z21rsSetElementAt_uchar413rs_allocationPKDv4_hj_Z21rsSetElementAt_uchar413rs_allocationPKDv4_hjj_Z21rsSetElementAt_uchar413rs_allocationPKDv4_hjjj_Z21rsSetElementAt_ulong213rs_allocationPKDv2_mj_Z21rsSetElementAt_ulong213rs_allocationPKDv2_mjj_Z21rsSetElementAt_ulong213rs_allocationPKDv2_mjjj_Z21rsSetElementAt_ulong213rs_allocationPKDv2_yj_Z21rsSetElementAt_ulong213rs_allocationPKDv2_yjj_Z21rsSetElementAt_ulong213rs_allocationPKDv2_yjjj_Z21rsSetElementAt_ulong313rs_allocationPKDv3_mj_Z21rsSetElementAt_ulong313rs_allocationPKDv3_mjj_Z21rsSetElementAt_ulong313rs_allocationPKDv3_mjjj_Z21rsSetElementAt_ulong313rs_allocationPKDv3_yj_Z21rsSetElementAt_ulong313rs_allocationPKDv3_yjj_Z21rsSetElementAt_ulong313rs_allocationPKDv3_yjjj_Z21rsSetElementAt_ulong413rs_allocationPKDv4_mj_Z21rsSetElementAt_ulong413rs_allocationPKDv4_mjj_Z21rsSetElementAt_ulong413rs_allocationPKDv4_mjjj_Z21rsSetElementAt_ulong413rs_allocationPKDv4_yj_Z21rsSetElementAt_ulong413rs_allocationPKDv4_yjj_Z21rsSetElementAt_ulong413rs_allocationPKDv4_yjjj_Z21rsSetElementAt_ushort13rs_allocationPKtj_Z21rsSetElementAt_ushort13rs_allocationPKtjj_Z21rsSetElementAt_ushort13rs_allocationPKtjjj_Z22rsGetElementAt_double213rs_allocationPDv2_dj_Z22rsGetElementAt_double213rs_allocationPDv2_djj_Z22rsGetElementAt_double213rs_allocationPDv2_djjj_Z22rsGetElementAt_double313rs_allocationPDv3_dj_Z22rsGetElementAt_double313rs_allocationPDv3_djj_Z22rsGetElementAt_double313rs_allocationPDv3_djjj_Z22rsGetElementAt_double413rs_allocationPDv4_dj_Z22rsGetElementAt_double413rs_allocationPDv4_djj_Z22rsGetElementAt_double413rs_allocationPDv4_djjj_Z22rsGetElementAt_ushort213rs_allocationPDv2_tj_Z22rsGetElementAt_ushort213rs_allocationPDv2_tjj_Z22rsGetElementAt_ushort213rs_allocationPDv2_tjjj_Z22rsGetElementAt_ushort313rs_allocationPDv3_tj_Z22rsGetElementAt_ushort313rs_allocationPDv3_tjj_Z22rsGetElementAt_ushort313rs_allocationPDv3_tjjj_Z22rsGetElementAt_ushort413rs_allocationPDv4_tj_Z22rsGetElementAt_ushort413rs_allocationPDv4_tjj_Z22rsGetElementAt_ushort413rs_allocationPDv4_tjjj_Z22rsSendToClientBlockingi_Z22rsSendToClientBlockingiPKvj_Z22rsSetElementAt_double213rs_allocationPKDv2_dj_Z22rsSetElementAt_double213rs_allocationPKDv2_djj_Z22rsSetElementAt_double213rs_allocationPKDv2_djjj_Z22rsSetElementAt_double313rs_allocationPKDv3_dj_Z22rsSetElementAt_double313rs_allocationPKDv3_djj_Z22rsSetElementAt_double313rs_allocationPKDv3_djjj_Z22rsSetElementAt_double413rs_allocationPKDv4_dj_Z22rsSetElementAt_double413rs_allocationPKDv4_djj_Z22rsSetElementAt_double413rs_allocationPKDv4_djjj_Z22rsSetElementAt_ushort213rs_allocationPKDv2_tj_Z22rsSetElementAt_ushort213rs_allocationPKDv2_tjj_Z22rsSetElementAt_ushort213rs_allocationPKDv2_tjjj_Z22rsSetElementAt_ushort313rs_allocationPKDv3_tj_Z22rsSetElementAt_ushort313rs_allocationPKDv3_tjj_Z22rsSetElementAt_ushort313rs_allocationPKDv3_tjjj_Z22rsSetElementAt_ushort413rs_allocationPKDv4_tj_Z22rsSetElementAt_ushort413rs_allocationPKDv4_tjj_Z22rsSetElementAt_ushort413rs_allocationPKDv4_tjjj_Z23rsAllocationCopy1DRange13rs_allocationjjjS_jj_Z23rsAllocationCopy2DRange13rs_allocationjjj26rs_allocation_cubemap_facejjS_jjjS0__Z6rsTimePi_Z7rsDebugPKcDv2_x_Z7rsDebugPKcDv2_y_Z7rsDebugPKcDv3_x_Z7rsDebugPKcDv3_y_Z7rsDebugPKcDv4_x_Z7rsDebugPKcDv4_y_Z7rsDebugPKcPK12rs_matrix2x2_Z7rsDebugPKcPK12rs_matrix3x3_Z7rsDebugPKcPK12rs_matrix4x4_Z7rsDebugPKcPKDv2_c_Z7rsDebugPKcPKDv2_d_Z7rsDebugPKcPKDv2_f_Z7rsDebugPKcPKDv2_fPKDv2_t_Z7rsDebugPKcPKDv2_h_Z7rsDebugPKcPKDv2_i_Z7rsDebugPKcPKDv2_j_Z7rsDebugPKcPKDv2_l_Z7rsDebugPKcPKDv2_m_Z7rsDebugPKcPKDv2_s_Z7rsDebugPKcPKDv2_t_Z7rsDebugPKcPKDv2_x_Z7rsDebugPKcPKDv2_y_Z7rsDebugPKcPKDv3_c_Z7rsDebugPKcPKDv3_d_Z7rsDebugPKcPKDv3_f_Z7rsDebugPKcPKDv3_fPKDv3_t_Z7rsDebugPKcPKDv3_h_Z7rsDebugPKcPKDv3_i_Z7rsDebugPKcPKDv3_j_Z7rsDebugPKcPKDv3_l_Z7rsDebugPKcPKDv3_m_Z7rsDebugPKcPKDv3_s_Z7rsDebugPKcPKDv3_t_Z7rsDebugPKcPKDv3_x_Z7rsDebugPKcPKDv3_y_Z7rsDebugPKcPKDv4_c_Z7rsDebugPKcPKDv4_d_Z7rsDebugPKcPKDv4_f_Z7rsDebugPKcPKDv4_fPKDv4_t_Z7rsDebugPKcPKDv4_h_Z7rsDebugPKcPKDv4_i_Z7rsDebugPKcPKDv4_j_Z7rsDebugPKcPKDv4_l_Z7rsDebugPKcPKDv4_m_Z7rsDebugPKcPKDv4_s_Z7rsDebugPKcPKDv4_t_Z7rsDebugPKcPKDv4_x_Z7rsDebugPKcPKDv4_y_Z7rsDebugPKcPKv_Z7rsDebugPKcc_Z7rsDebugPKcd_Z7rsDebugPKcf_Z7rsDebugPKcff_Z7rsDebugPKcfff_Z7rsDebugPKcffff_Z7rsDebugPKcft_Z7rsDebugPKch_Z7rsDebugPKci_Z7rsDebugPKcj_Z7rsDebugPKcl_Z7rsDebugPKcm_Z7rsDebugPKcs_Z7rsDebugPKct_Z7rsDebugPKcx_Z7rsDebugPKcy_Z7rsGetDtv_Z9rsForEach9rs_script13rs_allocationS0__Z9rsForEach9rs_script13rs_allocationS0_PKv_Z9rsForEach9rs_script13rs_allocationS0_PKvPK14rs_script_call_Z9rsForEach9rs_script13rs_allocationS0_PKvj_Z9rsForEach9rs_script13rs_allocationS0_PKvjPK14rs_script_call_ZN7android12renderscript10rsrForEachEPNS0_7ContextEPNS0_6ScriptEjjPPNS0_10AllocationES6_PKvjPK12RsScriptCall_ZN7android12renderscript11rsrToClientEPNS0_7ContextEiPKvi_ZN7android12renderscript12rsrLocalTimeEPNS0_7ContextEP2tmPl_ZN7android12renderscript12rsrSetObjectEPKNS0_7ContextEPNS0_14rs_object_baseEPKNS0_10ObjectBaseE_ZN7android12renderscript13rsrTypeCreateEPNS0_7ContextEPvjjjbbj_ZN7android12renderscript14rsrClearObjectEPNS0_14rs_object_baseE_ZN7android12renderscript14rsrUptimeNanosEPNS0_7ContextE_ZN7android12renderscript15rsrUptimeMillisEPNS0_7ContextE_ZN7android12renderscript16rsrElementCreateEPNS0_7ContextE10RsDataType10RsDataKindbj_ZN7android12renderscript19rsrAllocationIoSendEPNS0_7ContextEPNS0_10AllocationE_ZN7android12renderscript19rsrToClientBlockingEPNS0_7ContextEiPKvi_ZN7android12renderscript22rsrAllocationIoReceiveEPNS0_7ContextEPNS0_10AllocationE_ZN7android12renderscript24rsrAllocationCopy1DRangeEPNS0_7ContextEPNS0_10AllocationEjjjS4_jj_ZN7android12renderscript24rsrAllocationCopy2DRangeEPNS0_7ContextEPNS0_10AllocationEjjjjjjS4_jjjj_ZN7android12renderscript24rsrAllocationCreateTypedEPNS0_7ContextEPv25RsAllocationMipmapControljj_ZN7android12renderscript7rsrTimeEPNS0_7ContextEPl_ZN7android12renderscript8rsrGetDtEPNS0_7ContextEPKNS0_6ScriptE_ZNK7android12renderscript10ObjectBase10decUserRefEv_Z32rsdScriptGroupUpdateCachedObjectPKN7android12renderscript7ContextEPKNS0_11ScriptGroupEPNS0_15rs_script_groupE_edata__bss_start_endlibRS_internal.solibRSCpuRef.soliblog.solibnativewindow.solibc++.so U P/@dd  B "K0$L db)p "@*`V fHB0 I pȠe[@@0@J@GB BUBSŦ ,H!Q RS 61( P@@" sE!# hƖ`bH@8&:C2#! j<" 0ƀ!@4bE) 3(Xr(D$8AOpN2A@% E0n@) ` d5 B8CEZ @P z,ed $$@ A!)C8"@6 H[ $C  Q dI AADUVWYZ[_bdehijkmptvwyz{~     !"&'(+-02568;<@BCFHIKLOPQSTUVYZ_`abcefijknoqtuvwz|      $%'*,/2347;<=>?@BDGHJNOPQRTUXZ\]_aceklnotvwxz{|}~ oc B=zQ#ƃE! ]^#QYg,8ծq*7ѧxz[_KSU}2_Ӡ;Y*6KqxE>tVTJLy W[]xJ '`{{2h/l `YMzXʼnK((fWj#7i$~>؞9V}gk&sLp, &v&S +Td6d|ړZ-/|21E96د{kt9e7ǯ$ äj㓷HhC4I$:̽oCK 3[_]@|W]7B6(8Dyt9eӻ'&~Ic(k;(l!n2_֤+@V'iU =d=bT=bfNgFSUa0H&%.^$l=5>&* mIլr̷/ mI#w-X mI mIMҳø1mImImIfmIA$#S3Ɓ- ꨻-61mIJmI,h +D}mIltmIЮa[Q}0~b s~>{=LfD+",Oe@giQ4Y4>S}2$yՁb^4bTaS1&i\CEc(`c(`="hAHmIe(`KmIg(`LmIi(`si(`]9 3OmINmIk(`hGPmIm(`SmIO$SmITޗ7-s(`оg.S*Nvs(`DYmI?r[mI:w(`y(`ު#_mIaN_mIYbyh4"wgpzguOQ4h9e|O/6aiU~/C Q]Nj6\fP;1ZDmImIa8NJLmI?u;mI\hlmImIIs[mI{umId7Ϡ՚A jpoNBrmIڂtXmIL8lE ܫmImI;Z1 7:PRPl[R-ahIZ3xkC, "1*GjPl"\Hщ]W\n`cf@P΅Q38]YH}*[OI 08v$_j b 8"p|)+ *Л}2c#M'2GyV-J(MuR&r0as Ca$ŠiRx&H9[A3[1b7"=+ gOgD׶.OKoZ@]MCס%.v8rDe0 R=eX)㾻|^%a(P8sL# % ĺV`e%4ApECy `k0&\i]!_u5ёy3%2U L}YRd.4ڋ[3Ψ#Ck- >s/!, C>N4S\#Qi#;m\RГ>x{,l)?q1CwVtΖtݕi_yRqXނ 4eE!sVwlN=vZ?B'-W-*2QeeĚDuK+ui$0ۍ9%ĝ@xJD]ݲh/C0c >i H c  c S?Sf8fmmm m}mem&m mmmm1m/m-m,m.m3m2mmmJmm(mAm=m<mmPm6m5n7n:nB n@n;n8n4n9 nV$n(n?,n>0n4nV8n !F4kGPA"GFGT8I9J #FyDzDv  4I4J yDzDj;01I1J yDzD`/I# yDZ -IyD h h S`O6 [RS-DIAG] Launching GPU kernel intrinsic : id(%d) GPU launch failure, error code: %dAttempting CPU fallback -  L  [  N F#jj h I yD h Error: no memory object set to this allocation: %p pFH FxDhh tHxDh0baȹ 0j.bFG`K I J yDzD HxDhh@ FpnMdM  MpF H%FxDhh h)ah " ci0h& `F"G[IJ yDzDX`h `!)pIyD h h  pML  zLF |ch 40F#GFc I J yDzD`h!  th  FHxDh0 `hahkG FKFHxDh0 `hahkG F@KpG-OlFoFFKHFFF#ZFxDhh  ZF#$#ZFhjjjO3"VL !" 8 G--JjJo!" *:  "*#  "*! * * `G8"j"HC!`G HxDhh@ѧ FNK;GJpFF F FT 2FbHxDh0 ` F-h)FO FpDJ𵅰FH F#xDhhj F9( 0j:F+Fhh F1F61h F: IyD h hJI-OUH FxDhhG SH xDhhG &o#%!GW<3FJ %FfomUR!GF+GHxDhhW@0ѧ (FF !( 0!F-,`0G(`W P'J# zD?%AIIRenderScriptIntrinsic 3dLUT: cannot create program with source: %d-O1Intrinsic Blur BUILD_LOG: %s lut3dFH pG,0FF F(hB 5,,` F@5𵁰F0 FFF39F2F+F5FHxDh0 `j b FrGFHxDh0 `j F@NG𵅰F%HF FF#xDhhe F(, F(&  jhQ'phk"8h!G :F# :F#9h F IyD h h`GFpFF F FjH2FxDh0 ` F-h)F Fp|FZ * !% "!"+#+ ! !!! ! !! ! ! ! !# O-O\HFxDhhG ZH xDhhG p`#$!GW<3MP %It`mD X!GF+G?l @ApFF F F Pp2F@ HxDh0 ` F-h)F @I yD b Fp@rsdGpuScriptIntrinsicBlur: Init Gpu failed-O-FH t P%A xDhh) $U  " F# J& #h zD"G'K"vIvJyDzD^#  "hp"iI"yDG'+ cIcJyDzD29P @ AF " F)F h&"#h RJzD"G'˱"LILJyDzD JIyD h h)w*AFFeF !0B8 ( FAF(# (h F hh F ((( FAF# v(Ah F @hh F w((TG (TG (TG  >@4qpntp?-OFHF!FxDhhG p?(F H xDhhG p`#$!GW<;I yD &It`mD !GF+GE?x=-OFdHFxDhhG aH xDhhG p`#$!GW<;TIU yD %Jt`mD ^!GF+G&/FW $6N.,)F@" F@ U'   g7& /e l@ I yD@ MCvIntrinsic ColorMatrix: Global varible slot unknow slot: %d𵁰F FFF 9F2F+F #-GF,HFFOxDhh  jhQ)P!"(h`G(h`!"G  #*F *F#  (hpX!@"GO(hpAF"G)h F IyD h h /h/pFFF F &J!H$I"K|zDxDyDD{DD+JhzD0+J `,! ;L@ zD+JzD+#"@* *F. F6h1F  FpJ/sconv33_identityconv33_embossconv33_outline-CF6H,jF#xDhh F X(Kh9&/,)O(/$ FJF-8E&-&` (*}`ap`( F!JF &%JF@`(D`a F  F!JF ( IyD h h D.--OFlHFH9bFL!#RxDhhG X@+%01\)(F _H xDhhG p`#$!GW<3RU &It`mD _!GF+GinitGPU: SuccessL*jpGFF) Op I HyDxD D!F@rsAssert failed: %s, in %s at %islot == 1FHxDh0 ` lhIhG`l `d F @ O𵁰FH FFF 9F2F+F pFF F F D 2FbHxDh0 ` F-h)F  Fp\-CF+H<FF#xDhhj F %(3 F! _(-`j h(jhhk"8h!GHF)F:F# HF)F:F# j)F:F#HF 9h F IyD h h V-OUH FxDhhG SH xDhhG &o#%!GW<3FJ %FfomUR!GF+G@BCpr-OFHxDhhG F F F %Fvl !GFIyD %HxDhhW@@ـ (FFW (!! FQF"GF{I|yDW %(  O OG$ Z+@"#GFkIJ yDzD W  Eӆ@F Wl0F FZ+@2FCFGFnIoJ yDzDbIcJ 3FyDzD O_I@F "yD (Z+ "#@GF+YIYJ yDzDW$ ( WW l0F ZFZ+2F@FGFG$lFIGJ yDzDAIBJ 3FyDzD T G$ 0F6" ;(|W$ $@F !Z+P+W$ ( ?!@F  6>clGetPlatformIDsclGetPlatformInfoError on querying amount of platforms, error code: %dError on reading platforms, error code: %dARM |%lp FIFyD h h pFt1F"#(FG3I yD IyD h hp ZclCreateCommandQueue Failed creating CL command queue, error code: %d&IF yD g %s IF yD YCONTEXT CALLBACK: %s jj}pG-OFHF FxDhh (F 0(H  !& uF /(p9>u8`>q>s`j!(!&`'C("H(OA&#@ "G FcId yD B&HxDhh @@0F H0!hf FF"G F:a_I`yDL!#1FGF+ gIhJyDzD T0FG &  8iB&!* 0[D*0 1QD+"0h2DA;1B "8u9``8s8q(F!F 8iAF l j"l 11((F!F"# K (F N :B`0 E(F J 7ZGPU buffer creation failed, reverting to CPU, error code: %d \GPU buffer mapping failed, reverting to CPU, error code: %d 8 -C F:IFyD h h !jj9y)]9})Zi9h#% $ HS2"3HxDhh@q(F!F2F;FG (F!F <G(F!F2F;F (C(F!F TG 0#@p#"GcIJ yDzD  (F!F (F!F"  \  HxDhh@ u-A FF jj8{(T8}(QѸh@F 6`-h9h# m3#F"GFCI yD 6 hB&!* 0[D*0 1QD+"0h2DA;1B 8s 8q>Reading CL buffer failed, error code: %d-A FF jj0{0}0(F!FA θ/hi01hO\ "#GFK I  yD  (F!F JWriting CL buffer failed, error code: %dpF" FF 0F)F"Fp@ ະ" FF (F!F@ ܺpF" FF 0F)F"Fp@ ֺ𵁰F"FFF 8F1F*F#F@ ̺𵁰F"FFF 8F1F*F#F@ J𵁰F"FFF 8F1F*F#F@ 𵁰F"FFF 8F1F*F#F@ 𵁰F"FFF 8F1F*F#F@ 𵁰F"FFF 8F1F*F#F@ 𒺰" FF t(F!F@ 𵁰FF"FF b8F1F" ]8F1F*F#F@ 𵁰F F"FF J8F1F" E8F1F*F#F@ t𵁰F F"FF 28F1F" -8F1F*F#F@ d𵁰F"FFF 8F1F*F#F@ Z𵁰F"FFF 8F1F*F#F@ P" FF (F!F@ L𵇰F.H FxDhh!Oңho`GB _B8FdzO , o@G "*h@! 0 F BQo"D 9HxDhh@' $, FF`hR*`` `h)h(`a`hhhi``hh```h `𵃰IF#FyD, h hh jjh I #FyD  #(hP1F"GXIJ 3FyDzD  HxDhh@ Error: no memory object set to this allocation: %p jEj-AiFi i&U&pB36Bhp!F"GXIJ #FyDzD q -GFF)i)Li/Iкi.B $W$ B64B4jPE HF (FW$0FJF jPEQF3FO (hpJFGFjPE0F  I J SFyDzD I yD   ,rsdClSetPtrArgument: failed to malloc memory for bufferFFjjhhInG `HxDhh B`Pj0Pb iF HxDhh@ n-OF$.zDhh2^h*[(Y)W(Q jO`0.(F y)F` l,. O 2RF)O!0GF8,2yDe`RFt((FO #IFG(I,yD O(xDhh2@@3h(̀$! V (F:F=F 0GF; ڢyD %0 FD(@hA ` 0F F0F 0F PF2F  }',(fO @`"#P*2GA*XIJ SFyDzD  H0P* Q*tB"DEGA*KIJ yDzD  L`#Q* P*DGF+ ЪIJ yDzD    E /F $h1B:HJ)FxDzD *(F .(J D@FizD JDH MF(F EFphzI{J+FyDzD A!;F ;8F @sIsJ +FyDzD /D 8`h(?ծ&Ae"fp`hGF+ f bIbJ yDzD   (?`hp!" AfGF+ SITJ yDzD  5PJi@FzD DBLF F FGIHJ#FyDzD  !+F (F ?I@J #FyDzD D [ BFailed to create program %d, error code: %d rFailed to build program %d, error code: %d Cannot create root kernel, error code: %d 1wb+Xt2pE7Sf3@~v`-OF FF`@F~hd.V0"z_vN,{,^̲ =N>*NL1D*׽p FFFf$ I J HyDzDxD 5`pht`pĺ,tls-AFFFF ! 5+ ЫR0FAF"F#;hR0FAF"F0h8 I J #FyDzDk0F0FAF"F;FAϾp FF F)FF(F!F0F)Fp@-OFF P @0CIDJD@HyDzDxD .m@0@I@JA=HyDzDxD (L8hiGF 0 & hP&p9iiOi(P(imG9iExiii8hG _8F6^E oG hh jG`F0xаh0FHF9FO=ι6drvdc-AFF FF) j@i9hhBF3F)FA`G𵁰FFFF8j@i1h h"FKi)F@G FF(F!F@F h`ahB o#Da`` F𵇰F.H FxDhh!Oңho`GB _B8FdzO o@G~"*h@! 0 FBQo"D HxDhh@'n𵅰FF x # h SFfhB `QF3F Fw**FhgL!xpA!p``!9T F-OF FoF@BU x(h`OppB PE8h (O o0 0@FF0FRF,  ZF   XDQDID -j H!`1T$F@0BA0B<ܳzoG@pBnzo@pB@xDh`(A xDhApp(?Z{^bfHxDht@0BB@0BB@0B@kHxDhczp(QHxDhI@pBA0BDуHxDh=A0B'A0B8{HxDh1~HxDh-HxDh)|HxDh%vHxDh!oHxDhlHxDhiHxDhfHxDhjHxDh Z* _HxDHxD ` pGI yDy {HxDhxHxDhuHxDhrHxDhoHxDhlHxDhiHxDhfHxDhcHxDh`HxDh]HxDhZHxDhWHxDhTHxDhQHxDhNHxDhKHxDhHHxDhEHxDhBHxDh?HxDhHxDhg;HxDhc8HxDh_5HxDh[2HxDhW/HxDhS>HxDhO;HxDhKT`lx*d$ $0<HT`lxx ġERROR: unknown Mali RenderScript HAL API query, %iF F F@-CF9HF9IF9 #xDyDhh2 4h0F(Lі1 IB9F8FC AIA yDF"%9FG8F(%PCIDJ yDzDBHBIBF#xDyDh`@HIFxD?HxD>HxD0F(`a I yD 6IyD h h23 ARM mali (%d bit) RenderScript Compute DriverT[RS-DIAG] GPU Init failed.debug.rs.script.arm0Ğ[RS-DIAG] GPU Init succeeded.debug.rs.cpufbon˱d ZjhpG FFHxDhh060!( F I  yD0$, HxDhh@ F(j2Could not extract metadata from bitcode/vendor/lib/libclcore_neon.bc/vendor/lib/libclcore.bc2F F F@𵁰F FFFF8F2F+F-AFFFFFF0F:FCFTFFF(F"F@Ĺ FF)F"F@p FFFF0F*F#Fp@8 (IyDQ pG pG8(IyDQ pG pG p k p-G F ,!FFFճ.%b1p0@.ѡj F"3FO Fj`)Сj FRF3FEOpjIjJkKgHyDzD{DxD  `jA&H"cIcyD 'K!l)L@(XѠjTP)U=I>yD= @Q*HQ`jmm: *LK{DS" "*amm:*GK{DS" "(ja hw lx*j !#3#!!F"'8F@FIF2FjF|'(F'b'I'JyDzDjm0xCIJ yDzD'@FIF2F/w@FR (w@KĚUser-allocated buffers must not have multiple faces or LODsح"wYJؿrequiredAlignment must be power of 2-OFpHF!FF'FxDhhjPj0kG8ame0YFhCO3Fx`  PYlDP H! 0B B RFB B P-*F"PCBF"WL 9FD(v"Q;J"02BQ)gJO '&F F'm0 #F4F-BF  Q)F!"QCfO -(O @  7)E(0a!W! 1D@+!BaaF(HNxDhh@0F[8pFFF!FF%(F!F(Fp-OF`P8 F) Xa YY\VadEB&ReF>XPF)FZF#CPF!FZF#E F)F:FED>DXRD4(Q(EYa1!B  O OiS~C^B-Q8F F5heF F)F:F10h=D DFRFDE Hw h0+#׾-Omسn O %  FF25f G 6h>``nA8kpNn@tC`@ C eEӽ F ,!(!F b  p!H FxDhhm%jHX)1,` jAA(njirO1iFti0Fw0F| ahb(F bHxDhh@p0-O FFF,B0h'()F0 F+;F"F@FIFKF0h)F(F F;FEOp"I"J#KHyDzD{DxD E ١k  QCI yDOResize cannot be called on a USAGE_SHARED allocatione(VpF F-є,"%JzD%&j--@ IJKHyDzD{DxD G-jп[ pwpjmm )Ap@渧jTCannot only sync from RGBArender target!j!AwpG-AF%HFFxDhh4jgjicO1iFe@@F!|!i8F^8F` a`b"(F#n(@F!@F1F*FHeb"HxDhh@Error unlocking output buffer.Error canceling output buffer.Error setting IO output buffer geometry.l𵇰F+H FO1xDhh(jFF1F @" F!.'0h#0"( F!1hkjHC @ IJHyDzDxD 4&HxDhh@1Error dequeueing IO output buffer.Error Locking IO output buffer.B(alloc->mHal.drvState.lod[0].stride & 0xf) == 0B𵃰F&H FxDhh/j~j,4иi9O1; F!}R!2HxDhh@## F!@ji0F]( F!_ )F2Fa FHxDhh@Error unlock output buffer.Error sending IO output buffer.Sent IO buffer with no attached surface.\Hk1pG#(F,SO\2 2#S3#(D1EYc1#B! S  2 0 ,HC P  0@-IJ HyDzDxD  +e0-C FFj'jmX F1FBF#% F)FBF#'(F1FJFb xw-OFXk- ]F(p0jP Edd Y$ Y ?hwwtDEID  B*GF\>( 0F9FRF#0F!FRF# F9FZF/DDD  E  h'h qkDE Hw FFjmFkB2ZC-C8k pF/ lToT| qB B (F1FBF hMD?DI yDCAdd code to readback from non-script memory-OkoC,m)lhGh1F8F8F!F F1FBF hw8 ( A@Xlh-O- F(𷁠j"P)"!  Al)OmD(&( (@p;F/y O 00 Os@ OG(S  ! 5hBC6hDh ^% h! 0 X Yl  ͲdUdUdUTD%AUQUPUP 0,C CN :  0 0 E p;F/O   !0 1  O9O~n(O ! BChhh ^C# 5+D& h"Dh ( 3Lp3)D<\   YDD !!! A!AEE!DE!DE!DC( :  0 0 E`)]O F0# 2 6 !6hzC D4D,DJ@  JK03  E FFjAl9Ee F3F BV pG`pGJF B К JIzDyD#B(#c JIzDyD .IB Fframeworks/rs/rsCppUtils.hP͗false && "Must be power of 2 for rounding up"0false && "Overflow of rounding operation"h F@ihh  !FGF(h!Fh(FG !` b (hl(FG h F@ih i!FGF(h!Fh(FG !` b (hl(FG pH+xDhh TjhhhTFF#GHxDhh@pjhh IyD h hF"#p@`Gmjhh  FFcFGjhhhGjhhIiGjhhiGjhhFFcFF@`GjhhFFcFGjhhFFcFF@`Gjhh FFcFF@`Gjhh$ FFcFGjhh(FF`Gjhh,FF`G F%j(hhlG(F bjhh kFG`pG"` ` pG𵁰FF !FFM F9F2F%`hi@&  &   &(F&`l0FRenderScriptRsdCpuReference::create for driver hal failed.Calloc for driver hal failed.h@ihhGF%hMhihIhG(F' `! pGpG`pGFHxDhh8F=F(F"F!jFh i)FGIyD h hHFF(F!F@rsAllocationIoSend°FH FxDhhA h@ihiG(  (FV FA*F IyD h hABError: Call to unsupported function %s in kernelFF(F!F@rsAllocationIoReceive-AFF FF F` ! CFr F1F*FrsAllocationCopy1DRange-GFF FFgFe*A28F1F*FCFS rsAllocationCopy2DRange-AFHFFFxDhh.1F:FCFF7F - (h!Fi(FG(Fi(F.IyD h h 4-AFEHFFFxDhhF /"+{+ eIfJyDzD`#+EY`1 B Ա6-)?I@yDMMINJyDzDHUIUJ yDzDC((I(yD:.%KIKJyDzD35)EIEJyDzD+5(?I?JyDzD##ع)A@F9F2F+FF hAFi FG F F*I+JyDzD y 2IyD h ht>~rs_type creation error: Cube maps require 2D typesd~rs_type creation error: mipmap control requires 2D types\~}^~t"~K2~P~~Ӳ-AFHFFFxDhh F .-IyD$I$yD  ,IyD h h F1FBF+FF-(h!Fi(FG(F+(F|rs_allocation creation error: Invalid usage flag|rs_allocation creation error: Invalid type( pGƺ FF)F"F@úp FjF"m1!3F FF(F0!F0Fp-OŰFeHFFFxDhhDbjF-XxlPW ]cl`I yDGxkhBH8xlP[ RXaoI yD1khXEF <xlP_ FLUfI yDkh@EG{Rz-Bd$O ##{F#!"cF F!!!" F!"! FF"#F!)xpIJK yDzD{D@x+O 0##' FF"#F!`x(pIJK yDzD{D@JxO ## FF"#F!(h0 0 0AppIJK yDzD{D@uw+ݍO ## FF"#F!h0 0 0 hp (pIJK yDzD{D@3BwO ## FF"#F!)h`IJK yDzD{D@v7O \##S FF"#F!lh(`IJK yDzD{D@úbvNj{O ,### FF"#F!4)h`IJK yDzD{D@uW]O ## FF"#F!h(`IJK yDzD{D@SuEO ## FF"#F!)xpIJK yDzD{D@uw'O ## FF"#F!x(pIJK yDzD{D@tO l##c FF"#F!Th0 0 0AppIJK yDzD{D@t݌O 2##) FF"#F!h0 0 0 hp (pIJK yDzD{D@_sO ## FF"#F!)h`IJK yDzD{D@'*sO ## FF"#F!h(`IJK yDzD{D@rO ## FF"#F!`)h`IJK yDzD{D@JreO h##_ FF"#F!(h(`IJK yDzD{D@q?OO 8##/ FF"#F!)IJK yDzD{D@Gjqφ3O ## FF"#F!(IJK yDzD{D@p_O ## FF"#F!)h`IJK yDzD{D@׿pO ## FF"#F!Hh(`IJK yDzD{D@pO x##o FF"#F!( IJK yDzD{D@eo NjO F##= FF"#F!( IJK yDzD{D@+2oO ##  FF"#F!( IJK yDzD{D@n#O ## FF"#F!b( IJK yDzD{D@JnsO ## FF"#F !()IJK yDzD{D@m?WO ##w FF"#F !(IJK yDzD{D@GjmςAO P##G FF"#F !)h`IJK yDzD{D@l_%O  ## FF"#F !h(`IJK yDzD{D@׽lO ## FF"#F !H( IJK yDzD{D@l{O ## FF"#F !( IJK yDzD{D@ckۊO ## FF"#F !( IJK yDzD{D@).kO Z##Q FF"#F !( IJK yDzD{D@jO (## FF"#F!`)h`IJK yDzD{D@JjO ## FF"#F!(h(`IJK yDzD{D@i?oO ## FF"#F!( IJK yDzD{D@Efi~IO ## FF"#F!( IJK yDzD{D@ hW~+O d##[ FF"#F!|(e @ IJK yDzD{D@ѻ~h}O 2##) FF"#F!B(` E IJK yDzD{D@ ho}O ## FF"#F!(e @ IJK yDzD{D@]g|ʼnO ## FF"#F!(` E IJK yDzD{D@#"g|O ## FF"#F !)h`IJK yDzD{D@f|O l##c FF"#F !\h(`IJK yDzD{D@Bf{mO <##3 FF"#F !$( IJK yDzD{D@ye3{IO  ## FF"#F !( IJK yDzD{D@?Zez-O ## FF"#F !(e @ IJK yDzD{D@dKz O ## FF"#F !v(` E IJK yDzD{D@˹rdyO t##k FF"#F !<(e @ IJK yDzD{D@ccy͈O B##9 FF"#F !(` E IJK yDzD{D@WcxO ## FF"#F!(IJK yDzD{D@c{xO ## FF"#F!(IJK yDzD{D@bxqO ## FF"#F!T(e @ IJK yDzD{D@.bwMO z##q FF"#F!(` E IJK yDzD{D@oaw1O H##? FF"cF!Hd d*@ @*IJK yDzD{D@1>avO ##  FF"cF!H` `*D D*IJK yDzD{D@`'vO ## FF"cF!dHd d*@ @*IJK yDzD{D@F`uO ## FF"cF!&H` `*D D*IJK yDzD{D@w_/uO p##g FF"#F !(IJK yDzD{D@=V_tsO >##5 FF"#F !(IJK yDzD{D@^GtWO  ## FF"#F !t(e @ IJK yDzD{D@ɾn^s5O ## FF"#F !:(` E IJK yDzD{D@]_sO ## FF"cF !Hd d*@ @*IJK yDzD{D@Q~]rO r##i FF"cF !H` `*D D*IJK yDzD{D@]grцO <##3 FF"cF !Hd d*@ @*IJK yDzD{D@ս\qO ## FF"cF !FH` `*D D*IJK yDzD{D@ \oqO ## FF"#F!)IJK yDzD{D@_[pkO ## FF"#F!(IJK yDzD{D@'*[pQO p##g FF"#F!)h`IJK yDzD{D@Zp1O @##7 FF"#F!`h(`IJK yDzD{D@JZoO ## FF"cF!(HaÀA!IJK yDzD{D@yY3oO ## FF"cF!HAÈ〢a IJK yDzD{D@;RYn˅O ## FF"cF!HaÀA!IJK yDzD{D@X;nO n##e FF"cF!nHAÈ〢a IJK yDzD{D@ZXm}O 8##/ FF"#F!0)h`IJK yDzD{D@WOm_O ## FF"#F!h(`IJK yDzD{D@OzWlGO ## FF"#F!( IJK yDzD{D@Wkl%O ## FF"#F!( IJK yDzD{D@ۺVk O t##k FF"#F!L(e @ IJK yDzD{D@VkO B##9 FF"#F!(` E IJK yDzD{D@gUkфO ## FF"#F!(e @ IJK yDzD{D@-6UjO ## FF"#F!(` E IJK yDzD{D@T'jO ## FF"#F!d( IJK yDzD{D@NTiwO z##q FF"#F!*( IJK yDzD{D@S?i]O H##? FF"#F!(e @ IJK yDzD{D@EfSh=O ##  FF"#F!(` E IJK yDzD{D@ RWh%O ## FF"cF!|Hd d*@ @*IJK yDzD{D@͸vRgO ## FF"cF!>H` `*D D*IJK yDzD{D@Q_g߃O x##o FF"cF!Hd d*@ @*IJK yDzD{D@Q~QfO B##9 FF"cF!H` `*D D*IJK yDzD{D@QgfO  ##׾ݾ %-OmF8HFF FxDhh HFF)F"F@`5@a.@b'F,F F@^IF yDL%s %f, 0x%08x*I JF   yDzDL* I JF: * cF+ yDzDVLQ*IJ F: * cF+ :yDzD;L I JF  yDzD |K~! I JF * yDzD+ ^K~! I JF : *yDzD;+ <\Kg~IF yD&K%s {%f} {0x%hx} ? IF :*0 yD;+? J%s {%f %f} {0x%hx 0x%hx} ! F+ *:!I ;+yD bJ%s {%f %f %f} {0x%hx 0x%hx 0x%hx} ! F +;*:J2 I cFK;+yD I%s {%f %f %f %f} {0x%hx 0x%hx 0x%hx 0x%hx}IF## cFyD`I%s %f, 0x%08llxa IJF yDzDM FdIS|a JFhF+I@ zD+yD -.I1|a JFhFa*I@ zD@*yD H |𵉰 F?NF? * :#F~D:1F *;+   * :0 1F#F:F: *;+   *: 1F:F#F: *;+   *: 1F#F: *;+  H%s {%f, %f, %f, %f%s %f, %f, %f, %f%s %f, %f, %f, %f}p F&NF& *#F~D*1F + L * 1F#F* + 4 * 1F#F* + pG%s {%f, %f, %f%s %f, %f, %f%s %f, %f, %f}p FNF +F~D1F     1F+F  pF%s {%f, %f%s %f, %f}IF yDHF%s %hhd 0x%hhx IF 0 yD0 ?? F%s {%hhd, %hhd} 0x%hhx 0x%hhx?F0 +ɲҲ!I yDdE%s {%hhd, %hhd, %hhd} 0x%hhx 0x%hhx 0x%hhx?F0 +;ɲҲ۲22I cFyD("E%s {%hhd, %hhd, %hhd, %hhd} 0x%hhx 0x%hhx 0x%hhx 0x%hhxIF yDD%s %hhu 0x%hhx IF 0 yD0 ?? D%s {%hhu, %hhu} 0x%hhx 0x%hhx?F0 +ɲҲ!I yDD%s {%hhu, %hhu, %hhu} 0x%hhx 0x%hhx 0x%hhx?F0 +;ɲҲ۲22I cFyDhC%s {%hhu, %hhu, %hhu, %hhu} 0x%hhx 0x%hhx 0x%hhx 0x%hhxIF yD;HC%s %hd 0x%hx?F 0 0  I yDC%s {%hd, %hd} 0x%hx 0x%hx F0 +!I yDB%s {%hd, %hd, %hd} 0x%hx 0x%hx 0x%hx F0 0+;22I cFyDBB%s {%hd, %hd, %hd, %hd} 0x%hx 0x%hx 0x%hx 0x%hxIF yDA%s %hu 0x%hx? IF 0 yD?? jA%s {%hu, %hu} 0x%hx 0x%hx F +!I yDATA%s {%hu, %hu, %hu} 0x%hx 0x%hx 0x%hx F +;22I cFyD @%s {%hu, %hu, %hu, %hu} 0x%hx 0x%hx 0x%hx 0x%hxIF yD@%s %d 0x%x IF ?yD? h@%s {%d, %d} 0x%x 0x%xa F I ?yD??? @%s {%d, %d, %d} 0x%x 0x%x 0x%xa FI?yD??? `?%s {%d, %d, %d, %d} 0x%x 0x%x 0x%x 0x%xIF yD;H?%s %u 0x%x IF ?yD? ?%s {%u, %u} 0x%x 0x%xa F I ?yD??? >%s {%u, %u, %u} 0x%x 0x%x 0x%xa FI?yD??? l>%s {%u, %u, %u, %u} 0x%x 0x%x 0x%x 0x%xFJI zDyD=qI##JF cFyDzDy=pa FIJ?yDzD??? O=pa FJa*? IzD( yD?( ??(?(?  <=pa FJa*?IzD8yD?8 ( ?(  ? ?8?8(?(? <PpF JI zDyDL<pI##JF cFyDzD"<oa FIJ?yDzD??? };oa FJa*? IzD( yD?( ??(?(? E ;oa FJa*?IzD8yD?8 ( ?(  ? ?8?8(?(? (;@o FIC0+JyDcF??zD? ?:m C1++FIJ? cF yDzD? (?(??(?(  D:mF a*IC0+JcF yD? ?zD??8?8 ( ?(8?8(?( S9:m FIC0+JyDcF??zD? ?&P9l C1++FIJ? cF yDzD? (?(??(?(  8lF a*IC0+JcF yD? ?zD??8?8 ( ?(8?8(?( v8rlIF yD7%s %p pG I"yD h h1OqOsHxDhh@uzz pGpG`pGh F@ihJi!FG( b pGpGjh hGjhhGpG`pG pGpG`pGh+ # `#P `pGl ммL#O O #`pG pG-OFHFF FxDhhE& 'B7i\@0CHF !BF] !rHFjFVFHF!    HF !BF;O@1D4 8 (2AQ|B)FKFFFO 6\A*F!#8F\@FFF  FF6oBHF !sJ( HFBF  "O@# 9AHF!R)F Raaaa#(\1?@oF& JBH҇*\ *C@HF!#^ (7ч")\ BB*(1#ABD_ p !# !#0@"HFFE6 HF!xHF! IyD h hzwpFF)  -0F)F"FO++-0F- q0F"Fr( p@^0 I K HyD{DxDHxDh F Fjxjj_Unwind_VRS_Get_InternalpFF)  ,2h(F!F ++,(Fr, #q(F/( pOJp I K HyD{DxDHxDh F F0Ajwij_Unwind_VRS_Set-CFJHFF)xDhh-)oѦ F !p'O  &(B2F#!h F .O!C6.< F( F !nP *IyD h hIر/ 3FFQ+!@h F:FC(7B. F !N( OnpIK HyD{DxD HxDh F F\vuhvhh_Unwind_VRS_PopFFFH FxDhh  FaPh  FRF ) IyD h h  uttFFοFFʿp mF%H exDhh(F )F` 0FviF0F nFcH  !F*FG((F1F( IyD h h( mp( ( ` "F#  t4t-G FF:HFFxDhh F)& 8  (jBH>O'>*i FO1eO FoBFF FIF+ H 8F)F"FG(( (  FO1%(a FHxDhh@Op IK HyD{DxDHxDh F FjsfLtrffunwind_phase2Fsm eF(F(F"F#b @!KHI{DxDyDh F F/_Unwind_Resume() can't returnsef_Unwind_Resume IyD h h F)I yD h h  &Fr,r IyD h h F)I yD h h   rqFh  GpGF !(!F4Ѐ8</ / / /ဵJzD hJk!G wFH%b`JPHP@"xD0 `EP!! Pp|QxQ@ p@ FQopFF F0hh0FG(0h)Fh0FG ` pNr`pppFF F hh FG( h)F2Fi FGh% Nreu h!%Bk FG(FppFF F0hBi0FG(0h)Fi0FG pNr`pp-AFFFF8hBi8FGH8h1F*FCFi8FG Nr`ph jG FhJj!FGah )Nk`phjGNt`phkG9  pG ))) ) JD0<0@0 )0hpG@P I K HyD{DxDHxDh F- F2cpbcgetRegister ))) ) KBdpGcpGdpG )`pG@P I K HyD{DxDHxDh F FgWcFplb/csetRegister!  pG FF$а&ѕIع IPHJ0 J``@` I K HyD{DxDHxDh F FboakbgetFloatRegister𵃰FFFF&а&їIع IPHoJ0 Jl`m`T@` I K HyD{DxDHxDh F* F/an`asetFloatRegisterF F( h!Bk FG(hh ho% FGh  FoG pPPPP F F@pG𵇰FIFFyD h hhhO1GFhX0FO1*F8 ` IyD h hRllFH FO1xDhh hh FG ahjFEHxD8(jF F)F HxDhh@Zl7l0!HpGFH$xDhhDx#F(F(H$xDhh@ FkkFAH@HnsB(@rF I yD h h1F"cFHxDhh@Rk%s6k-CF`HF%xDhh Xhbh b7!7(yаBw h(FP(h/%hF/YOCG7C>?X?GO 8C0FP?0FO O 0Fc+++HMK'{D&0K&{D/K&{DC+F' /4OCxACHB)D.JDV!X~&&%la9HxDhh @(F )@/1.H.IxDyD @60I HyDxD HxDh  F4 F9k|kxk^j\getInfoFromEHABISectionexternal/libunwind_llvm/src/UnwindCursor.hppkunknown personality routinej ]ihhB pG𵃉O !O$hwYB/Fhh7DhBmi=DBO8%N FhO h7DVh7`[iWh{`4BO -AONt UX6pG> (int4) 15; /* get XdYdZd (difference between x and point above and below) */ int4 weight2 = baseCoord & 0x7fff; /* first 15 bits */ int4 weight1 = (int4) 0x8000 - weight2; /* 1 - original */ /* point below needed point in all axis */ const int lut_offset = (coord1.x * 4) + (coord1.y * stride.y) + (coord1.z * stride.z); /* v000 = v000.lo v100 = v000.hi */ uint8 v000 = convert_uint8(vload8(0, lut + lut_offset)); /* v010 = v010.lo v110 = v010.hi */ uint8 v010 = convert_uint8(vload8(0, lut + lut_offset + stride.y)); /* v001 = v001.lo v101 = v001.hi */ uint8 v001 = convert_uint8(vload8(0, lut + lut_offset + stride.z)); /* v011 = v011.lo v111 = v011.hi */ uint8 v011 = convert_uint8(vload8(0, lut + lut_offset + stride.y + stride.z)); /* Interpolate along x */ uint4 yz00 = ((v000.lo * weight1.x) + (v000.hi * weight2.x)) >> (int4)7; uint4 yz10 = ((v010.lo * weight1.x) + (v010.hi * weight2.x)) >> (int4)7; uint4 yz01 = ((v001.lo * weight1.x) + (v001.hi * weight2.x)) >> (int4)7; uint4 yz11 = ((v011.lo * weight1.x) + (v011.hi * weight2.x)) >> (int4)7; /* Interpolate along y */ uint4 z0 = ((yz00 * weight1.y) + (yz10 * weight2.y)) >> (int4)15; uint4 z1 = ((yz01 * weight1.y) + (yz11 * weight2.y)) >> (int4)15; /* Interpolate along z */ uint4 v = ((z0 * weight1.z) + (z1 * weight2.z)) >> (int4)15; uint4 v2 = (v + 0x7f) >> (int4)8; /* part of integer math */ uchar4 ret = convert_uchar4(v2); ret.w = in.w; output[offset] = ret; } /* * Copyright: * ---------------------------------------------------------------------------- * This confidential and proprietary software may be used only as authorized * by a licensing agreement from ARM Limited. * (C) COPYRIGHT 2013 ARM Limited, ALL RIGHTS RESERVED * The entire notice above must be reproduced on all authorized copies and * copies may only be made to the extent permitted by a licensing agreement * from ARM Limited. * ---------------------------------------------------------------------------- */ __kernel void blend_clear( uint width, /* Width of the in/out allocation in uchar4. */ __global const uchar * restrict input_data, /* Pointer to the input-buffer */ __global uchar * restrict output_data /* Pointer to the output buffer */ ) { const uint col = get_global_id(0) * 16; const uint row = get_global_id(1); const uint offset = row * width * 4 + col; const uchar16 load = 0; vstore16(load, 0, output_data + offset); } __kernel void blend_src( uint width, /* Width of the in/out allocation. */ __global const uchar * input_data, /* Pointer to the input-buffer */ __global uchar * output_data /* Pointer to the output buffer */ ) { const uint col = get_global_id(0) * 16; const uint row = get_global_id(1); const uint offset = row * width * 4 + col; const uchar16 load = vload16(0, input_data + offset); vstore16(load, 0, output_data + offset); } __kernel void blend_src_over( uint width, /* Width of the in/out allocation. */ __global const uchar * input_data, /* Pointer to the input-buffer */ __global uchar * output_data /* Pointer to the output buffer */ ) { const uint col = get_global_id(0) * 16; const uint row = get_global_id(1); const uint offset = (row * width * 4 + col ); /* BLEND_SRC_OVER */ short16 in_s = convert_short16( vload16(0, input_data + offset)); const short16 out_s = convert_short16( vload16(0, output_data + offset)); short16 alpha = (short16)255 - in_s.s33337777BBBBFFFF ; in_s = in_s + ((out_s * alpha) >> (short16)8); vstore16(convert_uchar16(in_s), 0, output_data + offset); } __kernel void blend_dst_over( uint width, /* Width of the in/out allocation. */ __global const uchar * input_data, /* Pointer to the input-buffer */ __global uchar * output_data /* Pointer to the output buffer */ ) { const uint col = get_global_id(0) * 16; const uint row = get_global_id(1); const uint offset = (row * width * 4 + col ); /* BLEND_DST_OVER */ const short16 in_s = convert_short16( vload16(0, input_data + offset)); short16 out_s = convert_short16( vload16(0, output_data + offset)); short16 alpha = (short16) 255 - out_s.s33337777BBBBFFFF; out_s = out_s + ((in_s * alpha )>> (short16)8); vstore16(convert_uchar16(out_s), 0, output_data + offset); } __kernel void blend_src_in( uint width, /* Width of the in/out allocation. */ __global const uchar * input_data, /* Pointer to the input-buffer */ __global uchar * output_data /* Pointer to the output buffer */ ) { const uint col = get_global_id(0) * 16; const uint row = get_global_id(1); const uint offset = (row * width * 4 + col ); /* BLEND_SRC_IN */ short16 in_s = convert_short16(vload16(0, input_data + offset)); const short16 out_s = convert_short16(vload16(0,output_data + offset)); in_s = (in_s * out_s.s33337777BBBBFFFF) >> (short16)8; vstore16(convert_uchar16(in_s), 0, output_data + offset); } __kernel void blend_dst_in( uint width, /* Width of the in/out allocation. */ __global const uchar * input_data, /* Pointer to the input-buffer */ __global uchar * output_data /* Pointer to the output buffer */ ) { const uint col = get_global_id(0) * 16; const uint row = get_global_id(1); const uint offset = (row * width * 4 + col ); /* BLEND_DST_IN */ short16 out_s = convert_short16(vload16(0, output_data + offset)); const short16 in_s = convert_short16(vload16(0, input_data + offset)); out_s = (out_s * in_s.s33337777BBBBFFFF) >> (short16)8; vstore16(convert_uchar16(out_s), 0, output_data + offset); } __kernel void blend_src_out( uint width, /* Width of the in/out allocation. */ __global const uchar * input_data, /* Pointer to the input-buffer */ __global uchar * output_data /* Pointer to the output buffer */ ) { const uint col = get_global_id(0) * 16; const uint row = get_global_id(1); const uint offset = (row * width * 4 + col ); /* BLEND_SRC_OUT */ short16 in_s = convert_short16(vload16(0, input_data + offset)); const short16 out_s = convert_short16(vload16(0, output_data + offset)); in_s = (in_s * ((short16)255 - out_s.s33337777BBBBFFFF)) >> (short16)8; vstore16(convert_uchar16(in_s), 0, output_data + offset); } __kernel void blend_dst_out( uint width, /* Width of the in/out allocation. */ __global const uchar * input_data, /* Pointer to the input-buffer */ __global uchar * output_data /* Pointer to the output buffer */ ) { const uint col = get_global_id(0) * 16; const uint row = get_global_id(1); const uint offset = (row * width * 4 + col ); /* BLEND_SRC_OUT */ const short16 in_s = convert_short16( vload16( 0, input_data + offset)); short16 out_s = convert_short16(vload16(0, output_data + offset)); out_s = (out_s * ((short16)255 - in_s.s33337777BBBBFFFF)) >> (short16)8; vstore16(convert_uchar16(out_s), 0, output_data + offset); } __kernel void blend_src_atop( uint width, /* Width of the in/out allocation. */ __global const uchar * input_data, /* Pointer to the input-buffer */ __global uchar * output_data /* Pointer to the output buffer */ ) { const uint col = get_global_id(0) * 16; const uint row = get_global_id(1); const uint offset = row * width * 4 + col ; /* BLEND_SRC_ATOP */ const short16 in_s = convert_short16(vload16(0, input_data + offset)); short16 out_s = convert_short16(vload16(0, output_data + offset)); const short4 alphas = out_s.s37BF; const short16 dsta = out_s.s33337777BBBBFFFF; const short16 srca = in_s.s33337777BBBBFFFF; out_s = ((in_s * dsta) + (out_s * ((short16)255 - srca))) >> (short16)8; out_s.s37BF = alphas; vstore16(convert_uchar16(out_s), 0, output_data + offset); } __kernel void blend_dst_atop( uint width, /* Width of the in/out allocation. */ __global const uchar * input_data, /* Pointer to the input-buffer */ __global uchar * output_data /* Pointer to the output buffer */ ) { const uint col = get_global_id(0) * 16; const uint row = get_global_id(1); const uint offset = row * width * 4 + col ; /* BLEND_SRC_ATOP */ const short16 in_s = convert_short16(vload16(0, input_data + offset)); short16 out_s = convert_short16(vload16(0, output_data + offset)); const short4 alphas = in_s.s37BF; const short16 dsta = out_s.s33337777BBBBFFFF; const short16 srca = in_s.s33337777BBBBFFFF; out_s = ((out_s * srca) + (in_s * ((short16)255 - dsta))) >> (short16)8; out_s.s37BF = alphas; vstore16(convert_uchar16(out_s), 0, output_data + offset); } __kernel void blend_xor( uint width, /* Width of the in/out allocation in pixels. */ __global const uchar * input_data, /* Pointer to the input-buffer */ __global uchar * output_data /* Pointer to the output buffer */ ) { const uint col = get_global_id(0) * 16; const uint row = get_global_id(1); const uint offset = row * width * 4 + col; const uchar16 loadin = vload16(0, input_data + offset); uchar16 loadout = vload16(0, output_data + offset); loadout = loadout ^ loadin; vstore16(loadout, 0, output_data + offset); } __kernel void blend_multiply( uint width, /* Width of the in/out allocation. */ __global const uchar * input_data, /* Pointer to the input-buffer */ __global uchar * output_data /* Pointer to the output buffer */ ) { const uint col = get_global_id(0) * 16; const uint row = get_global_id(1); const uint offset = row * width * 4 + col; /* BLEND_MULTIPLY */ const short16 loadin = convert_short16(vload16( 0, input_data + offset)); const short16 loadout = convert_short16(vload16( 0, output_data + offset)); const uchar16 res = convert_uchar16( (loadin * loadout) >> (short16)8); vstore16(res, 0, output_data + offset); } __kernel void blend_add( uint width, /* Width of the in/out allocation. */ __global const uchar * input_data, /* Pointer to the input-buffer */ __global uchar * output_data /* Pointer to the output buffer */ ) { const uint col = get_global_id(0) * 16; const uint row = get_global_id(1); const uint offset = row * width * 4 + col; /* BLEND_ADD */ const short16 loadin = convert_short16(vload16(0, input_data + offset)); short16 loadout = convert_short16(vload16(0, output_data + offset)); loadout = loadout + loadin; const uchar16 res = convert_uchar16(min( (short16) 255, loadout)); vstore16(res, 0, output_data + offset); } __kernel void blend_subtract( uint width, /* Width of the in/out allocation. */ __global const uchar * input_data, /* Pointer to the input-buffer */ __global uchar * output_data /* Pointer to the output buffer */ ) { const uint col = get_global_id(0) * 16; const uint row = get_global_id(1); const uint offset = row * width * 4 + col; /* BLEND_SUBTRACT */ const short16 loadin = convert_short16(vload16(0, input_data + offset)); short16 loadout = convert_short16(vload16(0, output_data + offset)); loadout = loadout - loadin; const uchar16 res = convert_uchar16(max( (short16) 0, loadout)); vstore16(res, 0, output_data + offset); } /* * Copyright: * ---------------------------------------------------------------------------- * This confidential and proprietary software may be used only as authorized * by a licensing agreement from ARM Limited. * (C) COPYRIGHT 2016 ARM Limited, ALL RIGHTS RESERVED * The entire notice above must be reproduced on all authorized copies and * copies may only be made to the extent permitted by a licensing agreement * from ARM Limited. * ---------------------------------------------------------------------------- */ #pragma OPENCL EXTENSION cl_khr_fp16 : enable __constant sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_LINEAR; __kernel void gauss_blur_1d_vert(__read_only image2d_t src, __write_only image2d_t dst, const int num_coeffs, __global float* weights ) { int2 pos = {get_global_id(0), get_global_id(1)}; int i; half4 rgb_out = 0.0f; const float2 fpos = convert_float2(pos); float2 off_weight; off_weight = vload2(0, weights); for( i=0; i> 8; uchar4 res = convert_uchar4(fsum); vstore4(res, 0, dst); } ????/* Copyright: * ---------------------------------------------------------------------------- * This confidential and proprietary software may be used only as authorized * by a licensing agreement from ARM Limited. * (C) COPYRIGHT 2013 ARM Limited, ALL RIGHTS RESERVED * The entire notice above must be reproduced on all authorized copies and * copies may only be made to the extent permitted by a licensing agreement * from ARM Limited. * ---------------------------------------------------------------------------- */ /* Kernel to handle 4 horizontally adjacent pixels - clips to y access only, assumes x-1 & xmax+1 are accessible */ __kernel void conv33_block( const uint width, /* Width of the in/out allocation in pixels (uchar4) */ const uint height, /* Height of the in/out allocation in pixels (uchar4) */ __global const uchar * restrict input_data, /* Pointer to the input buffer */ __global uchar * restrict output_data, /* Pointer to the output buffer */ const float16 weights, const int stride ) { const size_t pixel_size = 4; float16 sum; float16 l, m, r; const int2 pos = (int2)( get_global_offset(0) + 4 * ( get_global_id(0) - get_global_offset(0) ), get_global_id(1) ); const int col2_offset = pos.x * pixel_size; const int col1_offset = col2_offset - pixel_size; /* Safe - we're never called with pos.x < 1 */ /* No need for col3_offset - it's within the block we're computing */ /* Row offsets incorporate offset to col1_offset (one pixel left of target block) */ const int row2_offset = col1_offset + pos.y * stride; const int row1_offset = max( row2_offset - stride, col1_offset ); const int row3_offset = col1_offset + ( min( (uint)(pos.y + 1), (uint)(height - 1) ) * stride ); l = convert_float16( vload16( 0, input_data + row1_offset ) ); r = convert_float16( vload16( 0, input_data + row1_offset + 2 * pixel_size ) ); m = (float16)(l.s456789ab, r.s456789ab); sum = weights.s0 * l; sum += weights.s1 * m; sum += weights.s2 * r; l = convert_float16( vload16( 0, input_data + row2_offset ) ); r = convert_float16( vload16( 0, input_data + row2_offset + 2 * pixel_size ) ); m = (float16)( l.s456789ab, r.s456789ab ); sum += weights.s3 * l; sum += weights.s4 * m; sum += weights.s5 * r; l = convert_float16( vload16( 0, input_data + row3_offset ) ); r = convert_float16( vload16( 0, input_data + row3_offset + 2 * pixel_size ) ); m = (float16)( l.s456789ab, r.s456789ab ); sum += weights.s6 * l; sum += weights.s7 * m; sum += weights.s8 * r; /* Must add pixel_size to row2_offset to point to col2 - first target pixel of block */ vstore16( convert_uchar16( clamp( sum, 0.f, 255.f ) ), 0, output_data + row2_offset + pixel_size ); } /* Kernel to handle left and right sections either side of middle section handled by conv33. Clips in x & y. */ __kernel void conv33( const uint width, /* Width of the in/out allocation in pixels (uchar4) */ const uint height, /* Height of the in/out allocation in pixels (uchar4) */ __global const uchar * restrict input_data, /* Pointer to the input buffer */ __global uchar * restrict output_data, /* Pointer to the output buffer */ const float16 weights, /* convolution kernel */ const int stride ) { const size_t pixel_size = 4; const int2 pos = (int2)( get_global_id(0), get_global_id(1) ); const size_t col2_offset = pos.x * pixel_size; const size_t col1_offset = max( (int)col2_offset - (int)pixel_size, (int)0 ); const size_t col3_offset = min( (uint)(pos.x + 1), (uint)(width - 1) ) * pixel_size; const size_t row2_offset = pos.y * stride; const size_t row1_offset = max( (int)row2_offset - stride, (int)0 ); const size_t row3_offset = min( (uint)pos.y + 1, (uint)height - 1 ) * stride; float4 sum = weights.s0 * convert_float4( vload4( 0, input_data + row1_offset + col1_offset ) ) + weights.s1 * convert_float4( vload4( 0, input_data + row1_offset + col2_offset ) ) + weights.s2 * convert_float4( vload4( 0, input_data + row1_offset + col3_offset ) ) + weights.s3 * convert_float4( vload4( 0, input_data + row2_offset + col1_offset ) ) + weights.s4 * convert_float4( vload4( 0, input_data + row2_offset + col2_offset ) ) + weights.s5 * convert_float4( vload4( 0, input_data + row2_offset + col3_offset ) ) + weights.s6 * convert_float4( vload4( 0, input_data + row3_offset + col1_offset ) ) + weights.s7 * convert_float4( vload4( 0, input_data + row3_offset + col2_offset ) ) + weights.s8 * convert_float4( vload4( 0, input_data + row3_offset + col3_offset ) ); vstore4( convert_uchar4( clamp( sum, 0.f, 255.f ) ), 0, output_data + row2_offset + col2_offset ); } __kernel void conv33_sharpen_block( const uint width, /* Width of the in/out allocation in pixels (uchar4) */ const uint height, /* Height of the in/out allocation in pixels (uchar4) */ __global const uchar * restrict input_data, /* Pointer to the input buffer */ __global uchar * restrict output_data, /* Pointer to the output buffer */ const int stride ) { const size_t pixel_size = 4; const int2 pos = (int2)( get_global_offset(0) + 4 * ( get_global_id(0) - get_global_offset(0) ), get_global_id(1) ); const int col2_offset = pos.x * pixel_size; const int col1_offset = col2_offset - pixel_size; /* Safe - we're never called with pos.x < 1 */ /* No need for col3_offset - it's within the block we're computing */ /* Row offsets incorporate offset to col1 (one pixel left of target block) */ const int row2_offset = col1_offset + pos.y * stride; const int row1_offset = max( row2_offset - stride, col1_offset ); const int row3_offset = col1_offset + ( min( (uint)(pos.y + 1), (uint)(height - 1) ) * stride ); /* Read data for centre row at left, middle and right offsets, synthesizing data already read */ short16 l = convert_short16( vload16( 0, input_data + row2_offset ) ); short16 r = (short16)(l.hi, convert_short8( vload8( 0, input_data + row2_offset + 4 * pixel_size ) ) ); short16 m = (short16)(l.s456789ab, r.s456789ab); short16 sum = (short)5 * m - l - r - convert_short16( vload16( 0, input_data + row1_offset + pixel_size ) ) - convert_short16( vload16( 0, input_data + row3_offset + pixel_size ) ); vstore16( convert_uchar16( clamp( sum, (short)0, (short)255 ) ), 0, output_data + row2_offset + pixel_size ); } /* Kernel for Sharpen: 0 |-1 | 0 -1| 5 |-1 0 |-1 | 0 */ __kernel void conv33_sharpen( const uint width, /* Width of the in/out allocation in pixels (uchar4) */ const uint height, /* Height of the in/out allocation in pixels (uchar4) */ __global const uchar * restrict input_data, /* Pointer to the input buffer */ __global uchar * restrict output_data, /* Pointer to the output buffer */ const int stride ) { const size_t pixel_size = 4; const int2 pos = (int2)( get_global_id(0), get_global_id(1) ); const size_t col2_offset = pos.x * pixel_size; const size_t col1_offset = max( (int)col2_offset - (int)pixel_size, (int)0 ); const size_t col3_offset = min( (uint)(pos.x + 1), (uint)(width - 1) ) * pixel_size; const size_t row2_offset = pos.y * stride; const size_t row1_offset = max( (int)row2_offset - stride, (int)0 ); const size_t row3_offset = min( (uint)pos.y + 1, (uint)height - 1 ) * stride; short4 sum = ((short)5) * convert_short4(vload4( 0, input_data + row2_offset + col2_offset )) - convert_short4(vload4( 0, input_data + row1_offset + col2_offset )) - convert_short4(vload4( 0, input_data + row2_offset + col1_offset )) - convert_short4(vload4( 0, input_data + row2_offset + col3_offset )) - convert_short4(vload4( 0, input_data + row3_offset + col2_offset )) ; vstore4( convert_uchar4( clamp( sum, (short)0, (short)255 ) ), 0, output_data + row2_offset + col2_offset ); } /* Kernel for Sobel top: 1 | 2 | 1 0 | 0 | 0 -1|-2 |-1 */ __kernel void conv33_sobel_top( const uint width, /* Width of the in/out allocation in pixels (uchar4) */ const uint height, /* Height of the in/out allocation in pixels (uchar4) */ __global const uchar * restrict input_data, /* Pointer to the input buffer */ __global uchar * restrict output_data, /* Pointer to the output buffer */ const int stride ) { const size_t pixel_size = 4; const int2 pos = (int2)( get_global_id(0), get_global_id(1) ); const size_t col2_offset = pos.x * pixel_size; const size_t col1_offset = max( (int)col2_offset - (int)pixel_size, (int)0 ); const size_t col3_offset = min( (uint)(pos.x + 1), (uint)(width - 1) ) * pixel_size; const size_t row2_offset = pos.y * stride; const size_t row1_offset = max( (int)row2_offset - stride, (int)0 ); const size_t row3_offset = min( (uint)pos.y + 1, (uint)height - 1 ) * stride; short4 sum = convert_short4( vload4( 0, input_data + row1_offset + col1_offset ) ) + (short)2 * convert_short4( vload4( 0, input_data + row1_offset + col2_offset ) ) + convert_short4( vload4( 0, input_data + row1_offset + col3_offset ) ) - convert_short4( vload4( 0, input_data + row3_offset + col1_offset ) ) - (short)2 * convert_short4( vload4( 0, input_data + row3_offset + col2_offset ) ) - convert_short4( vload4( 0, input_data + row3_offset + col3_offset ) ); vstore4( convert_uchar4( clamp( sum, (short)0, (short)255 ) ), 0, output_data + row2_offset + col2_offset ); } /* Kernel for Sobel bottom: -1|-2 |-1 0 | 0 | 0 1 | 2 | 1 */ __kernel void conv33_sobel_bottom( const uint width, /* Width of the in/out allocation in pixels (uchar4) */ const uint height, /* Height of the in/out allocation in pixels (uchar4) */ __global const uchar * restrict input_data, /* Pointer to the input buffer */ __global uchar * restrict output_data, /* Pointer to the output buffer */ const int stride ) { const size_t pixel_size = 4; const int2 pos = (int2)( get_global_id(0), get_global_id(1) ); const size_t col2_offset = pos.x * pixel_size; const size_t col1_offset = max( (int)col2_offset - (int)pixel_size, (int)0 ); const size_t col3_offset = min( (uint)(pos.x + 1), (uint)(width - 1) ) * pixel_size; const size_t row2_offset = pos.y * stride; const size_t row1_offset = max( (int)row2_offset - stride, (int)0 ); const size_t row3_offset = min( (uint)pos.y + 1, (uint)height - 1 ) * stride; short4 sum = convert_short4( vload4( 0, input_data + row3_offset + col1_offset ) ) + (short)2 * convert_short4( vload4( 0, input_data + row3_offset + col2_offset ) ) + convert_short4( vload4( 0, input_data + row3_offset + col3_offset ) ) - convert_short4( vload4( 0, input_data + row1_offset + col1_offset ) ) - (short)2 * convert_short4( vload4( 0, input_data + row1_offset + col2_offset ) ) - convert_short4( vload4( 0, input_data + row1_offset + col3_offset ) ); vstore4( convert_uchar4( clamp( sum, (short)0, (short)255 ) ), 0, output_data + row2_offset + col2_offset ); } /* Kernel for Sobel left: 1 | 0 |-1 2 | 0 |-2 1 | 0 |-1 */ __kernel void conv33_sobel_left( const uint width, /* Width of the in/out allocation in pixels (uchar4) */ const uint height, /* Height of the in/out allocation in pixels (uchar4) */ __global const uchar * restrict input_data, /* Pointer to the input buffer */ __global uchar * restrict output_data, /* Pointer to the output buffer */ const int stride ) { const size_t pixel_size = 4; const int2 pos = (int2)( get_global_id(0), get_global_id(1) ); const size_t col2_offset = pos.x * pixel_size; const size_t col1_offset = max( (int)col2_offset - (int)pixel_size, (int)0 ); const size_t col3_offset = min( (uint)(pos.x + 1), (uint)(width - 1) ) * pixel_size; const size_t row2_offset = pos.y * stride; const size_t row1_offset = max( (int)row2_offset - stride, (int)0 ); const size_t row3_offset = min( (uint)pos.y + 1, (uint)height - 1 ) * stride; short4 sum = convert_short4( vload4( 0, input_data + row1_offset + col1_offset ) ) - convert_short4( vload4( 0, input_data + row1_offset + col3_offset ) ) + (short)2 * convert_short4( vload4( 0, input_data + row2_offset + col1_offset ) ) - (short)2 * convert_short4( vload4( 0, input_data + row2_offset + col3_offset ) ) + convert_short4( vload4( 0, input_data + row3_offset + col1_offset ) ) - convert_short4( vload4( 0, input_data + row3_offset + col3_offset ) ); vstore4( convert_uchar4( clamp( sum, (short)0, (short)255 ) ), 0, output_data + row2_offset + col2_offset ); } /* Kernel for Sobel right: -1| 0 | 1 -2| 0 | 2 -1| 0 | 1 */ __kernel void conv33_sobel_right( const uint width, /* Width of the in/out allocation in pixels (uchar4) */ const uint height, /* Height of the in/out allocation in pixels (uchar4) */ __global const uchar * restrict input_data, /* Pointer to the input buffer */ __global uchar * restrict output_data, /* Pointer to the output buffer */ const int stride ) { const size_t pixel_size = 4; const int2 pos = (int2)( get_global_id(0), get_global_id(1) ); const size_t col2_offset = pos.x * pixel_size; const size_t col1_offset = max( (int)col2_offset - (int)pixel_size, (int)0 ); const size_t col3_offset = min( (uint)(pos.x + 1), (uint)(width - 1) ) * pixel_size; const size_t row2_offset = pos.y * stride; const size_t row1_offset = max( (int)row2_offset - stride, (int)0 ); const size_t row3_offset = min( (uint)pos.y + 1, (uint)height - 1 ) * stride; short4 sum = convert_short4( vload4( 0, input_data + row1_offset + col3_offset ) ) - convert_short4( vload4( 0, input_data + row1_offset + col1_offset ) ) + (short)2 * convert_short4( vload4( 0, input_data + row2_offset + col3_offset ) ) - (short)2 * convert_short4( vload4( 0, input_data + row2_offset + col1_offset ) ) + convert_short4( vload4( 0, input_data + row3_offset + col3_offset ) ) - convert_short4( vload4( 0, input_data + row3_offset + col1_offset ) ); vstore4( convert_uchar4( clamp( sum, (short)0, (short)255 ) ), 0, output_data + row2_offset + col2_offset ); } /* Kernel for Identity: 0 | 0 | 0 0 | 1 | 0 0 | 0 | 0 */ __kernel void conv33_identity( const uint width, /* Width of the in/out allocation in pixels (uchar4) */ const uint height, /* Height of the in/out allocation in pixels (uchar4) */ __global const uchar * restrict input_data, /* Pointer to the input buffer */ __global uchar * restrict output_data, /* Pointer to the output buffer */ const int stride ) { const size_t pixel_size = 4; const int2 pos = (int2)( get_global_id(0), get_global_id(1) ); const size_t col2_offset = pos.x * pixel_size; const size_t row2_offset = pos.y * stride; vstore4( vload4( 0, input_data + row2_offset + col2_offset ) , 0, output_data + row2_offset + col2_offset ); } /* Kernel for Emboss: -2|-1 | 0 -1| 1 | 1 0| 1 | 2 */ __kernel void conv33_emboss( const uint width, /* Width of the in/out allocation in pixels (uchar4) */ const uint height, /* Height of the in/out allocation in pixels (uchar4) */ __global const uchar * restrict input_data, /* Pointer to the input buffer */ __global uchar * restrict output_data, /* Pointer to the output buffer */ const int stride ) { const size_t pixel_size = 4; const int2 pos = (int2)( get_global_id(0), get_global_id(1) ); const size_t col2_offset = pos.x * pixel_size; const size_t col1_offset = max( (int)col2_offset - (int)pixel_size, (int)0 ); const size_t col3_offset = min( (uint)(pos.x + 1), (uint)(width - 1) ) * pixel_size; const size_t row2_offset = pos.y * stride; const size_t row1_offset = max( (int)row2_offset - stride, (int)0 ); const size_t row3_offset = min( (uint)pos.y + 1, (uint)height - 1 ) * stride; short4 sum = convert_short4( vload4( 0, input_data + row2_offset + col2_offset ) ) + convert_short4( vload4( 0, input_data + row2_offset + col3_offset ) ) + convert_short4( vload4( 0, input_data + row3_offset + col2_offset ) ) + (short)2 * convert_short4( vload4( 0, input_data + row3_offset + col3_offset ) ) - (short)2 * convert_short4( vload4( 0, input_data + row1_offset + col1_offset ) ) - convert_short4( vload4( 0, input_data + row1_offset + col2_offset ) ) - convert_short4( vload4( 0, input_data + row2_offset + col1_offset ) ); vstore4( convert_uchar4( clamp( sum, (short)0, (short)255 ) ), 0, output_data + row2_offset + col2_offset ); } /* Kernel for Outline: -1 |-1 |-1 -1 | 8 |-1 -1 |-1 |-1 */ __kernel void conv33_outline( const uint width, /* Width of the in/out allocation in pixels (uchar4) */ const uint height, /* Height of the in/out allocation in pixels (uchar4) */ __global const uchar * restrict input_data, /* Pointer to the input buffer */ __global uchar * restrict output_data, /* Pointer to the output buffer */ const float16 weights, /* convolution kernel */ const int stride ) { const size_t pixel_size = 4; const int2 pos = (int2)( get_global_id(0), get_global_id(1) ); const size_t col2_offset = pos.x * pixel_size; const size_t col1_offset = max( (int)col2_offset - (int)pixel_size, (int)0 ); const size_t col3_offset = min( (uint)(pos.x + 1), (uint)(width - 1) ) * pixel_size; const size_t row2_offset = pos.y * stride; const size_t row1_offset = max( (int)row2_offset - stride, (int)0 ); const size_t row3_offset = min( (uint)pos.y + 1, (uint)height - 1 ) * stride; short4 sum = (short)8 * convert_short4( vload4( 0, input_data + row2_offset + col2_offset ) ) - convert_short4( vload4( 0, input_data + row1_offset + col1_offset ) ) - convert_short4( vload4( 0, input_data + row1_offset + col2_offset ) ) - convert_short4( vload4( 0, input_data + row1_offset + col3_offset ) ) - convert_short4( vload4( 0, input_data + row2_offset + col1_offset ) ) - convert_short4( vload4( 0, input_data + row2_offset + col3_offset ) ) - convert_short4( vload4( 0, input_data + row3_offset + col1_offset ) ) - convert_short4( vload4( 0, input_data + row3_offset + col2_offset ) ) - convert_short4( vload4( 0, input_data + row3_offset + col3_offset ) ); vstore4( convert_uchar4( clamp( sum, (short)0, (short)255 ) ), 0, output_data + row2_offset + col2_offset ); } /* Copyright: * ---------------------------------------------------------------------------- * This confidential and proprietary software may be used only as authorized * by a licensing agreement from ARM Limited. * (C) COPYRIGHT 2013 ARM Limited, ALL RIGHTS RESERVED * The entire notice above must be reproduced on all authorized copies and * copies may only be made to the extent permitted by a licensing agreement * from ARM Limited. * ---------------------------------------------------------------------------- */ __kernel void conv_55( __global const uchar * restrict input_data, /* Pointer to the inputbuffer */ __global uchar * restrict output_data, /* Pointer to the output buffer */ __global const float * weights, const int stride ) { float16 sum = 0; float16 l0, l1, l2, l3, l4; const int2 pos = (int2)(get_global_id(0)*16, get_global_id(1)); int offset = (pos.y - 2 )*stride + pos.x - 8; l0 = convert_float16(vload16( 0, input_data + offset)); l4 = convert_float16(vload16( 0, input_data + offset + 16)); l2 = (float16)(l0.hi, l4.lo); l1 = (float16)(l0.s456789ab, l2.s456789ab); l3 = (float16)(l2.s456789ab, l4.s456789ab); float16 we = vload16(0, weights); sum += we.s0 * l0 ; sum += we.s1 * l1 ; sum += we.s2 * l2 ; sum += we.s3 * l3 ; sum += we.s4 * l4 ; offset += stride; l0 = convert_float16(vload16( 0, input_data + offset)); l4 = convert_float16(vload16( 0, input_data + offset + 16)); l2 = (float16)(l0.hi, l4.lo); l1 = (float16)(l0.s456789ab, l2.s456789ab); l3 = (float16)(l2.s456789ab, l4.s456789ab); sum += we.s5 * l0 ; sum += we.s6 * l1 ; sum += we.s7 * l2 ; sum += we.s8 * l3 ; sum += we.s9 * l4 ; offset += stride; l0 = convert_float16(vload16( 0, input_data + offset)); l4 = convert_float16(vload16( 0, input_data + offset + 16)); l2 = (float16)(l0.hi, l4.lo); l1 = (float16)(l0.s456789ab, l2.s456789ab); l3 = (float16)(l2.s456789ab, l4.s456789ab); sum += we.sa * l0 ; sum += we.sb * l1 ; sum += we.sc * l2 ; sum += we.sd * l3 ; sum += we.se * l4 ; offset += stride; l0 = convert_float16(vload16( 0, input_data + offset)); l4 = convert_float16(vload16( 0, input_data + offset + 16)); l2 = (float16)(l0.hi, l4.lo); l1 = (float16)(l0.s456789ab, l2.s456789ab); l3 = (float16)(l2.s456789ab, l4.s456789ab); we = vload16(0, weights + 15); sum += we.s0 * l0 ; sum += we.s1 * l1 ; sum += we.s2 * l2 ; sum += we.s3 * l3 ; sum += we.s4 * l4 ; offset += stride; l0 = convert_float16(vload16( 0, input_data + offset)); l4 = convert_float16(vload16( 0, input_data + offset + 16)); l2 = (float16)(l0.hi, l4.lo); l1 = (float16)(l0.s456789ab, l2.s456789ab); l3 = (float16)(l2.s456789ab, l4.s456789ab); sum += we.s5 * l0 ; sum += we.s6 * l1 ; sum += we.s7 * l2 ; sum += we.s8 * l3 ; sum += we.s9 * l4 ; vstore16(convert_uchar16(clamp(sum, 0.f, 255.f)), 0, output_data + pos.y * stride + pos.x); } /* * Copyright: * ---------------------------------------------------------------------------- * This confidential and proprietary software may be used only as authorized * by a licensing agreement from ARM Limited. * (C) COPYRIGHT 2017 ARM Limited, ALL RIGHTS RESERVED * The entire notice above must be reproduced on all authorized copies and * copies may only be made to the extent permitted by a licensing agreement * from ARM Limited. * ---------------------------------------------------------------------------- */ #pragma OPENCL EXTENSION cl_arm_core_id : enable #define MAX_NUM_HISTOGRAM_BUFS 16 /* Ideally this should be a power of 2. This must match kMaxNumHistogramBufs */ #define BYTES_PER_FULL_KERNEL 64 /* 32, 64 or 128 : This must match kBytesPerFullKernel */ #define VATOMIC_INC4(histogram, vals4) do { \ atomic_inc(histogram + (vals4.s0)); \ atomic_inc(histogram + (vals4.s1)); \ atomic_inc(histogram + (vals4.s2)); \ atomic_inc(histogram + (vals4.s3)); \ } while (0) #define VATOMIC_INC8(histogram, vals8) do { \ uint4 vals8_lo = convert_uint4((vals8).lo); \ VATOMIC_INC4((histogram), (vals8_lo)); \ uint4 vals8_hi = convert_uint4((vals8).hi); \ VATOMIC_INC4((histogram), (vals8_hi)); \ } while (0) #define VATOMIC_INC16(histogram, vals16) do { \ ushort8 vals16_lo = convert_ushort8((vals16).lo); \ VATOMIC_INC8((histogram), (vals16_lo)); \ ushort8 vals16_hi = convert_ushort8((vals16).hi); \ VATOMIC_INC8((histogram), (vals16_hi)); \ } while (0) __kernel void histogram_uchar_full( __global const uchar * restrict input_image, /* Pointer to the input-buffer */ __global uint * restrict histogram, /* Pointer to the histogram buffer */ uint stride /* stride/16 */ ) { histogram += ((arm_get_core_id() % MAX_NUM_HISTOGRAM_BUFS) * 256); // The offset is scaled by 16 #if BYTES_PER_FULL_KERNEL == 32 const uint2 offset = (uint2)((get_global_id(1) * stride) + (get_global_id(0) * BYTES_PER_FULL_KERNEL/16)) + (uint2)(0,1); #elif BYTES_PER_FULL_KERNEL == 64 const uint4 offset = (uint4)((get_global_id(1) * stride) + (get_global_id(0) * BYTES_PER_FULL_KERNEL/16)) + (uint4)(0,1,2,3); #elif BYTES_PER_FULL_KERNEL == 128 const uint8 offset = (uint8)((get_global_id(1) * stride) + (get_global_id(0) * BYTES_PER_FULL_KERNEL/16)) + (uint8)(0,1,2,3,4,5,6,7); #else #error "Invalid BYTES_PER_FULL_KERNEL value" #endif const uchar16 vals0 = vload16(offset.s0, input_image); const uchar16 vals1 = vload16(offset.s1, input_image); #if BYTES_PER_FULL_KERNEL >= 64 const uchar16 vals2 = vload16(offset.s2, input_image); const uchar16 vals3 = vload16(offset.s3, input_image); #endif #if BYTES_PER_FULL_KERNEL >= 128 const uchar16 vals4 = vload16(offset.s4, input_image); const uchar16 vals5 = vload16(offset.s5, input_image); const uchar16 vals6 = vload16(offset.s6, input_image); const uchar16 vals7 = vload16(offset.s7, input_image); #endif VATOMIC_INC16(histogram, vals0); VATOMIC_INC16(histogram, vals1); #if BYTES_PER_FULL_KERNEL >= 64 VATOMIC_INC16(histogram, vals2); VATOMIC_INC16(histogram, vals3); #endif #if BYTES_PER_FULL_KERNEL >= 128 VATOMIC_INC16(histogram, vals4); VATOMIC_INC16(histogram, vals5); VATOMIC_INC16(histogram, vals6); VATOMIC_INC16(histogram, vals7); #endif } __kernel void histogram_uchar_n( __global const uchar * restrict input_image, /* Pointer to the input buffer */ __global uint * restrict histogram, /* Pointer to the histogram buffer */ const uint stride, /* stride in bytes */ const uint start_offset, /* starting offset in bytes */ const uchar4 n /* number of uchar's to process in multiples - s3(16):s2(8):s1(4):s0(1) */ ) { histogram += ((arm_get_core_id() % MAX_NUM_HISTOGRAM_BUFS) * 256); input_image += (get_global_id(1) * stride) + start_offset; uchar i = n.s3; while (i-- > 0) { const uchar16 vals = vload16(0, input_image); VATOMIC_INC16(histogram, vals); input_image+=16; } i = n.s2; while (i-- > 0) { const uchar8 vals = vload8(0, input_image); VATOMIC_INC8(histogram, vals); input_image+=8; } i = n.s1; while (i-- > 0) { const uchar4 vals = vload4(0, input_image); VATOMIC_INC4(histogram, vals); input_image+=4; } i = n.s0; while (i-- > 0) { atomic_inc(histogram + *input_image++); } } #undef VATOMIC_INC4 #define VATOMIC_INC4(histogram, vals4) do { \ atomic_inc(histogram + ((vals4.s0) * 4)); \ atomic_inc(histogram + ((vals4.s1) * 4) + 1); \ atomic_inc(histogram + ((vals4.s2) * 4) + 2); \ atomic_inc(histogram + ((vals4.s3) * 4) + 3); \ } while (0) __kernel void histogram_uchar4_full( __global const uchar * restrict input_image, /* Pointer to the input buffer */ __global uint * restrict histogram, /* Pointer to the histogram buffer */ uint stride /* stride/16 */ ) { histogram += ((arm_get_core_id() % MAX_NUM_HISTOGRAM_BUFS) * 256 * 4); // The offset is scaled by 16 #if BYTES_PER_FULL_KERNEL == 32 const uint2 offset = (uint2)((get_global_id(1) * stride) + (get_global_id(0) * BYTES_PER_FULL_KERNEL/16)) + (uint2)(0,1); #elif BYTES_PER_FULL_KERNEL == 64 const uint4 offset = (uint4)((get_global_id(1) * stride) + (get_global_id(0) * BYTES_PER_FULL_KERNEL/16)) + (uint4)(0,1,2,3); #elif BYTES_PER_FULL_KERNEL == 128 const uint8 offset = (uint8)((get_global_id(1) * stride) + (get_global_id(0) * BYTES_PER_FULL_KERNEL/16)) + (uint8)(0,1,2,3,4,5,6,7); #else #error "Invalid BYTES_PER_FULL_KERNEL value" #endif const uchar16 vals0 = vload16(offset.s0, input_image); const uchar16 vals1 = vload16(offset.s1, input_image); #if BYTES_PER_FULL_KERNEL >= 64 const uchar16 vals2 = vload16(offset.s2, input_image); const uchar16 vals3 = vload16(offset.s3, input_image); #endif #if BYTES_PER_FULL_KERNEL >= 128 const uchar16 vals4 = vload16(offset.s4, input_image); const uchar16 vals5 = vload16(offset.s5, input_image); const uchar16 vals6 = vload16(offset.s6, input_image); const uchar16 vals7 = vload16(offset.s7, input_image); #endif VATOMIC_INC16(histogram, vals0); VATOMIC_INC16(histogram, vals1); #if BYTES_PER_FULL_KERNEL >= 64 VATOMIC_INC16(histogram, vals2); VATOMIC_INC16(histogram, vals3); #endif #if BYTES_PER_FULL_KERNEL >= 128 VATOMIC_INC16(histogram, vals4); VATOMIC_INC16(histogram, vals5); VATOMIC_INC16(histogram, vals6); VATOMIC_INC16(histogram, vals7); #endif } __kernel void histogram_uchar4_n( __global const uchar * restrict input_image, /* Pointer to the input buffer */ __global uint * restrict histogram, /* Pointer to the histogram buffer */ const uint stride, /* stride in bytes */ const uint start_offset, /* starting offset in bytes */ const uchar4 n /* number of uchar4's to process in multiples - s3(4):s2(2):s1(1) */ ) { histogram += ((arm_get_core_id() % MAX_NUM_HISTOGRAM_BUFS) * 256 * 4); input_image += (get_global_id(1) * stride) + start_offset; uchar i = n.s3; while (i-- > 0) { const uchar16 vals = vload16(0, input_image); VATOMIC_INC16(histogram, vals); input_image+=16; } i = n.s2; while (i-- > 0) { const uchar8 vals = vload8(0, input_image); VATOMIC_INC8(histogram, vals); input_image+=8; } i = n.s1; while (i-- > 0) { const uchar4 vals = vload4(0, input_image); VATOMIC_INC4(histogram, vals); input_image+=4; } } /* Copyright: * ---------------------------------------------------------------------------- * This confidential and proprietary software may be used only as authorized * by a licensing agreement from ARM Limited. * (C) COPYRIGHT 2013 ARM Limited, ALL RIGHTS RESERVED * The entire notice above must be reproduced on all authorized copies and * copies may only be made to the extent permitted by a licensing agreement * from ARM Limited. * ---------------------------------------------------------------------------- */ __kernel void lut( uint stride, /* Stride of the in/out allocation in uchar4. */ __global const uchar * restrict input_data, /* Pointer to the inputbuffer */ __global uchar * restrict output_data, /* Pointer to the output buffer */ __global const uchar * restrict luttables /*Look up table of values 256 for each in RGBA order */ ) { const int2 pos = (int2)(get_global_id(0)*8, get_global_id(1)); const int offset = pos.y * stride * 4 + pos.x; const uchar8 in = vload8(0, input_data + offset); uchar8 out; out.s0 = luttables[in.s0]; out.s1 = luttables[in.s1 + 256]; out.s2 = luttables[in.s2 + 512]; out.s3 = luttables[in.s3 + 768]; out.s4 = luttables[in.s4]; out.s5 = luttables[in.s5 + 256]; out.s6 = luttables[in.s6 + 512]; out.s7 = luttables[in.s7 + 768]; vstore8(out, 0, output_data + offset); } c43  5) ?  ="a"((((()))-)<)K)U)d)n)y//**++,&,H,X,l,,BBFF....}PP@@ZM]Ej $oJl ,2 #oPۧSHfoo`oxooo pe^hQ2%EGCC: (GNU) 4.9.x 20150123 (prerelease)Android clang version 5.0.300080 (based on LLVM 5.0.300080) GNUgold 1.11ACaeabi9Cortex-A53A   "$&*D.shstrtab.note.android.ident.note.gnu.build-id.dynsym.dynstr.gnu.hash.gnu.version.gnu.version_d.gnu.version_r.rel.dyn.rel.plt.text.ARM.exidx.rodata.ARM.extab.fini_array.data.rel.ro.dynamic.got.data.bss.comment.note.gnu.gold-version.ARM.attributes 44LL 2 ll0:,2,2#BoPPLo``Yoxxhow $$ B  hp &&h%LL`fVfVh|h|X0iYTp`$$p$`0$`e`p`D`