B \W @sddlmZmZddlmZmZmZmZedZ ddZ ddZ dd Z d d Z d d ZddZddZddZddZddZddZddZddZGdddeZd d!ZGd"d#d#eZd$d%ZGd&d'd'eZd(S)))print_functionabsolute_import)typesirtypingmacrozThis is a stub.cOstdS)z OpenCL get_global_id() N) _stub_error)argskargsr .lib/python3.7/site-packages/numba/roc/stubs.py get_global_idsr cOstdS)z OpenCL get_local_id() N)r)r r r r r get_local_idsrcOstdS)z" OpenCL get_global_size() N)r)r r r r r get_global_sizesrcOstdS)z! OpenCL get_local_size() N)r)r r r r r get_local_sizesrcOstdS)z OpenCL get_group_id() N)r)r r r r r get_group_id$srcOstdS)z! OpenCL get_num_groups() N)r)r r r r r get_num_groups+srcOstdS)z OpenCL get_work_dim() N)r)r r r r r get_work_dim2srcOstdS)a0 OpenCL barrier() Example: # workgroup barrier + local memory fence hsa.barrier(hsa.CLK_LOCAL_MEM_FENCE) # workgroup barrier + global memory fence hsa.barrier(hsa.CLK_GLOBAL_MEM_FENCE) # workgroup barrier + global memory fence hsa.barrier() N)r)r r r r r barrier9srcOstdS)z OpenCL mem_fence() Example: # local memory fence hsa.mem_fence(hsa.CLK_LOCAL_MEM_FENCE) # global memory fence hsa.mem_fence(hsa.CLK_GLOBAL_MEM_FENCE) N)r)r r r r r mem_fenceJs rcCstdS)z HSAIL wavebarrier N)rr r r r wavebarrierXsrcCstdS)z- HSAIL activelanepermute_wavewidth_* N)r)srcZlaneidZidentityZ useidentityr r r activelanepermute_wavewidth_srcCstdS)zG AMDGCN Data Share intrinsic forwards permute (push semantics) N)r)src_lane dest_laner r r ds_permutefsrcCstdS)zH AMDGCN Data Share intrinsic backwards permute (pull semantics) N)r)rrr r r ds_bpermutemsrc@s(eZdZdZdZdZddZddZdS) StubzkA stub object to represent special objects which is meaningless outside the context of HSA-python. zr cCstd|dS)Nz%s is not instantiable)NotImplementedError)clsr r r __new__{sz Stub.__new__cCs|jS)N) _description_)selfr r r __repr__~sz Stub.__repr__N)__name__ __module__ __qualname____doc__r! __slots__r r#r r r r rts rcCsPt|}t|}d}t||d}t|ttj|tj}t j ||||fdS)Nzhsail.smem.allocC)r ) _legalize_shapelenrZArrayrZ signatureZUniTupleZintpZAnyrZ Intrinsic)shapedtypendimfnameZrestypeZsigr r r shared_arrays r0c@s*eZdZdZdZejdedddgdZdS) sharedzshared namespace zz shared.arrayTr,r-)callableZargnamesN) r$r%r&r'r!rZMacror0Zarrayr r r r r1s r1cCs4t|tr|St|tr|fStdt|dS)Nzinvalid type for shape; got {0}) isinstancetupleint TypeErrorformattype)r,r r r r*s   r*c@s$eZdZdZdZGdddeZdS)atomiczatomic namespace zc@seZdZdZdS)z atomic.addzCadd(ary, idx, val) Perform atomic ary[idx] += val N)r$r%r&r'r r r r addsr:N)r$r%r&r'r!rr:r r r r r9sr9N)Z __future__rrZnumbarrrrrrr rrrrrrrrrrrrobjectrr0r1r*r9r r r r s(