def __init__(self, arg_count, py_func, pysig):
self._tm = default_type_manager
# A mapping of signatures to entry points
self.overloads = utils.OrderedDict()
# A mapping of signatures to compile results
self._compileinfos = utils.OrderedDict()
self.py_func = py_func
# other parts of Numba assume the old Python 2 name for code object
self.func_code = get_code_object(py_func)
# but newer python uses a different name
self.__code__ = self.func_code
self._pysig = pysig
argnames = tuple(self._pysig.parameters)
defargs = self.py_func.__defaults__ or ()
try:
lastarg = list(self._pysig.parameters.values())[-1]
except IndexError:
has_stararg = False
else:
has_stararg = lastarg.kind == lastarg.VAR_POSITIONAL
_dispatcher.Dispatcher.__init__(self, self._tm.get_pointer(),
arg_count, self._fold_args, argnames,
defargs, has_stararg)
self.doc = py_func.__doc__
self._compile_lock = utils.NonReentrantLock()
utils.finalize(self, self._make_finalizer())
def __init__(self):
logsz = int(os.environ.get('NUMBAPRO_CUDA_LOG_SIZE', 1024))
linkerinfo = (c_char * logsz)()
linkererrors = (c_char * logsz)()
options = {
enums.CU_JIT_INFO_LOG_BUFFER: addressof(linkerinfo),
enums.CU_JIT_INFO_LOG_BUFFER_SIZE_BYTES: c_void_p(logsz),
enums.CU_JIT_ERROR_LOG_BUFFER: addressof(linkererrors),
enums.CU_JIT_ERROR_LOG_BUFFER_SIZE_BYTES: c_void_p(logsz),
enums.CU_JIT_LOG_VERBOSE: c_void_p(1),
}
raw_keys = list(options.keys()) + [enums.CU_JIT_TARGET_FROM_CUCONTEXT]
raw_values = list(options.values())
del options
option_keys = (drvapi.cu_jit_option * len(raw_keys))(*raw_keys)
option_vals = (c_void_p * len(raw_values))(*raw_values)
self.handle = handle = drvapi.cu_link_state()
driver.cuLinkCreate(len(raw_keys), option_keys, option_vals,
byref(self.handle))
utils.finalize(self, driver.cuLinkDestroy, handle)
self.linker_info_buf = linkerinfo
self.linker_errors_buf = linkererrors
self._keep_alive = [linkerinfo, linkererrors, option_keys, option_vals]
def __init__(self, arg_count, py_func, pysig, can_fallback,
exact_match_required):
self._tm = default_type_manager
# A mapping of signatures to compile results
self.overloads = collections.OrderedDict()
self.py_func = py_func
# other parts of Numba assume the old Python 2 name for code object
self.func_code = get_code_object(py_func)
# but newer python uses a different name
self.__code__ = self.func_code
argnames = tuple(pysig.parameters)
default_values = self.py_func.__defaults__ or ()
defargs = tuple(OmittedArg(val) for val in default_values)
try:
lastarg = list(pysig.parameters.values())[-1]
except IndexError:
has_stararg = False
else:
has_stararg = lastarg.kind == lastarg.VAR_POSITIONAL
_dispatcher.Dispatcher.__init__(self, self._tm.get_pointer(),
arg_count, self._fold_args, argnames,
defargs, can_fallback, has_stararg,
exact_match_required)
self.doc = py_func.__doc__
self._compiling_counter = _CompilingCounter()
utils.finalize(self, self._make_finalizer())
def __init__(self, arg_count, py_func):
self._tm = default_type_manager
#_dispatcher.Dispatcher.__init__(self, self._tm.get_pointer(), arg_count)
# A mapping of signatures to entry points
self.overloads = {}
# A mapping of signatures to compile results
self._compileinfos = {}
# A list of nopython signatures
self._npsigs = []
self.py_func = py_func
# other parts of Numba assume the old Python 2 name for code object
self.func_code = get_code_object(py_func)
# but newer python uses a different name
self.__code__ = self.func_code
self._pysig = utils.pysignature(self.py_func)
_argnames = tuple(self._pysig.parameters)
_dispatcher.Dispatcher.__init__(self, self._tm.get_pointer(),
arg_count, _argnames)
self.doc = py_func.__doc__
self._compile_lock = utils.NonReentrantLock()
utils.finalize(self, self._make_finalizer())
def __init__(self):
# Lazily load the libHLC library
bitcode_path = os.path.join(sys.prefix, 'share', 'rocmtools')
assert os.path.exists(bitcode_path) and os.path.isdir(bitcode_path)
self.bitcode_path = bitcode_path
dev_ctx = devices.get_context()
target_cpu = dev_ctx.agent.name
self.target_cpu = target_cpu
if self.hlc is None:
try:
hlc = CDLL(os.path.join(sys.prefix, 'lib', 'librocmlite.so'))
except OSError:
raise ImportError("librocmlite.so cannot be found. Please "
"install the roctools package by: "
"conda install -c numba roctools")
else:
hlc.ROC_ParseModule.restype = moduleref_ptr
hlc.ROC_ParseBitcode.restype = moduleref_ptr
hlc.ROC_ModuleEmitBRIG.restype = c_size_t
hlc.ROC_Initialize()
utils.finalize(hlc, hlc.ROC_Finalize)
hlc.ROC_SetCommandLineOption.argtypes = [
c_int,
c_void_p,
]
type(self).hlc = hlc
def _finalize_final_module(self):
"""
Make the underlying LLVM module ready to use.
"""
self._finalize_dyanmic_globals()
# Remember this on the module, for the object cache hooks
self._final_module.__library = weakref.proxy(self)
# It seems add_module() must be done only here and not before
# linking in other modules, otherwise get_pointer_to_function()
# could fail.
cleanup = self._codegen._add_module(self._final_module)
if cleanup:
utils.finalize(self, cleanup)
self._finalize_specific()
self._finalized = True
if config.DUMP_OPTIMIZED:
dump("OPTIMIZED DUMP %s" % self._name, self.get_llvm_str())
if config.DUMP_ASSEMBLY:
# CUDA backend cannot return assembly this early, so don't
# attempt to dump assembly if nothing is produced.
asm = self.get_asm_str()
if asm:
dump("ASSEMBLY %s" % self._name, self.get_asm_str())
def __init__(self):
ex = drvapi.hsa_executable_t()
hsa.hsa_executable_create(enums.HSA_PROFILE_FULL,
enums.HSA_EXECUTABLE_STATE_UNFROZEN, None,
ctypes.byref(ex))
self._id = ex
self._as_parameter_ = self._id
utils.finalize(self, hsa.hsa_executable_destroy, self._id)
def __init__(self, model=enums.HSA_MACHINE_MODEL_LARGE,
profile=enums.HSA_PROFILE_FULL,
rounding_mode=enums.HSA_DEFAULT_FLOAT_ROUNDING_MODE_DEFAULT,
options=None):
self._id = drvapi.hsa_ext_program_t()
assert options is None
hsa.hsa_ext_program_create(model, profile, rounding_mode,
options, ctypes.byref(self._id))
self._as_parameter_ = self._id
utils.finalize(self, hsa.hsa_ext_program_destroy, self._id)
def __init__(self, device, handle, finalizer=None):
self.device = device
self.handle = handle
self.external_finalizer = finalizer
self.trashing = TrashService("cuda.device%d.context%x.trash" %
(self.device.id, self.handle.value))
self.allocations = utils.UniqueDict()
self.modules = utils.UniqueDict()
utils.finalize(self, self._make_finalizer())
# For storing context specific data
self.extras = {}
def __init__(self, context, owner, pointer, size, finalizer=None):
self.context = context
self.owned = owner
self.size = size
self.host_pointer = pointer
self.is_managed = finalizer is not None
self.handle = self.host_pointer
# For buffer interface
self._buflen_ = self.size
self._bufptr_ = self.host_pointer.value
if finalizer is not None:
utils.finalize(self, finalizer)
def finalize(self):
"""
Finalize the library. After this call, nothing can be added anymore.
Finalization involves various stages of code optimization and
linking.
"""
self._raise_if_finalized()
if config.DUMP_FUNC_OPT:
dump("FUNCTION OPTIMIZED DUMP %s" % self._name,
self.get_llvm_str())
# Link libraries for shared code
for library in self._linking_libraries:
self._final_module.link_in(library._get_module_for_linking(),
preserve=True)
for library in self._codegen._libraries:
self._final_module.link_in(library._get_module_for_linking(),
preserve=True)
# Optimize the module after all dependences are linked in above,
# to allow for inlining.
self._optimize_final_module()
self._final_module.verify()
# It seems add_module() must be done only here and not before
# linking in other modules, otherwise get_pointer_to_function()
# could fail.
cleanup = self._codegen._add_module(self._final_module)
if cleanup:
utils.finalize(self, cleanup)
self._finalize_specific()
self._finalized = True
if config.DUMP_OPTIMIZED:
dump("OPTIMIZED DUMP %s" % self._name, self.get_llvm_str())
if config.DUMP_ASSEMBLY:
# CUDA backend cannot return assembly this early, so don't
# attempt to dump assembly if nothing is produced.
asm = self.get_asm_str()
if asm:
dump("ASSEMBLY %s" % self._name, self.get_asm_str())
def __init__(self, memptr, view=None):
self._mem = memptr
if view is None:
self._view = self._mem
else:
assert not view.is_managed
self._view = view
mem = self._mem
def deref():
mem.refct -= 1
assert mem.refct >= 0
if mem.refct == 0:
mem.free()
self._mem.refct += 1
utils.finalize(self, deref)
def __init__(self, arg_count, py_func):
self.tm = default_type_manager
_dispatcher.Dispatcher.__init__(self, self.tm.get_pointer(), arg_count)
# A mapping of signatures to entry points
self.overloads = {}
# A mapping of signatures to types.Function objects
self._function_types = {}
# A mapping of signatures to compile results
self._compileinfos = {}
self.py_func = py_func
# other parts of Numba assume the old Python 2 name for code object
self.func_code = get_code_object(py_func)
# but newer python uses a different name
self.__code__ = self.func_code
self.doc = py_func.__doc__
self._compiling = False
utils.finalize(self, self._make_finalizer())
def __init__(self, context, pointer, size, finalizer=None, owner=None):
self.context = context
self.device_pointer = pointer
self.size = size
self._cuda_memsize_ = size
self.is_managed = finalizer is not None
self.refct = 0
self.handle = self.device_pointer
self._owner = owner
if finalizer is not None:
self._finalizer = utils.finalize(self, finalizer)
def __init__(self):
# Lazily load the libHLC library
if self.hlc is None:
try:
hlc = CDLL(os.path.join(sys.prefix, 'lib', 'libHLC.so'))
except OSError:
raise ImportError(
"libHLC.so cannot be found. Please install the libhlc "
"package by: conda install -c numba libhlc")
else:
hlc.HLC_ParseModule.restype = moduleref_ptr
hlc.HLC_ModuleEmitBRIG.restype = c_size_t
hlc.HLC_Initialize()
utils.finalize(hlc, hlc.HLC_Finalize)
hlc.HLC_SetCommandLineOption.argtypes = [
c_int,
c_void_p,
]
type(self).hlc = hlc
def __init__(self, devnum):
got_devnum = c_int()
driver.cuDeviceGet(byref(got_devnum), devnum)
assert devnum == got_devnum.value, "Driver returned another device"
self.id = got_devnum.value
self.trashing = trashing = TrashService("cuda.device%d.trash" % self.id)
self.attributes = {}
# Read compute capability
cc_major = c_int()
cc_minor = c_int()
driver.cuDeviceComputeCapability(byref(cc_major), byref(cc_minor),
self.id)
self.compute_capability = (cc_major.value, cc_minor.value)
# Read name
bufsz = 128
buf = (c_char * bufsz)()
driver.cuDeviceGetName(buf, bufsz, self.id)
self.name = buf.value
def finalizer():
trashing.clear()
utils.finalize(self, finalizer)
def __init__(self, context, handle, finalizer=None):
self.context = context
self.handle = handle
if finalizer is not None:
utils.finalize(self, finalizer)
def __init__(self, code_object):
self._id = code_object
self._as_parameter_ = self._id
utils.finalize(self, hsa.hsa_code_object_destroy, self._id)
def __init__(self, signal_id):
self._id = signal_id
self._as_parameter_ = self._id
utils.finalize(self, hsa.hsa_signal_destroy, self._id)
pass
moduleref_ptr = POINTER(OpaqueModuleRef)
try:
hlc = CDLL(os.path.join(sys.prefix, 'lib', 'libHLC.so'))
except OSError:
raise ImportError("libHLC.so cannot be found. Please install the libhlc "
"package by: conda install -c numba libhlc")
else:
hlc.HLC_ParseModule.restype = moduleref_ptr
hlc.HLC_ModuleEmitBRIG.restype = c_size_t
hlc.HLC_Initialize()
utils.finalize(hlc, hlc.HLC_Finalize)
hlc.HLC_SetCommandLineOption.argtypes = [
c_int,
c_void_p,
]
def set_option(*opt):
"""
Use this for setting debug flags to libHLC using the same options
available to LLVM.
E.g -debug-pass=Structure
"""
inp = [
create_string_buffer(x.encode('ascii')) for x in (('libhlc', ) + opt)
def __init__(self, context, handle, info_log, finalizer=None):
self.context = context
self.handle = handle
self.info_log = info_log
if finalizer is not None:
self._finalizer = utils.finalize(self, finalizer)