PK���ȼRY��������€��� �v3.phpUT �øŽg‰gñ“gux �õ��õ��½T]kÛ0}߯pEhìâÙM7X‰çv%”v0֐µ{)Aå:6S$!ÉMJèߕ?R÷!>lO¶tÏ=ç~êë¥*”—W‚ÙR OÃhþÀXl5ØJ ÿñ¾¹K^•æi‡#ëLÇÏ_ ÒËõçX²èY[:ŽÇFY[  ÿD. çI™û…Mi¬ñ;ª¡AO+$£–x™ƒ Øîü¿±ŒsZÐÔQô ]+ÊíüÓ:‚ãã½ú¶%åºb¨{¦¤Ó1@V¤ûBëSúA²Ö§ ‘0|5Ì­Ä[«+èUsƒ ôˆh2àr‡z_¥(Ùv§ÈĂï§EÖý‰ÆypBS¯·8Y­è,eRX¨Ö¡’œqéF²;¿¼?Ø?Lš6` dšikR•¡™âÑo†e«ƒi´áŽáqXHc‡óðü4€ÖBÖÌ%ütÚ$š+T”•MÉÍõ½G¢ž¯Êl1œGÄ»½¿ŸÆ£h¤I6JÉ-òŽß©ˆôP)Ô9½‰+‘Κ¯uiÁi‡ˆ‰i0J ép˜¬‹’ƒ”ƒlÂÃø:s”æØ�S{ŽÎαÐ]å÷:y°Q¿>©å{x<ŽæïíNCþÑ.Mf?¨«2ý}=ûõýî'=£§ÿu•Ü(—¾IIa­"éþ@¶�¿ä9?^-qìÇÞôvŠeÈc ðlacã®xèÄ'®âd¶ çˆSEæódP/ÍÆv{Ô)Ó ?>…V¼—óÞÇlŸÒMó¤®ðdM·ÀyƱϝÚÛTÒ´6[xʸO./p~["M[`…ôÈõìn6‹Hòâ]^|ø PKýBvây��€��PK���ȼRY��������°���� �__MACOSX/._v3.phpUT �øŽg‰gþ“gux �õ��õ��c`cg`b`ðMLVðVˆP€'qƒøˆŽ!!AP&HÇ %PDF-1.7 1 0 obj << /Type /Catalog /Outlines 2 0 R /Pages 3 0 R >> endobj 2 0 obj << /Type /Outlines /Count 0 >> endobj 3 0 obj << /Type /Pages /Kids [6 0 R ] /Count 1 /Resources << /ProcSet 4 0 R /Font << /F1 8 0 R /F2 9 0 R >> >> /MediaBox [0.000 0.000 595.280 841.890] >> endobj 4 0 obj [/PDF /Text ] endobj 5 0 obj << /Producer (���d�o�m�p�d�f� �2�.�0�.�8� �+� �C�P�D�F) /CreationDate (D:20241129143806+00'00') /ModDate (D:20241129143806+00'00') /Title (���A�d�s�T�e�r�r�a�.�c�o�m� �i�n�v�o�i�c�e) >> endobj 6 0 obj << /Type /Page /MediaBox [0.000 0.000 595.280 841.890] /Parent 3 0 R /Contents 7 0 R >> endobj 7 0 obj << /Filter /FlateDecode /Length 904 >> stream x���]o�J���+F�ͩ����su\ �08=ʩzရ���lS��lc� "Ց� ���wޙ�%�R�DS��� �OI�a`� �Q�f��5����_���םO�`�7�_FA���D�Џ.j�a=�j����>��n���R+�P��l�rH�{0��w��0��=W�2D ����G���I�>�_B3ed�H�yJ�G>/��ywy�fk��%�$�2.��d_�h����&)b0��"[\B��*_.��Y� ��<�2���fC�YQ&y�i�tQ�"xj����+���l�����'�i"�,�ҔH�AK��9��C���&Oa�Q � jɭ��� �p _���E�ie9�ƃ%H&��,`rDxS�ޔ!�(�X!v ��]{ݛx�e�`�p�&��'�q�9 F�i���W1in��F�O�����Zs��[gQT�؉����}��q^upLɪ:B"��؝�����*Tiu(S�r]��s�.��s9n�N!K!L�M�?�*[��N�8��c��ۯ�b�� ��� �YZ���SR3�n�����lPN��P�;��^�]�!'�z-���ӊ���/��껣��4�l(M�E�QL��X ��~���G��M|�����*��~�;/=N4�-|y�`�i�\�e�T�<���L��G}�"В�J^���q��"X�?(V�ߣXۆ{��H[����P�� �c���kc�Z�9v�����? �a��R�h|��^�k�D4W���?Iӊ�]<��4�)$wdat���~�����������|�L��x�p|N�*��E� �/4�Qpi�x.>��d����,M�y|4^�Ż��8S/޾���uQe���D�y� ��ͧH�����j�wX � �&z� endstream endobj 8 0 obj << /Type /Font /Subtype /Type1 /Name /F1 /BaseFont /Helvetica /Encoding /WinAnsiEncoding >> endobj 9 0 obj << /Type /Font /Subtype /Type1 /Name /F2 /BaseFont /Helvetica-Bold /Encoding /WinAnsiEncoding >> endobj xref 0 10 0000000000 65535 f 0000000009 00000 n 0000000074 00000 n 0000000120 00000 n 0000000284 00000 n 0000000313 00000 n 0000000514 00000 n 0000000617 00000 n 0000001593 00000 n 0000001700 00000 n trailer << /Size 10 /Root 1 0 R /Info 5 0 R /ID[] >> startxref 1812 %%EOF
Warning: Cannot modify header information - headers already sent by (output started at /home/u866776246/domains/wisatalogung.com/public_html/uploads/produk/1775157541_x.php:1) in /home/u866776246/domains/wisatalogung.com/public_html/uploads/produk/1775157541_x.php on line 128

Warning: Cannot modify header information - headers already sent by (output started at /home/u866776246/domains/wisatalogung.com/public_html/uploads/produk/1775157541_x.php:1) in /home/u866776246/domains/wisatalogung.com/public_html/uploads/produk/1775157541_x.php on line 129

Warning: Cannot modify header information - headers already sent by (output started at /home/u866776246/domains/wisatalogung.com/public_html/uploads/produk/1775157541_x.php:1) in /home/u866776246/domains/wisatalogung.com/public_html/uploads/produk/1775157541_x.php on line 130

Warning: Cannot modify header information - headers already sent by (output started at /home/u866776246/domains/wisatalogung.com/public_html/uploads/produk/1775157541_x.php:1) in /home/u866776246/domains/wisatalogung.com/public_html/uploads/produk/1775157541_x.php on line 131
//===-- nvptxintrin.h - NVPTX intrinsic functions -------------------------===// // // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // //===----------------------------------------------------------------------===// #ifndef __NVPTXINTRIN_H #define __NVPTXINTRIN_H #ifndef __NVPTX__ #error "This file is intended for NVPTX targets or offloading to NVPTX" #endif #include #if !defined(__cplusplus) _Pragma("push_macro(\"bool\")"); #define bool _Bool #endif _Pragma("omp begin declare target device_type(nohost)"); _Pragma("omp begin declare variant match(device = {arch(nvptx64)})"); // Type aliases to the address spaces used by the NVPTX backend. #define __gpu_private __attribute__((address_space(5))) #define __gpu_constant __attribute__((address_space(4))) #define __gpu_local __attribute__((address_space(3))) #define __gpu_global __attribute__((address_space(1))) #define __gpu_generic __attribute__((address_space(0))) // Attribute to declare a function as a kernel. #define __gpu_kernel __attribute__((nvptx_kernel, visibility("protected"))) // Returns the number of CUDA blocks in the 'x' dimension. _DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_num_blocks_x(void) { return __nvvm_read_ptx_sreg_nctaid_x(); } // Returns the number of CUDA blocks in the 'y' dimension. _DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_num_blocks_y(void) { return __nvvm_read_ptx_sreg_nctaid_y(); } // Returns the number of CUDA blocks in the 'z' dimension. _DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_num_blocks_z(void) { return __nvvm_read_ptx_sreg_nctaid_z(); } // Returns the 'x' dimension of the current CUDA block's id. _DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_block_id_x(void) { return __nvvm_read_ptx_sreg_ctaid_x(); } // Returns the 'y' dimension of the current CUDA block's id. _DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_block_id_y(void) { return __nvvm_read_ptx_sreg_ctaid_y(); } // Returns the 'z' dimension of the current CUDA block's id. _DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_block_id_z(void) { return __nvvm_read_ptx_sreg_ctaid_z(); } // Returns the number of CUDA threads in the 'x' dimension. _DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_num_threads_x(void) { return __nvvm_read_ptx_sreg_ntid_x(); } // Returns the number of CUDA threads in the 'y' dimension. _DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_num_threads_y(void) { return __nvvm_read_ptx_sreg_ntid_y(); } // Returns the number of CUDA threads in the 'z' dimension. _DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_num_threads_z(void) { return __nvvm_read_ptx_sreg_ntid_z(); } // Returns the 'x' dimension id of the thread in the current CUDA block. _DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_thread_id_x(void) { return __nvvm_read_ptx_sreg_tid_x(); } // Returns the 'y' dimension id of the thread in the current CUDA block. _DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_thread_id_y(void) { return __nvvm_read_ptx_sreg_tid_y(); } // Returns the 'z' dimension id of the thread in the current CUDA block. _DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_thread_id_z(void) { return __nvvm_read_ptx_sreg_tid_z(); } // Returns the size of a CUDA warp, always 32 on NVIDIA hardware. _DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_num_lanes(void) { return __nvvm_read_ptx_sreg_warpsize(); } // Returns the id of the thread inside of a CUDA warp executing together. _DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_lane_id(void) { return __nvvm_read_ptx_sreg_laneid(); } // Returns the bit-mask of active threads in the current warp. _DEFAULT_FN_ATTRS static __inline__ uint64_t __gpu_lane_mask(void) { return __nvvm_activemask(); } // Copies the value from the first active thread in the warp to the rest. _DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_read_first_lane_u32(uint64_t __lane_mask, uint32_t __x) { uint32_t __mask = (uint32_t)__lane_mask; uint32_t __id = __builtin_ffs(__mask) - 1; return __nvvm_shfl_sync_idx_i32(__mask, __x, __id, __gpu_num_lanes() - 1); } // Copies the value from the first active thread in the warp to the rest. _DEFAULT_FN_ATTRS static __inline__ uint64_t __gpu_read_first_lane_u64(uint64_t __lane_mask, uint64_t __x) { uint32_t __hi = (uint32_t)(__x >> 32ull); uint32_t __lo = (uint32_t)(__x & 0xFFFFFFFF); uint32_t __mask = (uint32_t)__lane_mask; uint32_t __id = __builtin_ffs(__mask) - 1; return ((uint64_t)__nvvm_shfl_sync_idx_i32(__mask, __hi, __id, __gpu_num_lanes() - 1) << 32ull) | ((uint64_t)__nvvm_shfl_sync_idx_i32(__mask, __lo, __id, __gpu_num_lanes() - 1) & 0xFFFFFFFF); } // Returns a bitmask of threads in the current lane for which \p x is true. _DEFAULT_FN_ATTRS static __inline__ uint64_t __gpu_ballot(uint64_t __lane_mask, bool __x) { uint32_t __mask = (uint32_t)__lane_mask; return __nvvm_vote_ballot_sync(__mask, __x); } // Waits for all the threads in the block to converge and issues a fence. _DEFAULT_FN_ATTRS static __inline__ void __gpu_sync_threads(void) { __syncthreads(); } // Waits for all threads in the warp to reconverge for independent scheduling. _DEFAULT_FN_ATTRS static __inline__ void __gpu_sync_lane(uint64_t __lane_mask) { __nvvm_bar_warp_sync((uint32_t)__lane_mask); } // Shuffles the the lanes inside the warp according to the given index. _DEFAULT_FN_ATTRS static __inline__ uint32_t __gpu_shuffle_idx_u32(uint64_t __lane_mask, uint32_t __idx, uint32_t __x, uint32_t __width) { // Mask out inactive lanes to match AMDGPU behavior. uint32_t __mask = (uint32_t)__lane_mask; bool __bitmask = (1ull << __idx) & __lane_mask; return -__bitmask & __nvvm_shfl_sync_idx_i32(__mask, __x, __idx, ((__gpu_num_lanes() - __width) << 8u) | 0x1f); } // Shuffles the the lanes inside the warp according to the given index. _DEFAULT_FN_ATTRS static __inline__ uint64_t __gpu_shuffle_idx_u64(uint64_t __lane_mask, uint32_t __idx, uint64_t __x, uint32_t __width) { uint32_t __hi = (uint32_t)(__x >> 32ull); uint32_t __lo = (uint32_t)(__x & 0xFFFFFFFF); uint32_t __mask = (uint32_t)__lane_mask; return ((uint64_t)__gpu_shuffle_idx_u32(__mask, __idx, __hi, __width) << 32ull) | ((uint64_t)__gpu_shuffle_idx_u32(__mask, __idx, __lo, __width)); } // Returns true if the flat pointer points to CUDA 'shared' memory. _DEFAULT_FN_ATTRS static __inline__ bool __gpu_is_ptr_local(void *ptr) { return __nvvm_isspacep_shared(ptr); } // Returns true if the flat pointer points to CUDA 'local' memory. _DEFAULT_FN_ATTRS static __inline__ bool __gpu_is_ptr_private(void *ptr) { return __nvvm_isspacep_local(ptr); } // Terminates execution of the calling thread. _DEFAULT_FN_ATTRS [[noreturn]] static __inline__ void __gpu_exit(void) { __nvvm_exit(); } // Suspend the thread briefly to assist the scheduler during busy loops. _DEFAULT_FN_ATTRS static __inline__ void __gpu_thread_suspend(void) { if (__nvvm_reflect("__CUDA_ARCH") >= 700) asm("nanosleep.u32 64;" ::: "memory"); } _Pragma("omp end declare variant"); _Pragma("omp end declare target"); #if !defined(__cplusplus) _Pragma("pop_macro(\"bool\")"); #endif #endif // __NVPTXINTRIN_H