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