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())