def num_cycle_dws_conv(self, T_H, T_W, T_C, P_H, P_W, P_C):
num_cycle = [
int_ceil(T_H, P_H),
int_ceil(T_W, P_W),
int_ceil(T_C, P_C),
]
return np.prod(num_cycle)
def run_matmul(self, layer: MatmulLayer):
num_tile = [
int_ceil(layer.num_cols, self.params.T_COL),
int_ceil(layer.num_rows, self.params.T_ROW)
]
total_num_tiles = np.prod(num_tile)
return total_num_tiles * self.run_matmul_tile(layer)
def num_cycle_std_conv(self, T_H, T_W, T_C, T_F, P_H, P_W, P_C, P_F):
num_cycle = [
int_ceil(T_H, P_H),
int_ceil(T_W, P_W),
int_ceil(T_C, P_C),
int_ceil(T_F, P_F)
]
return np.prod(num_cycle)
def run_dpt_conv(self, layer: ConvLayer):
num_tile = [
int_ceil(layer.H, self.params.T_H),
int_ceil(layer.W, self.params.T_W),
int_ceil(layer.C, self.params.T_C)
]
total_num_tiles = np.prod(num_tile)
return self.run_dpt_conv_tile(layer) * total_num_tiles
def run_matmul_tile(self, layer: MatmulLayer):
T_ROW = min(self.params.T_ROW, layer.num_rows)
T_COL = min(self.params.T_COL, layer.num_cols)
num_cycle = [
int_ceil(T_ROW, self.params.P_ROW),
int_ceil(T_COL, self.params.P_COL)
]
total_num_cycles = np.prod(num_cycle)
total_data = (T_ROW * T_COL + T_COL + T_ROW)
return np.array([total_num_cycles, total_data])
def run_dpt_conv_tile(self, layer):
T_C = min(self.params.T_C, layer.C)
num_cycle = [
int_ceil(self.params.T_H, self.params.P_H),
int_ceil(self.params.T_W, self.params.P_W),
int_ceil(T_C, self.params.P_C * self.params.P_F)
]
total_num_cycles = np.prod(num_cycle)
total_num_cycles = int(math.ceil(total_num_cycles / (layer.stride**2)))
total_data = ((T_C * (self.params.K**2)) +
(self.params.T_H * self.params.T_W * T_C) * 2)
return np.array([total_num_cycles, total_data])
def run_depthwise_separable_block(self, block: DepthwiseSeparableBlock):
if not isinstance(block, DepthwiseSeparableBlock):
raise TypeError(
'block should be a DepthwiseSeparableBlock instance, got %s' %
type(block))
num_tile = [
int_ceil(block.layers[0].H, self.params.T_H),
int_ceil(block.layers[0].W, self.params.T_W),
int_ceil(block.layers[0].C, self.params.T_C),
int_ceil(block.layers[1].F, self.params.T_F)
]
total_num_tiles = np.prod(num_tile)
result = self.run_depthwise_separable_block_tile(
block) * total_num_tiles
return result
def run_separable_bottleneck_block(self, block: SeparableBottleneckBlock):
if not isinstance(block, SeparableBottleneckBlock):
raise TypeError(
'block should be a SeparableBottleneckBlock instance, got %s' %
type(block))
num_tile = [
int_ceil(block.layers[0].H, self.params.T_H),
int_ceil(block.layers[0].W, self.params.T_W),
int_ceil(block.layers[0].C, self.params.T_C1),
int_ceil(block.layers[1].C, self.params.T_C2),
int_ceil(block.layers[2].F, self.params.T_FF)
]
total_num_tiles = np.prod(num_tile)
result = self.run_separable_bottleneck_block_tile(
block) * total_num_tiles
return result
def run_bottleneck_block(self, block: BottleneckBlock):
num_tile = [
int_ceil(block.layers[0].H, self.params.T_H),
int_ceil(block.layers[0].W, self.params.T_W),
int_ceil(block.layers[0].C, self.params.T_C1),
int_ceil(block.layers[1].C, self.params.T_C2),
int_ceil(block.layers[2].C, self.params.T_C3),
int_ceil(block.layers[2].F, self.params.T_FF)
]
total_num_tiles = np.prod(num_tile)
result = self.run_bottleneck_block_tile(block) * total_num_tiles
if len(block.layers) == 4:
result += self.run_layer(block.layers[3])
return result
def get_BRAM(size): return int_ceil(size, StratixVPlatform.BRAM_BLOCK_SIZE)
def get_DSP(num_mult, bit_width): return num_mult / int_ceil(StratixVPlatform.DSP_BIT_WIDTH, bit_width)
