def sobel(A, Gx, Gy):
B = hcl.compute((height, width),
lambda x, y: A[x][y][0] + A[x][y][1] + A[x][y][2],
"B",
dtype=hcl.Float())
r = hcl.reduce_axis(0, 3)
c = hcl.reduce_axis(0, 3)
D = hcl.compute(
(height, width),
lambda x, y: hcl.select(
hcl.and_(x > 0, x < (height - 1), y > 0, y < (width - 1)),
hcl.sum(B[x + r, y + c] * Gx[r, c], axis=[r, c]), B[x, y]),
"D",
dtype=hcl.Float())
t = hcl.reduce_axis(0, 3)
g = hcl.reduce_axis(0, 3)
E = hcl.compute(
(height, width),
lambda x, y: hcl.select(
hcl.and_(x > 0, x < (height - 1), y > 0, y < (width - 1)),
hcl.sum(B[x + t, y + g] * Gy[t, g], axis=[t, g]), B[x, y]),
"E",
dtype=hcl.Float())
return hcl.compute((height, width),
lambda x, y: hcl.sqrt(D[x][y] * D[x][y] + E[x][y] * E[x]
[y]) / 4328 * 255,
dtype=hcl.Float())
def learn(k, hdTrainData, prototype, prototypeCounter):
#Find samples that have the label k
match = hcl.compute(
hdTrainData.shape,
lambda x, y: hcl.select(trainLabels[x] == k, hdTrainData[x][y], 0),
"match")
#Record the number of these samples
with hcl.for_(0, hdTrainData.shape[0]) as a:
with hcl.if_(trainLabels[a] == k):
max[k] += 1
#Do hdc sum on these samples' hdv
r = hcl.reduce_axis(0, hdTrainData.shape[0], 'r')
result = hcl.compute((hdTrainData.shape[1], ),
lambda y: hcl.sum(match[r][y], axis=r), "result")
#Do the binary voting
sum1 = hcl.compute((hdTrainData.shape[1], ), lambda x: 0, "sum1")
with hcl.if_(max[k] % 2 == 0):
hcl.update(
sum1, lambda x: hcl.select(
result[x] + rdv3[k][x] - max[k] / 2 > 0, 1, 0))
with hcl.else_():
hcl.update(sum1,
lambda x: hcl.select(result[x] - max[k] / 2 > 0, 1, 0))
#Push the binary sum to prototype and the original sum to prototypeCounter
with hcl.for_(0, hdTrainData.shape[1]) as t:
prototype[k][t] = sum1[t]
prototypeCounter[k][t] = result[t]
def update(l, prototype, prototypeCounter, max):
hcl.print((l + 1),
"%d:Use hard examples to update the prototype counters.\n")
###data preparation
distance = hcl.compute((hdTrainData.shape[1], ), lambda x: 0,
'distance')
hamming_dist = hcl.compute((numClasses, ), lambda x: 0, "hamming_dist")
m = hcl.reduce_axis(0, hdTrainData.shape[1], "m")
###
with hcl.for_(0, hdTrainData.shape[0]) as i:
with hcl.for_(0, numClasses) as n:
#Do hdc multiplication(XOR) on sample[i]'s hdv and prototype[n]'s hdv (elementwise on the high-bit data)
hcl.update(distance,
lambda x: hdTrainData[i][x] ^ prototype[n][x])
#Calculate the hamming distance of the two vectors by adding 1s
hamming_dist[n] = hcl.sum(distance[m], axis=m)
#Find the one having the least hamming distance and choose it's label as the predicted label
pred = hcl.scalar(0, 'pred')
with hcl.for_(0, hamming_dist.shape[0]) as j:
with hcl.if_(hamming_dist[j] < hamming_dist[pred]):
pred.v = j
#Adjust the proto vectors by adding the sample vector on its label proto hdv and substrct it on its predicted proto hdv
with hcl.if_(pred.v != trainLabels[i]):
max[trainLabels[i]] += 1
max[pred] -= 1
with hcl.for_(0, hdTrainData.shape[1]) as m:
prototypeCounter[trainLabels[i]][m] += hdTrainData[i][m]
prototypeCounter[pred][m] -= hdTrainData[i][m]
with hcl.if_(max[trainLabels[i]] % 2 == 0):
with hcl.if_(prototypeCounter[trainLabels[i]][m] -
max[trainLabels[i]] / 2 == 0):
prototype[trainLabels[i]][m] &= 1
with hcl.else_():
prototype[trainLabels[i]][m] = hcl.select(
prototypeCounter[trainLabels[i]][m] -
max[trainLabels[i]] / 2 > 0, 1, 0)
with hcl.if_(max[pred] % 2 == 0):
with hcl.if_(prototypeCounter[pred][m] -
max[pred] / 2 == 0):
prototype[pred][m] &= 1
with hcl.else_():
prototype[pred][m] = hcl.select(
prototypeCounter[pred][m] - max[pred] / 2 > 0, 1,
0)
#print the accuracy
hcl.mutate(
(1, ),
lambda x: test_hdc_accu(prototype, hdTrainData, trainLabels, 1),
'training_update')
hcl.mutate(
(1, ),
lambda x: test_hdc_accu(prototype, hdTestData, testLabels, 2),
'testing_update')
def sobelAlgo(A, Fx, Fy): B = hcl.compute((height, width), lambda x,y :A[x][y][0]+A[x][y][1]+A[x][y][2],"B", dtype=hcl.Float()) r = hcl.reduce_axis(0, 3) c = hcl.reduce_axis(0, 3) Gx = hcl.compute((height, width), lambda x,y: hcl.select(hcl.and_(x>0,x<(height-1),y>0,y<(width-1)), hcl.sum(B[x+r,y+c]*Fx[r,c],axis=[r,c]), B[x,y]), "Gx") t = hcl.reduce_axis(0, 3) g = hcl.reduce_axis(0, 3) Gy = hcl.compute((height, width), lambda x,y: hcl.select(hcl.and_(x>0,x<(height-1),y>0,y<(width-1)), hcl.sum(B[x+t,y+g]*Fy[t,g],axis=[t,g]), B[x,y]), "Gy") return hcl.compute((height, width), lambda x,y:(hcl.sqrt(Gx[x][y]*Gx[x][y]+Gy[x][y]*Gy[x][y]))/4328*255, dtype = hcl.Float())
def clip(x, a_min=0.0, a_max=1.0, name="clip"):
lower = hcl.compute(
x.shape,
lambda *args: hcl.select(x[args] <= a_max, x[args], a_max),
name=name + "_low")
return hcl.compute(
x.shape,
lambda *args: hcl.select(lower[args] >= a_min, lower[args], a_min),
name=name + "_high")
def kernel_select(a, b, c, d):
use_imm = hcl.scalar(1)
with hcl.for_(0, 10, name="i") as i:
src = hcl.select(use_imm == 1, hcl.cast(hcl.Int(16),
(c[i] + b[i])),
hcl.cast(hcl.Int(32), (c[i] - b[i])))
dst = hcl.cast(hcl.Int(32), (2 * (c[i] + b[i])))
d[i] = hcl.select(dst >= (-1 * src),
hcl.select(dst <= src, a[i], src), (-1 * src))
def update(l, prototype, prototypeCounter, max):
hcl.print((l+1),"%d:Use hard examples to update the prototype counters.\n")
###data preparation
distance = hcl.compute((in_train.shape[1],), lambda x: 0, 'distance', dtype=hcl.UInt(in_bw))
pre_dist = hcl.compute((in_train.shape[1],), lambda x: 0, "pre_dist")
hamming_dist = hcl.compute((numClasses,), lambda x: 0, "hamming_dist")
m = hcl.reduce_axis(0, in_train.shape[1], "m")
###
with hcl.for_(0, in_train.shape[0]) as i:
hcl.print((i),"%d suc\n")
# pack_proto = hcl.pack(prototype, axis=1, dtype=hcl.UInt(in_bw), name="pack_proto")
with hcl.for_(0, numClasses) as n:
#Do hdc multiplication(XOR) on sample[i]'s hdv and prototype[n]'s hdv (elementwise on the high-bit data)
hcl.update(distance, lambda x: in_train[i][x] ^ prototype[n][x])
#Calculate the hamming distance of the two vectors by adding 1s
hcl.update(pre_dist, lambda x: popcount(distance[x]))
hcl.print((),"sum of 1s suc")
hamming_dist[n] = hcl.sum(pre_dist[m], axis=m)
#Find the one having the least hamming distance and choose it's label as the predicted label
pred = hcl.scalar(0, 'pred')
with hcl.for_(0, hamming_dist.shape[0]) as j:
with hcl.if_(hamming_dist[j] < hamming_dist[pred]):
pred.v = j
#Adjust the proto vectors by adding the sample vector on its label proto hdv and substrct it on its predicted proto hdv
with hcl.if_(pred.v != trainLabels[i]):
max[trainLabels[i]] += 1
max[pred] -= 1
with hcl.for_(0, in_train.shape[1]) as m:
with hcl.for_(0, in_bw) as bit:
# with hcl.if_(in_train[i][m][bit] == 1):
# ###########
# prototypeCounter[trainLabels[i]][m*in_bw+bit] += 1
# prototypeCounter[pred][m*in_bw+bit] -= 1
prototypeCounter[trainLabels[i]][m*in_bw+bit] += in_train[i][m][bit]
prototypeCounter[pred][m*in_bw+bit] -= in_train[i][m][bit]
with hcl.if_(max[trainLabels[i]] % 2 == 0):
with hcl.if_(prototypeCounter[trainLabels[i]][m*in_bw+bit] - max[trainLabels[i]]/2 == 0):
prototype[trainLabels[i]][m][bit] &= 1
with hcl.else_():
prototype[trainLabels[i]][m][bit] = hcl.select(prototypeCounter[trainLabels[i]][m*in_bw+bit] - max[trainLabels[i]]/2 > 0, 1, 0)
with hcl.if_(max[pred] % 2 == 0):
with hcl.if_(prototypeCounter[pred][m*in_bw+bit] - max[pred]/2 == 0):
prototype[pred][m][bit] &= 1
with hcl.else_():
prototype[pred][m][bit] = hcl.select(prototypeCounter[pred][m*in_bw+bit] - max[pred]/2 > 0, 1, 0)
#print the accuracy
hcl.mutate((1,), lambda x: test_hdc_accu(prototype, in_train, trainLabels, 1), 'training_update')
hcl.mutate((1,), lambda x: test_hdc_accu(prototype, in_test, testLabels, 2), 'testing_update')
def broadcast_greater_equal(input1, input2, name='broadcast_greater_equal'):
input1_mod, input2_mod, switch = broadcast_set(input1, input2)
if (switch):
return hcl.compute(input1_mod.shape,
lambda *x: hcl.select(
input1_mod[x] >= input2_mod[_broadcast(
input2_mod.shape, x)], 1, 0),
name=name)
else:
return hcl.compute(input2_mod.shape,
lambda *x: hcl.select(
input1_mod[_broadcast(input1_mod.shape, x)] >=
input2_mod[x], 1, 0),
name=name)
def broadcast_left_shift(input1, input2, name='broadcast_left_shift'):
input1_mod, input2_mod, switch = broadcast_set(input1, input2)
if (switch):
return hcl.compute(input1_mod.shape,
lambda *x: hcl.select(
input1_mod[x] >> input2_mod[_broadcast(
input2_mod.shape, x)], 1, 0),
name=name)
else:
return hcl.compute(input1_mod.shape,
lambda *x: hcl.select(
input1_mod[_broadcast(input1_mod.shape, x)] >>
input2_mod[x], 1, 0),
name=name)
def top(input, filter, bias, ):
input_extent_3_required_s = (((((final_extent_2 + 31)//32) * final_extent_3) + -1)//hcl.select(((final_extent_2 + 31)//32) > 1, ((final_extent_2 + 31)//32), 1))
final_total_extent_1 = (hcl.cast(dtype = hcl.Int(bits = 64), expr = final_extent_1) * hcl.cast(dtype = hcl.Int(bits = 64), expr = final_extent_0))
final_total_extent_2 = (final_total_extent_1 * hcl.cast(dtype = hcl.Int(bits = 64), expr = final_extent_2))
final_total_extent_3 = (final_total_extent_2 * hcl.cast(dtype = hcl.Int(bits = 64), expr = final_extent_3))
f_conv_n_extent_realized_s = hcl.select(hcl.select((((final_extent_2 * final_extent_3) + -1)//hcl.select(final_extent_2 > 1, final_extent_2, 1)) > (final_extent_3 + -1), (((final_extent_2 * final_extent_3) + -1)//hcl.select(final_extent_2 > 1, final_extent_2, 1)), (final_extent_3 + -1)) > (((((final_extent_2 + 31)//32) * final_extent_3) + -1)//(hcl.select(((final_extent_2 + -1)//32) > 0, ((final_extent_2 + -1)//32), 0) + 1)), hcl.select((((final_extent_2 * final_extent_3) + -1)//hcl.select(final_extent_2 > 1, final_extent_2, 1)) > (final_extent_3 + -1), (((final_extent_2 * final_extent_3) + -1)//hcl.select(final_extent_2 > 1, final_extent_2, 1)), (final_extent_3 + -1)), (((((final_extent_2 + 31)//32) * final_extent_3) + -1)//(hcl.select(((final_extent_2 + -1)//32) > 0, ((final_extent_2 + -1)//32), 0) + 1)))
f_conv_z_extent_realized = hcl.select(((hcl.select(((final_extent_2 + -1)//32) > 0, ((final_extent_2 + -1)//32), 0) * 32) + 32) > final_extent_2, ((hcl.select(((final_extent_2 + -1)//32) > 0, ((final_extent_2 + -1)//32), 0) * 32) + 32), final_extent_2)
f_conv = hcl.compute((final_extent_0, ((((final_extent_1 + -1)//32) * 32) + 32), f_conv_z_extent_realized, (f_conv_n_extent_realized_s + 1)), lambda x, y, z, w: 0, name = "f_conv", dtype = hcl.Float(bits = 32))
with hcl.Stage("f_conv"):
with hcl.for_(0, (final_extent_2 * final_extent_3), name = "f_conv_s0_z_par") as f_conv_s0_z_par:
with hcl.for_(final_min_1, final_extent_1, name = "f_conv_s0_y") as f_conv_s0_y:
with hcl.for_(final_min_0, final_extent_0, name = "f_conv_s0_x") as f_conv_s0_x:
f_conv[f_conv_s0_x, f_conv_s0_y, ((f_conv_s0_z_par % hcl.select(final_extent_2 > 1, final_extent_2, 1)) + final_min_2), ((f_conv_s0_z_par//hcl.select(final_extent_2 > 1, final_extent_2, 1)) + final_min_3)] = bias[((f_conv_s0_z_par % hcl.select(final_extent_2 > 1, final_extent_2, 1)) + final_min_2)]
with hcl.for_(0, (((final_extent_2 + 31)//32) * final_extent_3), name = "f_conv_s1_z_z_par") as f_conv_s1_z_z_par:
f_conv_s1_z_z_t_base_s = (f_conv_s1_z_z_par % hcl.select(((final_extent_2 + 31)//32) > 1, ((final_extent_2 + 31)//32), 1))
with hcl.for_(0, 32, name = "f_conv_s1_r__z") as f_conv_s1_r__z:
with hcl.for_(0, ((final_extent_1 + 31)//32), name = "f_conv_s1_y_y") as f_conv_s1_y_y:
with hcl.for_(0, 32, name = "f_conv_s1_z_z_t") as f_conv_s1_z_z_t:
with hcl.for_(0, 32, name = "f_conv_s1_y_y_t") as f_conv_s1_y_y_t:
with hcl.for_(final_min_0, final_extent_0, name = "f_conv_s1_x") as f_conv_s1_x:
with hcl.for_(0, 3, name = "f_conv_s1_r__y_r21") as f_conv_s1_r__y_r21:
with hcl.for_(0, 3, name = "f_conv_s1_r__x_r20") as f_conv_s1_r__x_r20:
t51_s = (f_conv_s1_z_z_par//hcl.select(((final_extent_2 + 31)//32) > 1, ((final_extent_2 + 31)//32), 1))
f_conv[f_conv_s1_x, (((f_conv_s1_y_y * 32) + final_min_1) + f_conv_s1_y_y_t), (((f_conv_s1_z_z_t_base_s * 32) + final_min_2) + f_conv_s1_z_z_t), ((f_conv_s1_z_z_par//hcl.select(((final_extent_2 + 31)//32) > 1, ((final_extent_2 + 31)//32), 1)) + final_min_3)] = (f_conv[f_conv_s1_x, (((f_conv_s1_y_y * 32) + final_min_1) + f_conv_s1_y_y_t), (((f_conv_s1_z_z_t_base_s * 32) + final_min_2) + f_conv_s1_z_z_t), (final_min_3 + t51_s)] + (filter[f_conv_s1_r__x_r20, f_conv_s1_r__y_r21, f_conv_s1_r__z, (((f_conv_s1_z_z_t_base_s * 32) + final_min_2) + f_conv_s1_z_z_t)] * input[(f_conv_s1_r__x_r20 + f_conv_s1_x), ((((f_conv_s1_y_y * 32) + final_min_1) + f_conv_s1_y_y_t) + f_conv_s1_r__y_r21), f_conv_s1_r__z, (final_min_3 + t51_s)]))
final = hcl.compute((64, 64, 32, 4), lambda x, y, z, w: 0, name = "final", dtype = hcl.Float(bits = 32))
with hcl.Stage("final"):
with hcl.for_(final_min_3, final_extent_3, name = "final_s0_n") as final_s0_n:
with hcl.for_(final_min_2, final_extent_2, name = "final_s0_z") as final_s0_z:
with hcl.for_(final_min_1, final_extent_1, name = "final_s0_y") as final_s0_y:
with hcl.for_(final_min_0, final_extent_0, name = "final_s0_x") as final_s0_x:
final[final_s0_x, final_s0_y, final_s0_z, final_s0_n] = hcl.select(f_conv[final_s0_x, final_s0_y, final_s0_z, final_s0_n] > hcl.cast(dtype = hcl.Float(bits = 32), expr = 0.000000), f_conv[final_s0_x, final_s0_y, final_s0_z, final_s0_n], hcl.cast(dtype = hcl.Float(bits = 32), expr = 0.000000))
return final
def kernel(trainData, testData, itemMem, idMem, rdv1, rdv2):
def train_encoding(m, preTrainData):
train_temp = hcl.compute((trainData.shape[1], dim), lambda x, y: itemMem[trainData[m][x]][y] ^ idMem[x][y], name = "train_temp")
k1 = hcl.reduce_axis(0, trainData.shape[1], 'k1')
train_result = hcl.compute((dim,), lambda x: hcl.sum(train_temp[k1, x], axis = k1, dtype=hcl.Int()), name = "train_result")
with hcl.for_(0, dim) as n:
preTrainData[m][n] = train_result[n]
with hcl.if_((m + 1) % 1000 == 0):
hcl.print((m+1), "Finish encoding %d training data\n")
def test_encoding(m, preTestData):
test_temp = hcl.compute((testData.shape[1], dim), lambda x, y: itemMem[testData[m][x]][y]^idMem[x][y], name = "test_temp")
k2 = hcl.reduce_axis(0, testData.shape[1], 'k2')
test_result = hcl.compute((dim,), lambda x: hcl.sum(test_temp[k2, x], axis = k2, dtype=hcl.Int()), name = "test_result")
with hcl.for_(0, dim) as n:
preTestData[m][n] = test_result[n]
with hcl.if_((m+1)%100 == 0):
hcl.print((m+1), "Finish encoding %d testing data\n")
#Encoding
hcl.print((), "Encoding the training data into HDVs.\n")
preTrainData = hcl.compute((trainData.shape[0], dim), lambda x, y: 0, "preTrainData")
hcl.mutate((trainData.shape[0], ), lambda x: train_encoding(x, preTrainData))
hdTrainData = hcl.compute((trainData.shape[0], dim), lambda x, y: 0, "hdTrainData", dtype=hcl.UInt(1))
with hcl.Stage("S1"):
with hcl.if_(trainData.shape[1] % 2 == 0):
hcl.print((), "Use the random vector\n")
hcl.update(hdTrainData, lambda x, y: hcl.select(preTrainData[x][y] + rdv1[x][y] - trainData.shape[1]/2 > 0, 1, 0))
with hcl.else_():
hcl.update(hdTrainData, lambda x, y: hcl.select(preTrainData[x][y] - trainData.shape[1]/2 > 0, 1, 0))
hcl.print((),"Encoding the testing data into HDVs.\n")
preTestData = hcl.compute((testData.shape[0], dim), lambda x, y: 0, "preTestData")
hcl.mutate((testData.shape[0], ), lambda x: test_encoding(x, preTestData))
hdTestData = hcl.compute((testData.shape[0], dim), lambda x, y: 0, "hdTestData", dtype=hcl.UInt(1))
with hcl.Stage("S2"):
with hcl.if_(testData.shape[1] % 2 == 0):
hcl.print((), "Use the random vector\n")
hcl.update(hdTestData, lambda x, y: hcl.select(preTestData[x][y] + rdv2[x][y] - testData.shape[1]/2 > 0, 1, 0))
with hcl.else_():
hcl.update(hdTestData, lambda x, y: hcl.select(preTestData[x][y] - testData.shape[1]/2 > 0, 1, 0))
###data_packing
pack_train = hcl.pack(hdTrainData, axis=1, dtype=hcl.UInt(bw), name="pack_train")
pack_test = hcl.pack(hdTestData, axis=1, dtype=hcl.UInt(bw), name="pack_test")
return pack_train, pack_test
def genpack(nn, cc, hh, ww):
out = hcl.scalar(0, name=name + "_pack", dtype=hcl.UInt(bitwidth))
with hcl.for_(0, bitwidth) as k:
out[0][(k + 1):k] = hcl.select(
data[nn, cc * bitwidth + k, hh, ww] + alpha[cc * bitwidth + k]
> 0, 1, 0)
return out[0]
def RSign(data, alpha, name="rsign"):
assert data.shape[1] == alpha.shape[0]
return hcl.compute(data.shape,
lambda nn, cc, ww, hh: hcl.select(
data[nn, cc, ww, hh] + alpha[cc] > 0, 1, 0),
name=name,
dtype=hcl.UInt(1))
def kernel(A, B):
return hcl.compute(
(8, 8),
lambda y, x: hcl.select(x < 4, A[y][x], B[y][x]),
"C",
dtype=hcl.Int(8),
)
def elu(out, x, alpha):
assert len(x.shape) == 2, "only support 2-dim ELU"
m, n = x.shape
k = hcl.reduce_axis(0, n)
return hcl.update(
out, lambda i, j: hcl.select(x[i, j] < 0, alpha *
(hcl.exp(x[i, j]) - 1), x[i, j]))
def prelu(out, x, alpha):
assert len(x.shape) == 2, "only support 2-dim PReLU"
m, n = x.shape
k = hcl.reduce_axis(0, n)
return hcl.update(
out, lambda i, j: hcl.select(x[
i, j] < 0, hcl.cast(x.dtype, alpha[j] * x[i, j]), x[i, j]))
def zero_pad2d(inputs, padding=0):
"""Zero padding for 2d tensor
Args:
-----------------------------
inputs : hcl.tensor.Tensor
shape [batch, channel, height, width]
padding: (optional:0) int or tuple
expected: (h_pad_up, h_pad_down, w_pad_up, w_pad_down)
-----------------------------
Returns:
-----------------------------
hcl.tensor.Tensor
shape [batch, channel, padded_height, padded_width]
-----------------------------
"""
padding = (padding, padding, padding, padding) if isinstance(padding, (int, )) else padding
if len(padding) == 2:
padding = (padding[0], padding[0], padding[1], padding[1])
padding_zero = 0.0 if "float" in inputs.dtype else 0
batch_size, in_channel, height, width = inputs.shape
return hcl.compute(
(batch_size, in_channel, height + padding[0] + padding[1], width + padding[2] + padding[3]),
lambda b, c, h, w: hcl.select(
hcl.tvm.te.all(h >= padding[0], h < height + padding[0], w >= padding[2], w < width + padding[2]),
inputs[b, c, h - padding[0], w - padding[2]],
padding_zero
),
name='Padding'
)
def thresholdedrelu(out, x, theta):
assert len(x.shape) == 2, "only support 2-dim ThresholdedReLU"
m, n = x.shape
k = hcl.reduce_axis(0, n)
return hcl.update(
out, lambda i, j: hcl.select(x[i, j] > theta, x[i, j],
hcl.cast(x.dtype, 0)))
def broadcast_min(input1, input2, name='broadcast_min'):
input1_mod, input2_mod, switch = broadcast_set(input1, input2)
if (switch):
return hcl.compute(
input1_mod.shape,
lambda *x: hcl.select(
input1_mod[x] < input2_mod[_broadcast(input2_mod.shape, x)],
input1_mod[x], input2_mod[_broadcast(input2_mod.shape, x)]),
name=name)
else:
return hcl.compute(
input2_mod.shape,
lambda *x: hcl.select(
input1_mod[_broadcast(input1_mod.shape, x)] < input2_mod[x],
input1_mod[_broadcast(input1_mod.shape, x)], input2_mod[x]),
name=name)
def abs(array, dtype=None, name='abs'):
if dtype is None:
dtype = array.dtype
return hcl.compute(
array.shape,
lambda *x: hcl.select(array[x] < 0, -array[x], array[x]),
dtype=dtype,
name=name)
def RPReLU(data, x0, y0, beta, name="rprelu", dtype=None):
assert data.shape[1] == beta.shape[0] \
and x0.shape[0] == y0.shape[0] \
and beta.shape[0] == x0.shape[0]
dtype = data.dtype if dtype == None else dtype
return hcl.compute(data.shape, lambda nn, cc, ww, hh:
hcl.select(data[nn,cc,ww,hh] + x0[cc] > 0,
data[nn,cc,ww,hh] + x0[cc],
beta[cc] * (data[nn,cc,ww,hh] + x0[cc])) + y0[cc],
name=name, dtype=dtype)
def sobel(B, G):
r = hcl.reduce_axis(0, 3)
c = hcl.reduce_axis(0, 3)
return hcl.compute(
(height, width),
lambda x, y: hcl.select(
hcl.and_(x > 0, x < (height - 1), y > 0, y < (width - 1)),
hcl.sum(B[x + r, y + c] * G[r, c], axis=[r, c]), B[x, y]),
"D",
dtype=hcl.Float())
def top(input, ):
final_total_extent_1 = (
hcl.cast(dtype=hcl.Int(bits=64), expr=final_extent_1) *
hcl.cast(dtype=hcl.Int(bits=64), expr=final_extent_0))
max_local = hcl.compute((final_extent_0, final_extent_1),
lambda x, y: 0,
name="max_local",
dtype=hcl.UInt(bits=16))
with hcl.Stage("max_local"):
with hcl.for_(final_min_1, final_extent_1,
name="max_local_s0_y") as max_local_s0_y:
with hcl.for_(final_min_0, final_extent_0,
name="max_local_s0_x") as max_local_s0_x:
maximum = hcl.compute((1, 1),
lambda x, y: 0,
name="maximum",
dtype=hcl.UInt(bits=16))
with hcl.Stage("maximum"):
maximum[max_local_s0_x,
max_local_s0_y] = hcl.cast(dtype=hcl.UInt(bits=16),
expr=0)
with hcl.for_(
0, 3,
name="maximum_s1_box__y") as maximum_s1_box__y:
with hcl.for_(
0, 3,
name="maximum_s1_box__x") as maximum_s1_box__x:
maximum[max_local_s0_x,
max_local_s0_y] = hcl.select(
maximum[max_local_s0_x, max_local_s0_y]
> input[(max_local_s0_x +
maximum_s1_box__x),
(max_local_s0_y +
maximum_s1_box__y)],
maximum[max_local_s0_x,
max_local_s0_y],
input[(max_local_s0_x +
maximum_s1_box__x),
(max_local_s0_y +
maximum_s1_box__y)])
max_local[max_local_s0_x,
max_local_s0_y] = maximum[max_local_s0_x,
max_local_s0_y]
final = hcl.compute((640, 480),
lambda x, y: 0,
name="final",
dtype=hcl.UInt(bits=16))
with hcl.Stage("final"):
with hcl.for_(final_min_1, final_extent_1,
name="final_s0_y") as final_s0_y:
with hcl.for_(final_min_0, final_extent_0,
name="final_s0_x") as final_s0_x:
final[final_s0_x, final_s0_y] = max_local[final_s0_x,
final_s0_y]
return final
def guassian(A, G):
h = hcl.reduce_axis(0, size)
w = hcl.reduce_axis(0, size)
return hcl.compute(
(height, width),
lambda x, y: hcl.select(
hcl.and_(x > (size - 1), x < (height - size), y > (size - 1), y <
(width - size)),
hcl.sum(A[x + h, y + w] * G[h, w], axis=[h, w]), A[x, y]),
"F",
dtype=hcl.Float())
def Gaussian_Sobel_filters(A, G, Fx, Fy):
h = hcl.reduce_axis(0, kernel_size)
w = hcl.reduce_axis(0, kernel_size)
B = hcl.compute(
(height, width),
lambda y, x: hcl.select(
hcl.and_(y > (k - 1), y < (width - k), x > (k - 1), x <
(height - k)),
hcl.sum(A[y + w, x + h] * G[w, h], axis=[w, h]), B[y, x]),
"B",
dtype=hcl.Float())
# Sobel Filters
r = hcl.reduce_axis(0, 3)
c = hcl.reduce_axis(0, 3)
Gx = hcl.compute(
(height, width),
lambda y, x: hcl.select(
hcl.and_(y > (k - 1), y < (width - k), x > (k - 1), x <
(height - k)),
hcl.sum(B[y + r, x + c] * Fx[r, c], axis=[r, c]), B[y, x]),
"Gx",
dtype=hcl.Float())
t = hcl.reduce_axis(0, 3)
g = hcl.reduce_axis(0, 3)
Gy = hcl.compute(
(height, width),
lambda y, x: hcl.select(
hcl.and_(y > (k - 1), y < (width - k), x > (k - 1), x <
(height - k)),
hcl.sum(B[y + t, x + g] * Fy[t, g], axis=[t, g]), B[y, x]),
"Gy",
dtype=hcl.Float())
# return the intensity matrix and the edge direction matrix?
return hcl.compute(
(height, width),
lambda y, x:
(hcl.sqrt(Gx[y][x] * Gx[y][x] + Gy[y][x] * Gy[y][x])) / 4328 * 255,
dtype=hcl.Float())
def learn(k, in_train, prototype, prototypeCounter):
#Find samples that have the label k
match = hcl.compute(in_train.shape,
lambda x,y: hcl.select(trainLabels[x] == k, in_train[x][y], 0), "match", dtype=hcl.UInt(in_bw))
#Record the number of these samples
with hcl.for_(0, in_train.shape[0]) as a:
with hcl.if_(trainLabels[a] == k):
max[k] += 1
#Do hdc sum on these samples' hdv
r = hcl.reduce_axis(0, in_train.shape[0], 'r')
with hcl.for_(0, in_bw) as bit:
bit_sum = hcl.compute((in_train.shape[1],), lambda y: hcl.sum(match[r][y][bit], axis=r), "result")
#Do the binary voting
sum1 = hcl.compute((in_train.shape[1],), lambda x: 0, "sum1", dtype=hcl.UInt(1))
with hcl.if_(max[k] % 2 == 0):
hcl.update(sum1, lambda x: hcl.select(bit_sum[x] + pack_rdv3[k][x][bit] - max[k]/2 > 0, 1, 0))
with hcl.else_():
hcl.update(sum1, lambda x: hcl.select(bit_sum[x] - max[k]/2 > 0, 1, 0))
#Push the binary sum to prototype and the original sum to prototypeCounter
with hcl.for_(0, in_train.shape[1]) as t:
prototype[k][t][bit] = sum1[t]
prototypeCounter[k][t*in_bw+bit] = bit_sum[t]
def prelu(data, alpha, axis=1):
def _axis_ind(axis, ind):
ind = ind[0]
new_ind = []
for i in range(len(ind)):
if i == axis:
new_ind = ind[i]
return tuple(new_ind)
return hcl.compute(
data.shape, lambda *x: hcl.select(
data[x] < 0,
hcl.cast(data.dtype, alpha[_axis_ind(axis, x)] * data[x]), data[x])
)
def sobel_x(A, Gx):
B = hcl.compute((height, width),
lambda x, y: A[x][y][0] + A[x][y][1] + A[x][y][2],
"B",
dtype=hcl.Float())
r = hcl.reduce_axis(0, 3)
c = hcl.reduce_axis(0, 3)
return hcl.compute(
(height, width),
lambda x, y: hcl.select(
hcl.and_(x > 0, x < (height - 1), y > 0, y < (width - 1)),
hcl.sum(B[x + r, y + c] * Gx[r, c], axis=[r, c]), B[x, y]),
"X",
dtype=hcl.Float())
def top(input, ):
final_total_extent_1 = (hcl.cast(dtype = hcl.Int(bits = 64), expr = final_extent_1) * hcl.cast(dtype = hcl.Int(bits = 64), expr = final_extent_0))
final_total_extent_2 = (final_total_extent_1 * hcl.cast(dtype = hcl.Int(bits = 64), expr = final_extent_2))
linear = hcl.compute(((final_extent_0 + 2), (final_extent_1 + 2), final_extent_2), lambda x, y, z: 0, name = "linear", dtype = hcl.Float(bits = 32))
with hcl.Stage("linear"):
with hcl.for_(final_min_2, final_extent_2, name = "linear_s0_c") as linear_s0_c:
with hcl.for_(final_min_1, (final_extent_1 + 2), name = "linear_s0_y") as linear_s0_y:
with hcl.for_(final_min_0, (final_extent_0 + 2), name = "linear_s0_x") as linear_s0_x:
t4 = input[linear_s0_x, linear_s0_y, linear_s0_c]
linear[linear_s0_x, linear_s0_y, linear_s0_c] = hcl.select((hcl.cast(dtype = hcl.Float(bits = 32), expr = 0.040450) < t4), hcl.power(((t4 * hcl.cast(dtype = hcl.Float(bits = 32), expr = 0.947867)) + hcl.cast(dtype = hcl.Float(bits = 32), expr = 0.052133)), hcl.cast(dtype = hcl.Float(bits = 32), expr = 2.400000)), (t4 * hcl.cast(dtype = hcl.Float(bits = 32), expr = 0.077399)))
blur_x = hcl.compute((final_extent_0, (final_extent_1 + 2), final_extent_2), lambda x, y, z: 0, name = "blur_x", dtype = hcl.Float(bits = 32))
with hcl.Stage("blur_x"):
with hcl.for_(final_min_2, final_extent_2, name = "blur_x_s0_c") as blur_x_s0_c:
with hcl.for_(final_min_1, (final_extent_1 + 2), name = "blur_x_s0_y") as blur_x_s0_y:
with hcl.for_(final_min_0, final_extent_0, name = "blur_x_s0_x") as blur_x_s0_x:
blur_x[blur_x_s0_x, blur_x_s0_y, blur_x_s0_c] = ((linear[(blur_x_s0_x + 2), blur_x_s0_y, blur_x_s0_c] + (linear[blur_x_s0_x, blur_x_s0_y, blur_x_s0_c] + linear[(blur_x_s0_x + 1), blur_x_s0_y, blur_x_s0_c])) * hcl.cast(dtype = hcl.Float(bits = 32), expr = 0.333333))
blur_y = hcl.compute((final_extent_0, final_extent_1, final_extent_2), lambda x, y, z: 0, name = "blur_y", dtype = hcl.Float(bits = 32))
with hcl.Stage("blur_y"):
with hcl.for_(final_min_2, final_extent_2, name = "blur_y_s0_c") as blur_y_s0_c:
with hcl.for_(final_min_1, final_extent_1, name = "blur_y_s0_y") as blur_y_s0_y:
with hcl.for_(final_min_0, final_extent_0, name = "blur_y_s0_x") as blur_y_s0_x:
blur_y[blur_y_s0_x, blur_y_s0_y, blur_y_s0_c] = ((blur_x[blur_y_s0_x, (blur_y_s0_y + 2), blur_y_s0_c] + (blur_x[blur_y_s0_x, blur_y_s0_y, blur_y_s0_c] + blur_x[blur_y_s0_x, (blur_y_s0_y + 1), blur_y_s0_c])) * hcl.cast(dtype = hcl.Float(bits = 32), expr = 0.333333))
srgb = hcl.compute((final_extent_0, final_extent_1, final_extent_2), lambda x, y, z: 0, name = "srgb", dtype = hcl.Float(bits = 32))
with hcl.Stage("srgb"):
with hcl.for_(final_min_2, final_extent_2, name = "srgb_s0_c") as srgb_s0_c:
with hcl.for_(final_min_1, final_extent_1, name = "srgb_s0_y") as srgb_s0_y:
with hcl.for_(final_min_0, final_extent_0, name = "srgb_s0_x") as srgb_s0_x:
t5 = blur_y[srgb_s0_x, srgb_s0_y, srgb_s0_c]
srgb[srgb_s0_x, srgb_s0_y, srgb_s0_c] = hcl.select((hcl.cast(dtype = hcl.Float(bits = 32), expr = 0.003131) < t5), ((hcl.power(t5, hcl.cast(dtype = hcl.Float(bits = 32), expr = 0.416667)) * hcl.cast(dtype = hcl.Float(bits = 32), expr = 1.055000)) + hcl.cast(dtype = hcl.Float(bits = 32), expr = -0.055000)), (t5 * hcl.cast(dtype = hcl.Float(bits = 32), expr = 12.920000)))
final = hcl.compute((766, 1278, 3), lambda x, y, z: 0, name = "final", dtype = hcl.Float(bits = 32))
with hcl.Stage("final"):
with hcl.for_(final_min_2, final_extent_2, name = "final_s0_c") as final_s0_c:
with hcl.for_(final_min_1, final_extent_1, name = "final_s0_y") as final_s0_y:
with hcl.for_(final_min_0, final_extent_0, name = "final_s0_x") as final_s0_x:
final[final_s0_x, final_s0_y, final_s0_c] = srgb[final_s0_x, final_s0_y, final_s0_c]
return final
def sobel_y(A, Gy):
B = hcl.compute((height, width),
lambda x, y: A[x][y][0] + A[x][y][1] + A[x][y][2],
"B",
dtype=hcl.Float())
t = hcl.reduce_axis(0, 3)
g = hcl.reduce_axis(0, 3)
return hcl.compute(
(height, width),
lambda x, y: hcl.select(
hcl.and_(x > 0, x < (height - 1), y > 0, y < (width - 1)),
hcl.sum(B[x + t, y + g] * Gy[t, g], axis=[t, g]), B[x, y]),
"Y",
dtype=hcl.Float())
