class Cuda_blockDim(AttributeTemplate):
key = types.Module(cuda.blockDim)
def resolve_x(self, mod):
return types.Macro(Cuda_blockDim_x)
def resolve_y(self, mod):
return types.Macro(Cuda_blockDim_y)
def resolve_z(self, mod):
return types.Macro(Cuda_blockDim_z)
class Cuda_threadIdx(AttributeTemplate):
key = types.Module(cuda.threadIdx)
def resolve_x(self, mod):
return types.Macro(Cuda_threadIdx_x)
def resolve_y(self, mod):
return types.Macro(Cuda_threadIdx_y)
def resolve_z(self, mod):
return types.Macro(Cuda_threadIdx_z)
class CudaAtomicTemplate(AttributeTemplate):
key = types.Module(cuda.atomic)
def resolve_add(self, mod):
return types.Function(Cuda_atomic_add)
def resolve_max(self, mod):
return types.Function(Cuda_atomic_max)
def resolve_min(self, mod):
return types.Function(Cuda_atomic_min)
def resolve_compare_and_swap(self, mod):
return types.Function(Cuda_atomic_compare_and_swap)
class HsaModuleTemplate(AttributeTemplate):
key = types.Module(roc)
def resolve_get_global_id(self, mod):
return types.Function(Hsa_get_global_id)
def resolve_get_local_id(self, mod):
return types.Function(Hsa_get_local_id)
def resolve_get_global_size(self, mod):
return types.Function(Hsa_get_global_size)
def resolve_get_local_size(self, mod):
return types.Function(Hsa_get_local_size)
def resolve_get_num_groups(self, mod):
return types.Function(Hsa_get_num_groups)
def resolve_get_work_dim(self, mod):
return types.Function(Hsa_get_work_dim)
def resolve_get_group_id(self, mod):
return types.Function(Hsa_get_group_id)
def resolve_barrier(self, mod):
return types.Function(Hsa_barrier)
def resolve_mem_fence(self, mod):
return types.Function(Hsa_mem_fence)
def resolve_wavebarrier(self, mod):
return types.Function(Hsa_wavebarrier)
def resolve_activelanepermute_wavewidth(self, mod):
return types.Function(Hsa_activelanepermute_wavewidth)
def resolve_ds_permute(self, mod):
return types.Function(Hsa_ds_permute)
def resolve_ds_bpermute(self, mod):
return types.Function(Hsa_ds_bpermute)
def resolve_shared(self, mod):
return types.Module(roc.shared)
def resolve_atomic(self, mod):
return types.Module(roc.atomic)
class CudaModuleTemplate(AttributeTemplate):
key = types.Module(cuda)
def resolve_grid(self, mod):
return types.Macro(Cuda_grid)
def resolve_gridsize(self, mod):
return types.Macro(Cuda_gridsize)
def resolve_threadIdx(self, mod):
return types.Module(cuda.threadIdx)
def resolve_blockIdx(self, mod):
return types.Module(cuda.blockIdx)
def resolve_blockDim(self, mod):
return types.Module(cuda.blockDim)
def resolve_gridDim(self, mod):
return types.Module(cuda.gridDim)
def resolve_shared(self, mod):
return types.Module(cuda.shared)
def resolve_syncthreads(self, mod):
return types.Function(Cuda_syncthreads)
def resolve_threadfence(self, mod):
return types.Function(Cuda_threadfence_device)
def resolve_threadfence_block(self, mod):
return types.Function(Cuda_threadfence_block)
def resolve_threadfence_system(self, mod):
return types.Function(Cuda_threadfence_system)
def resolve_selp(self, mod):
return types.Function(Cuda_selp)
def resolve_atomic(self, mod):
return types.Module(cuda.atomic)
def resolve_const(self, mod):
return types.Module(cuda.const)
def resolve_local(self, mod):
return types.Module(cuda.local)
class NumpyModuleAttribute(AttributeTemplate):
# note: many unary ufuncs are added later on, using setattr
key = types.Module(numpy)
def resolve_arctan2(self, mod):
return types.Function(Numpy_arctan2)
def resolve_add(self, mod):
return types.Function(Numpy_add)
def resolve_subtract(self, mod):
return types.Function(Numpy_subtract)
def resolve_multiply(self, mod):
return types.Function(Numpy_multiply)
def resolve_divide(self, mod):
return types.Function(Numpy_divide)
class OclModuleTemplate(AttributeTemplate):
key = types.Module(dppy)
def resolve_get_global_id(self, mod):
return types.Function(Ocl_get_global_id)
def resolve_get_local_id(self, mod):
return types.Function(Ocl_get_local_id)
def resolve_get_global_size(self, mod):
return types.Function(Ocl_get_global_size)
def resolve_get_local_size(self, mod):
return types.Function(Ocl_get_local_size)
def resolve_get_num_groups(self, mod):
return types.Function(Ocl_get_num_groups)
def resolve_get_work_dim(self, mod):
return types.Function(Ocl_get_work_dim)
def resolve_get_group_id(self, mod):
return types.Function(Ocl_get_group_id)
def resolve_barrier(self, mod):
return types.Function(Ocl_barrier)
def resolve_mem_fence(self, mod):
return types.Function(Ocl_mem_fence)
def resolve_sub_group_barrier(self, mod):
return types.Function(Ocl_sub_group_barrier)
def resolve_atomic(self, mod):
return types.Module(dppy.atomic)
def resolve_local(self, mod):
return types.Module(dppy.local)
def resolve_private(self, mod):
return types.Module(dppy.private)
class NumbaTypesModuleAttribute(AttributeTemplate):
key = types.Module(types)
def resolve_int8(self, mod):
return types.Function(ToInt8)
def resolve_int16(self, mod):
return types.Function(ToInt16)
def resolve_int32(self, mod):
return types.Function(ToInt32)
def resolve_int64(self, mod):
return types.Function(ToInt64)
def resolve_uint8(self, mod):
return types.Function(ToUint8)
def resolve_uint16(self, mod):
return types.Function(ToUint16)
def resolve_uint32(self, mod):
return types.Function(ToUint32)
def resolve_uint64(self, mod):
return types.Function(ToUint64)
def resolve_float32(self, mod):
return types.Function(ToFloat32)
def resolve_float64(self, mod):
return types.Function(ToFloat64)
def resolve_complex64(self, mod):
return types.Function(ToComplex64)
def resolve_complex128(self, mod):
return types.Function(ToComplex128)
def resolve_local(self, mod):
return types.Module(cuda.local)
def resolve_gridDim(self, mod):
return types.Module(cuda.gridDim)
def resolve_const(self, mod):
return types.Module(cuda.const)
def resolve_threadIdx(self, mod):
return types.Module(cuda.threadIdx)
def resolve_blockDim(self, mod):
return types.Module(cuda.blockDim)
def resolve_local(self, mod):
return types.Module(dppy.local)
class OclPrivateTemplate(AttributeTemplate):
key = types.Module(dppy.private)
def resolve_array(self, mod):
return types.Function(OCL_private_array)
def resolve_shared(self, mod):
return types.Module(hsa.shared)
class HsaSharedTemplate(AttributeTemplate):
key = types.Module(hsa.shared)
def resolve_array(self, mod):
return types.Macro(Hsa_shared_array)
return types.Function(Hsa_get_num_groups)
def resolve_get_work_dim(self, mod):
return types.Function(Hsa_get_work_dim)
def resolve_get_group_id(self, mod):
return types.Function(Hsa_get_group_id)
def resolve_barrier(self, mod):
return types.Function(Hsa_barrier)
def resolve_mem_fence(self, mod):
return types.Function(Hsa_mem_fence)
def resolve_wavebarrier(self, mod):
return types.Function(Hsa_wavebarrier)
def resolve_activelanepermute_wavewidth(self, mod):
return types.Function(Hsa_activelanepermute_wavewidth)
def resolve_shared(self, mod):
return types.Module(hsa.shared)
def resolve_atomic(self, mod):
return types.Module(hsa.atomic)
# intrinsic
intrinsic_global(hsa, types.Module(hsa))
signature(types.boolean, types.float32),
signature(types.boolean, types.float64),
]
class Math_isinf(ConcreteTemplate):
key = math.isinf
cases = [
signature(types.boolean, types.int64),
signature(types.boolean, types.uint64),
signature(types.boolean, types.float32),
signature(types.boolean, types.float64),
]
builtin_global(math, types.Module(math))
builtin_global(math.fabs, types.Function(Math_fabs))
builtin_global(math.exp, types.Function(Math_exp))
if utils.PYVERSION > (2, 6):
builtin_global(math.expm1, types.Function(Math_expm1))
builtin_global(math.sqrt, types.Function(Math_sqrt))
builtin_global(math.log, types.Function(Math_log))
builtin_global(math.log1p, types.Function(Math_log1p))
builtin_global(math.log10, types.Function(Math_log10))
builtin_global(math.sin, types.Function(Math_sin))
builtin_global(math.cos, types.Function(Math_cos))
builtin_global(math.tan, types.Function(Math_tan))
builtin_global(math.sinh, types.Function(Math_sinh))
builtin_global(math.cosh, types.Function(Math_cosh))
builtin_global(math.tanh, types.Function(Math_tanh))
builtin_global(math.asin, types.Function(Math_asin))
key = types.float32
class ToFloat64(Caster):
key = types.float64
class ToComplex64(Caster):
key = types.complex64
class ToComplex128(Caster):
key = types.complex128
builtin_global(types, types.Module(types))
builtin_global(types.int8, types.Function(ToInt8))
builtin_global(types.int16, types.Function(ToInt16))
builtin_global(types.int32, types.Function(ToInt32))
builtin_global(types.int64, types.Function(ToInt64))
builtin_global(types.uint8, types.Function(ToUint8))
builtin_global(types.uint16, types.Function(ToUint16))
builtin_global(types.uint32, types.Function(ToUint32))
builtin_global(types.uint64, types.Function(ToUint64))
builtin_global(types.float32, types.Function(ToFloat32))
builtin_global(types.float64, types.Function(ToFloat64))
builtin_global(types.complex64, types.Function(ToComplex64))
builtin_global(types.complex128, types.Function(ToComplex128))
#------------------------------------------------------------------------------
@register_attribute
class StringModuleAttr(AttributeTemplate):
key = ntypes.Module(string)
def resolve_capitalize(self, mod):
return ntypes.Function(String_capitalize)
def resolve_split(self, mod):
return ntypes.Function(String_split)
class String_capitalize(ConcreteTemplate):
key = string.capitalize
cases = [signature(StringVal, StringVal)]
class String_split(ConcreteTemplate):
key = string.split
cases = [
signature(ntypes.Array(StringVal, 1, 'C'), StringVal),
signature(ntypes.Array(StringVal, 1, 'C'), StringVal, StringVal),
signature(ntypes.Array(StringVal, 1, 'C'), StringVal, StringVal,
IntVal)
]
register_global(string, ntypes.Module(string))
register_global(string.capitalize, ntypes.Function(String_capitalize))
register_global(string.split, ntypes.Function(String_split))
def resolve_private(self, mod):
return types.Module(dppy.private)
class MathModuleAttribute(AttributeTemplate):
key = types.Module(math)
def resolve_fabs(self, mod):
return types.Function(Math_fabs)
def resolve_exp(self, mod):
return types.Function(Math_exp)
def resolve_expm1(self, mod):
return types.Function(Math_expm1)
def resolve_sqrt(self, mod):
return types.Function(Math_sqrt)
def resolve_log(self, mod):
return types.Function(Math_log)
def resolve_log1p(self, mod):
return types.Function(Math_log1p)
def resolve_log10(self, mod):
return types.Function(Math_log10)
def resolve_sin(self, mod):
return types.Function(Math_sin)
def resolve_cos(self, mod):
return types.Function(Math_cos)
def resolve_tan(self, mod):
return types.Function(Math_tan)
def resolve_sinh(self, mod):
return types.Function(Math_sinh)
def resolve_cosh(self, mod):
return types.Function(Math_cosh)
def resolve_tanh(self, mod):
return types.Function(Math_tanh)
def resolve_asin(self, mod):
return types.Function(Math_asin)
def resolve_acos(self, mod):
return types.Function(Math_acos)
def resolve_atan(self, mod):
return types.Function(Math_atan)
def resolve_atan2(self, mod):
return types.Function(Math_atan2)
def resolve_asinh(self, mod):
return types.Function(Math_asinh)
def resolve_acosh(self, mod):
return types.Function(Math_acosh)
def resolve_atanh(self, mod):
return types.Function(Math_atanh)
def resolve_pi(self, mod):
return types.float64
def resolve_e(self, mod):
return types.float64
def resolve_floor(self, mod):
return types.Function(Math_floor)
def resolve_ceil(self, mod):
return types.Function(Math_ceil)
def resolve_trunc(self, mod):
return types.Function(Math_trunc)
def resolve_isnan(self, mod):
return types.Function(Math_isnan)
def resolve_isinf(self, mod):
return types.Function(Math_isinf)
def resolve_degrees(self, mod):
return types.Function(Math_degrees)
def resolve_radians(self, mod):
return types.Function(Math_radians)
# def resolve_hypot(self, mod):
# return types.Function(Math_hypot)
def resolve_copysign(self, mod):
return types.Function(Math_copysign)
def resolve_fmod(self, mod):
return types.Function(Math_fmod)
def resolve_pow(self, mod):
return types.Function(Math_pow)
def resolve_erf(self, mod):
return types.Function(Math_erf)
def resolve_erfc(self, mod):
return types.Function(Math_erfc)
def resolve_gamma(self, mod):
return types.Function(Math_gamma)
def resolve_lgamma(self, mod):
return types.Function(Math_lgamma)
def resolve_threadfence_block(self, mod):
return types.Function(Cuda_threadfence_block)
def resolve_threadfence_system(self, mod):
return types.Function(Cuda_threadfence_system)
def resolve_selp(self, mod):
return types.Function(Cuda_selp)
def resolve_atomic(self, mod):
return types.Module(cuda.atomic)
def resolve_const(self, mod):
return types.Module(cuda.const)
def resolve_local(self, mod):
return types.Module(cuda.local)
intrinsic_global(cuda, types.Module(cuda))
## Forces the use of the cuda namespace by not recognizing individual the
## following as globals.
# intrinsic_global(cuda.grid, types.Function(Cuda_grid))
# intrinsic_global(cuda.gridsize, types.Function(Cuda_gridsize))
# intrinsic_global(cuda.threadIdx, types.Module(cuda.threadIdx))
# intrinsic_global(cuda.shared, types.Module(cuda.shared))
# intrinsic_global(cuda.shared.array, types.Function(Cuda_shared_array))
# intrinsic_global(cuda.syncthreads, types.Function(Cuda_syncthreads))
# intrinsic_global(cuda.atomic, types.Module(cuda.atomic))
class HsaAtomicTemplate(AttributeTemplate):
key = types.Module(hsa.atomic)
def resolve_add(self, mod):
return types.Function(Hsa_atomic_add)
class CudaSharedModuleTemplate(AttributeTemplate):
key = types.Module(cuda.shared)
def resolve_array(self, mod):
return types.Macro(Cuda_shared_array)
def resolve_atomic(self, mod):
return types.Module(hsa.atomic)
class CudaConstModuleTemplate(AttributeTemplate):
key = types.Module(cuda.const)
def resolve_array_like(self, mod):
return types.Macro(Cuda_const_arraylike)
def _get_global_type(self, gv):
ty = self._lookup_global(gv)
if ty is not None:
return ty
if isinstance(gv, pytypes.ModuleType):
return types.Module(gv)
class CudaLocalModuleTemplate(AttributeTemplate):
key = types.Module(cuda.local)
def resolve_array(self, mod):
return types.Macro(Cuda_local_array)
