def replace_npu_address_range_with_address(npu_addr_range):
assert isinstance(npu_addr_range.address, tvm.tir.Load)
buffer = npu_addr_range.address.buffer_var
if buffer in scratch_region_map.keys():
return vapi.NpuAddressRange(
scratch_region_map[buffer].region,
scratch_region_map[buffer].offset,
npu_addr_range.length,
)
assert buffer in buffer_addresses.keys(), f"searching for buffer : {buffer}, but not found"
address, buffer_type = buffer_addresses[buffer]
return vapi.NpuAddressRange(_get_region(buffer_type), address, npu_addr_range.length)
def _create_npu_dma_op(serial_copy):
"""This is a helper function to capture the list of arguments
to create a NpuDmaOperation object"""
src = vapi.NpuAddressRange(
# region will be updated later
region=0,
address=serial_copy.read_address,
length=int(serial_copy.length.value),
)
dest = vapi.NpuAddressRange(
# region will be updated later
region=0,
address=serial_copy.write_address,
length=int(serial_copy.length.value),
)
return vapi.NpuDmaOperation(src, dest)
def replace_npu_address_range_with_address(npu_addr_range):
assert isinstance(npu_addr_range.address, tvm.tir.BufferLoad)
buffer = npu_addr_range.address.buffer.data
index = int(
npu_addr_range.address.indices[0]
* (np.iinfo(np.dtype(npu_addr_range.address)).bits // 8)
)
if buffer in scratch_region_map.keys():
return vapi.NpuAddressRange(
scratch_region_map[buffer].region,
scratch_region_map[buffer].offset + index,
npu_addr_range.length,
)
assert buffer in buffer_addresses.keys(), f"searching for buffer : {buffer}, but not found"
address, buffer_type = buffer_addresses[buffer]
return vapi.NpuAddressRange(_get_region(buffer_type), address, npu_addr_range.length)
def replace_npu_address_range_with_address(npu_addr_range):
assert isinstance(npu_addr_range.address, tvm.tir.Load)
buffer = npu_addr_range.address.buffer_var
assert buffer in buffer_addresses.keys()
address, buffer_type = buffer_addresses[buffer]
return vapi.NpuAddressRange(_REGION_MAP[buffer_type], address,
npu_addr_range.length)
def _create_npu_dma_op(serial_copy):
"""This is a helper function to capture the list of arguments
to create a NpuDmaOperation object"""
data_type_bytes = np.iinfo(np.dtype(serial_copy.read_address.dtype)).bits // 8
src = vapi.NpuAddressRange(
# region will be updated later
region=0,
address=serial_copy.read_address,
length=int(serial_copy.length.value) * data_type_bytes,
)
dest = vapi.NpuAddressRange(
# region will be updated later
region=0,
address=serial_copy.write_address,
length=int(serial_copy.length.value) * data_type_bytes,
)
return vapi.NpuDmaOperation(src, dest)
def _create_npu_address_range(serial_address_range):
"""This is a helper function to capture a list
of arguments to create Vela NpuAddressRange object
"""
addr_range = vapi.NpuAddressRange(
# region will be updated later
region=0,
address=serial_address_range.address,
length=int(serial_address_range.length.value),
)
return addr_range
