cgo ae update

Author: Jian Weng

apps/cpu/kernel/.gitignore

@@ -0,0 +1 @@
+input

apps/cpu/kernel/conv2d.py

@@ -10,38 +10,43 @@
 import topi
 from tvm.relay import op
 
-##ic, h, w, oc, _, kh, kw, sh, sw = map(int, input().split())
-#ic, h, w, oc, kh, sh = map(int, input().split())
-#kw = kh
-#sw = sh
-#
-#if ic % 4:
-#    ic += 4 - ic % 4
-#
-#if oc % 16:
-#    oc += 16 - oc % 16
-#
-#a = tvm.te.placeholder((1, ic // 16, h, w, 16), dtype='int8')
-#if ic % 16 == 0:
-#    b = tvm.te.placeholder((oc // 16, ic // 16, kh, kw, 4, 16, 4), dtype='int8')
-#else:
-#    assert ic % 4 == 0
-#    a = tvm.te.placeholder((1, ic // 4, h, w, 4), dtype='int8')
-#    b = tvm.te.placeholder((oc // 16, ic // 4, kh, kw, 1, 16, 4), dtype='int8')
+#ic, h, w, oc, _, kh, kw, sh, sw = map(int, input().split())
+_, ic, h, w, oc, _, kh, _, sh, _ = map(int, input().split())
+kw = kh
+sw = sh
 
-N, C, H, W, c, O, I, KH, KW, e, o, i, sh, sw = map(int, input().split())
+if ic % 4:
+    ic += 4 - ic % 4
 
-a = tvm.te.placeholder((N, C, H, W, c), dtype='int8')
-b = tvm.te.placeholder((O, I, KH, KW, e, o, i), dtype='int8')
+if oc % 16:
+    oc += 16 - oc % 16
+
+a = tvm.te.placeholder((1, ic // 16, h, w, 16), dtype='int8')
+if ic % 16 == 0:
+    b = tvm.te.placeholder((oc // 16, ic // 16, kh, kw, 4, 16, 4), dtype='int8')
+else:
+    assert ic % 4 == 0
+    a = tvm.te.placeholder((1, ic // 4, h, w, 4), dtype='int8')
+    b = tvm.te.placeholder((oc // 16, ic // 4, kh, kw, 1, 16, 4), dtype='int8')
+
+#N, C, H, W, c, O, I, KH, KW, e, o, i, sh, sw = map(int, input().split())
+
+#a = tvm.te.placeholder((N, C, H, W, c), dtype='int8')
+#b = tvm.te.placeholder((O, I, KH, KW, e, o, i), dtype='int8')
 
 passes = [(1, tensorizer.rewrite)]
 from tensorizer import tune
-tune.cpu_idx = 0
+tune.cpu_idx = -1
 target = -1
 results = []
-virgin = True
+result = 1e9
+
 while True:
     with tvm.transform.PassContext(opt_level=3, config={'tir.add_lower_pass': passes}), tvm.target.create('llvm -mcpu=cascadelake'):
+        if tune.cpu_idx == -1:
+            tune.cpu_idx = 0
+            tune.parallel_only = True
+
         conv = topi.nn.conv2d_NCHWc_int8(a, b, stride=(sh, sw), padding=0, dilation=1, out_dtype='int32',
                                          layout='NCHW4c', out_layout='NCHW16c')
         sch = tensorizer.INTRINSICS['vnni']['schedule']([conv], (sh, sw))
@@ -59,9 +64,16 @@
             res = fte(nd_a, nd_b, nd_c)
         results.append(res.mean)
 
+        if tune.parallel_only:
+            tune.cpu_idx = -1
+            tune.parallel_only = False
+
+        if res.mean < result:
+            target = tune.cpu_idx
+            result = res.mean
+
     relay.backend.compile_engine.get().clear()
     tune.cpu_idx += 1
-    break
     if tune.cpu_idx - target > 8:
         break
     if tune.cpu_idx >= tune.total_idx:

apps/cpu/kernel/cpu-tune.log

@@ -430,3 +430,19 @@
 (1, 2, 112, 112, 16) (1, 2, 1, 1, 4, 16, 4) (1, 1) [5.283333333333333e-06]
 (1, 128, 7, 7, 16) (32, 128, 1, 1, 4, 16, 4) (1, 1) [6.08829e-05]
 (1, 36, 14, 14, 16) (4, 36, 1, 1, 4, 16, 4) (1, 1) [1.0045566666666667e-05]
+(1, 18, 35, 35, 16) (24, 18, 3, 3, 4, 16, 4) (2, 2) [0.0003514215333333333, 9.668606666666667e-05, 0.00011963883333333332, 0.0001122919, 0.0002835423, 8.759403333333333e-05, 0.0001860404, 0.00014346666666666667, 0.00039427413333333327, 0.00039437946666666667, 0.00016885596666666667, 0.00028357029999999996, 0.000281523, 0.0004038882]
+(1, 10, 9, 9, 16) (14, 10, 3, 3, 4, 16, 4) (1, 1) [2.265226666666667e-05, 8.446466666666666e-06, 1.0891100000000001e-05, 1.31841e-05, 1.0353866666666666e-05, 1.2766766666666668e-05, 1.1706900000000001e-05, 1.1118366666666667e-05, 2.4449199999999996e-05, 2.41494e-05]
+(1, 66, 7, 7, 16) (12, 66, 1, 1, 4, 16, 4) (1, 1) [1.4224633333333333e-05, 6.878800000000001e-06, 8.318566666666667e-06, 9.253066666666668e-06, 8.469466666666667e-06, 9.702333333333332e-06, 8.547600000000001e-06, 9.105799999999998e-06, 1.4786199999999999e-05, 1.5234766666666663e-05]
+(1, 5, 73, 73, 16) (12, 5, 3, 3, 4, 16, 4) (1, 1) [0.0007445264666666666, 0.00036293423333333337, 0.0004282294666666667, 0.00040284786666666667, 0.00036327886666666663, 0.00042169116666666666, 0.00047538076666666675, 0.00042853806666666666, 0.0004993923, 0.00036451126666666665]
+(1, 8, 16, 16, 16) (8, 8, 3, 3, 4, 16, 4) (1, 1) [3.388693333333333e-05, 1.1279599999999998e-05, 1.0211000000000003e-05, 1.0718e-05, 4.224366666666667e-05, 2.1665866666666665e-05, 1.1584133333333333e-05, 3.2161733333333334e-05, 3.2217433333333334e-05, 3.2244733333333326e-05, 3.0116033333333334e-05]
+(1, 12, 16, 16, 16) (12, 12, 3, 3, 4, 16, 4) (1, 1) [6.998316666666667e-05, 1.8598966666666666e-05, 1.8014533333333332e-05, 1.8093833333333333e-05, 7.989396666666665e-05, 4.3606333333333336e-05, 1.838603333333333e-05, 7.076240000000001e-05, 7.021e-05, 7.07758e-05, 5.9090766666666673e-05]
+(1, 16, 16, 16, 16) (16, 16, 3, 3, 4, 16, 4) (1, 1) [0.0001271999, 2.9112400000000003e-05, 2.99537e-05, 3.089303333333333e-05, 0.00012914833333333334, 7.562293333333334e-05, 2.9749866666666664e-05, 0.00012641846666666662, 0.0001263090666666667, 0.00012664806666666664]
+(1, 64, 14, 14, 16) (32, 64, 1, 1, 4, 16, 4) (1, 1) [0.00010905626666666667, 2.81171e-05, 2.7718833333333336e-05, 2.8254033333333328e-05, 0.0001004875, 7.107966666666667e-05, 2.9853599999999995e-05, 9.300396666666667e-05, 0.00010172613333333334, 0.00010611056666666666, 3.0567733333333336e-05]
+(1, 8, 16, 16, 16) (10, 8, 3, 3, 4, 16, 4) (1, 1) [4.00728e-05, 1.3028833333333335e-05, 1.2608533333333333e-05, 1.2518066666666666e-05, 4.4550733333333334e-05, 2.5933866666666666e-05, 1.2598300000000001e-05, 4.01845e-05, 4.0060066666666666e-05, 3.942363333333333e-05, 3.655853333333333e-05, 3.844476666666666e-05]
+(1, 36, 14, 14, 16) (12, 36, 1, 1, 4, 16, 4) (1, 1) [2.2557133333333335e-05, 9.7363e-06, 1.0471766666666667e-05, 9.455833333333335e-06, 2.3511866666666668e-05, 1.597273333333333e-05, 9.989866666666667e-06, 2.2900466666666662e-05, 2.26288e-05, 2.29245e-05, 2.1308e-05, 2.338776666666667e-05]
+(1, 6, 16, 16, 16) (8, 6, 3, 3, 4, 16, 4) (1, 1) [2.5126666666666668e-05, 9.6417e-06, 9.640433333333333e-06, 9.442599999999998e-06, 3.2072099999999994e-05, 1.667773333333333e-05, 9.598066666666667e-06, 2.48075e-05, 2.417616666666667e-05, 2.4684066666666666e-05, 2.3512433333333335e-05, 2.6022999999999998e-05]
+(1, 64, 14, 14, 16) (16, 64, 1, 1, 4, 16, 4) (1, 1) [5.522673333333333e-05, 1.6551866666666663e-05, 1.7416066666666667e-05, 1.63774e-05, 5.077093333333333e-05, 3.9541566666666674e-05, 1.7163833333333333e-05, 5.5439433333333343e-05, 5.500383333333334e-05, 5.518076666666666e-05, 5.034303333333334e-05, 5.2097466666666674e-05]
+(1, 36, 14, 14, 16) (8, 36, 1, 1, 4, 16, 4) (1, 1) [1.65873e-05, 6.972e-06, 7.9898e-06, 8.129266666666667e-06, 1.7917033333333335e-05, 1.3090033333333333e-05, 8.531433333333334e-06, 1.59983e-05, 1.6602433333333336e-05, 1.6452366666666665e-05]
+(1, 4, 29, 29, 16) (6, 4, 3, 3, 4, 16, 4) (1, 1) [3.99176e-05, 2.10778e-05, 3.882036666666667e-05, 1.9675066666666664e-05, 2.0670900000000002e-05, 3.98366e-05, 3.8784699999999995e-05, 3.0413866666666667e-05, 2.063113333333333e-05, 3.92953e-05, 2.0253333333333335e-05, 4.06086e-05]
+(1, 4, 56, 56, 16) (8, 4, 1, 1, 4, 16, 4) (2, 2) [5.222933333333333e-06, 6.210866666666667e-06, 5.7302e-06, 6.315266666666668e-06, 1.03406e-05, 6.2129e-06, 5.113633333333333e-06, 6.020233333333333e-06, 7.795133333333333e-06, 7.409666666666667e-06, 6.312966666666667e-06, 5.339066666666666e-06, 1.11085e-05, 5.6410666666666675e-06, 6.255299999999999e-06]
+(1, 38, 14, 14, 16) (12, 38, 1, 1, 4, 16, 4) (1, 1) [2.3831933333333335e-05, 8.716766666666668e-06, 9.773e-06, 9.623833333333333e-06, 2.4328300000000002e-05, 1.7445633333333328e-05, 1.0776666666666668e-05, 2.384416666666666e-05, 2.3965666666666667e-05, 2.4390266666666666e-05]

apps/cpu/kernel/input

@@ -15,34 +15,28 @@
 #(1024,15,15,2048,1024,1,1,2,2),
 #(128,65,65,256,128,3,3,2,2)]
 
-workloads = [
-(256, 16, 16, 256, 3, 1),
-(512, 9, 9, 512, 3, 1),
-(128, 30, 30, 128, 3, 1),
-(64, 56, 56, 128, 1, 2),
-(64, 58, 58, 64, 3, 1),
-(128, 28, 28, 256, 1, 2),
-(256, 16, 16, 512, 3, 2),
-(64, 58, 58, 128, 3, 2),
-(4, 230, 230, 64, 7, 2),
-(128, 30, 30, 256, 3, 2),
-(256, 14, 14, 512, 1, 2)
-]
+#workloads = [
+#(256, 16, 16, 256, 3, 1),
+#(512, 9, 9, 512, 3, 1),
+#(128, 30, 30, 128, 3, 1),
+#(64, 56, 56, 128, 1, 2),
+#(64, 58, 58, 64, 3, 1),
+#(128, 28, 28, 256, 1, 2),
+#(256, 16, 16, 512, 3, 2),
+#(64, 58, 58, 128, 3, 2),
+#(4, 230, 230, 64, 7, 2),
+#(128, 30, 30, 256, 3, 2),
+#(256, 14, 14, 512, 1, 2)
+#]
+
+workloads = [(1, 288, 35, 35, 384, 288, 3, 3, 2, 2), (1, 160, 9, 9, 224, 160, 3, 3, 1, 1), (1, 1056, 7, 7, 192, 1056, 1, 1, 1, 1), (1, 80, 73, 73, 192, 80, 3, 3, 1, 1), (1, 128, 16, 16, 128, 128, 3, 3, 1, 1), (1, 192, 16, 16, 192, 192, 3, 3, 1, 1), (1, 256, 16, 16, 256, 256, 3, 3, 1, 1), (1, 1024, 14, 14, 512, 1024, 1, 1, 1, 1), (1, 128, 16, 16, 160, 128, 3, 3, 1, 1), (1, 576, 14, 14, 192, 576, 1, 1, 1, 1), (1, 96, 16, 16, 128, 96, 3, 3, 1, 1), (1, 1024, 14, 14, 256, 1024, 1, 1, 1, 1), (1, 576, 14, 14, 128, 576, 1, 1, 1, 1), (1, 64, 29, 29, 96, 64, 3, 3, 1, 1), (1, 64, 56, 56, 128, 64, 1, 1, 2, 2), (1, 608, 14, 14, 192, 608, 1, 1, 1, 1)]
 
 for i in workloads:
     exec_time = []
     with open('input', 'w') as f:
         f.write(' '.join(map(str, i)))
     try:
-        avg = []
-        for j in range(10):
-            res = subprocess.check_output('python ./conv2d.py < input', shell=True).decode('utf-8')
-            for k in res.split('\n'):
-                if k.startswith('exec: '):
-                    res = float(k.strip('exec: '))
-            avg.append(res)
-        exec_time.append(sum(avg) / len(avg))
+        res = subprocess.check_output('python ./conv2d.py < input', shell=True).decode('utf-8')
+        print(i, 'done')
     except:
         print(i, 'fails')
-    print('!!!', i, min(exec_time))
-

apps/cpu/kernel/run.py

@@ -39,7 +39,9 @@ def tracer(module, info, is_before):
 result = 1e9
 target = -1
 from tensorizer import tune
-tune.cpu_idx = 0
+tune.cpu_idx = -1
+results = []
+
 while True:
     with tvm.transform.PassContext(opt_level=3, trace=tracer, config={'tir.add_lower_pass': [(1, tensorizer.rewrite)]}):
         graph, lib, params = tvm.relay.build(module, target='llvm -mcpu=cascadelake')
@@ -52,6 +54,7 @@ def tracer(module, info, is_before):
 
         timer = func.module.time_evaluator('run', ctx=tvm.cpu(0), number=3, repeat=10)
         timed = timer()
+        results.append(timed.mean)
 
         if timed.mean < result:
             result = timed.mean
@@ -64,7 +67,7 @@ def tracer(module, info, is_before):
     if tune.cpu_idx >= tune.total_idx:
         break
 
-with open('/home/ubuntu/Tensorization-PoC/cpu-tune.log', 'a') as f:
+with open('./cpu-tune.log', 'a') as f:
     f.write(f'{tune.ashape} {tune.bshape} {tune.strides} {target}\n')
 
-print(result, target, tune.cpu_idx)
+#print(result, target, tune.cpu_idx)

apps/cpu/relay/conv2d.py

@@ -1,8 +1,8 @@
-256
+608
 14
 14
-512
+192
+1
+1
 1
 1
-2
-2

poc/vnni/input

@@ -14,36 +14,36 @@
 #(1024,15,15,2048,1024,1,1,2,2),
 #(128,65,65,256,128,3,3,2,2)]
 #
-#workloads = [(1, 288, 35, 35, 384, 288, 3, 3, 2, 2), (1, 160, 9, 9, 224, 160, 3, 3, 1, 1), (1, 1056, 7, 7, 192, 1056, 1, 1, 1, 1), (1, 80, 73, 73, 192, 80, 3, 3, 1, 1), (1, 128, 16, 16, 128, 128, 3, 3, 1, 1), (1, 192, 16, 16, 192, 192, 3, 3, 1, 1), (1, 256, 16, 16, 256, 256, 3, 3, 1, 1), (1, 1024, 14, 14, 512, 1024, 1, 1, 1, 1), (1, 128, 16, 16, 160, 128, 3, 3, 1, 1), (1, 576, 14, 14, 192, 576, 1, 1, 1, 1), (1, 96, 16, 16, 128, 96, 3, 3, 1, 1), (1, 1024, 14, 14, 256, 1024, 1, 1, 1, 1), (1, 576, 14, 14, 128, 576, 1, 1, 1, 1), (1, 64, 29, 29, 96, 64, 3, 3, 1, 1), (1, 64, 56, 56, 128, 64, 1, 1, 2, 2), (1, 608, 14, 14, 192, 608, 1, 1, 1, 1)]
-
-workloads = [
-(256, 16, 16, 256, 3, 1),
-(512, 9, 9, 512, 3, 1),
-(128, 30, 30, 128, 3, 1),
-(64, 56, 56, 128, 1, 2),
-(64, 58, 58, 64, 3, 1),
-(128, 28, 28, 256, 1, 2),
-(256, 16, 16, 512, 3, 2),
-(64, 58, 58, 128, 3, 2),
-(4, 230, 230, 64, 7, 2),
-(128, 30, 30, 256, 3, 2),
-(256, 14, 14, 512, 1, 2)
-]
+workloads = [(1, 288, 35, 35, 384, 288, 3, 3, 2, 2), (1, 160, 9, 9, 224, 160, 3, 3, 1, 1), (1, 1056, 7, 7, 192, 1056, 1, 1, 1, 1), (1, 80, 73, 73, 192, 80, 3, 3, 1, 1), (1, 128, 16, 16, 128, 128, 3, 3, 1, 1), (1, 192, 16, 16, 192, 192, 3, 3, 1, 1), (1, 256, 16, 16, 256, 256, 3, 3, 1, 1), (1, 1024, 14, 14, 512, 1024, 1, 1, 1, 1), (1, 128, 16, 16, 160, 128, 3, 3, 1, 1), (1, 576, 14, 14, 192, 576, 1, 1, 1, 1), (1, 96, 16, 16, 128, 96, 3, 3, 1, 1), (1, 1024, 14, 14, 256, 1024, 1, 1, 1, 1), (1, 576, 14, 14, 128, 576, 1, 1, 1, 1), (1, 64, 29, 29, 96, 64, 3, 3, 1, 1), (1, 64, 56, 56, 128, 64, 1, 1, 2, 2), (1, 608, 14, 14, 192, 608, 1, 1, 1, 1)]
+#
+#workloads = [
+#(256, 16, 16, 256, 3, 1),
+#(512, 9, 9, 512, 3, 1),
+#(128, 30, 30, 128, 3, 1),
+#(64, 56, 56, 128, 1, 2),
+#(64, 58, 58, 64, 3, 1),
+#(128, 28, 28, 256, 1, 2),
+#(256, 16, 16, 512, 3, 2),
+#(64, 58, 58, 128, 3, 2),
+#(4, 230, 230, 64, 7, 2),
+#(128, 30, 30, 256, 3, 2),
+#(256, 14, 14, 512, 1, 2)
+#]
 
 for i in workloads:
     exec_time = []
     with open('input', 'w') as f:
         lst = list(i)
-        #for j in (i[1:5] + i[6:]):
-        #    f.write(str(j) + '\n')
-        f.write(str(i[0]) + '\n')
-        f.write(str(i[1]) + '\n')
-        f.write(str(i[2]) + '\n')
-        f.write(str(i[3]) + '\n')
-        f.write(str(i[4]) + '\n')
-        f.write(str(i[4]) + '\n')
-        f.write(str(i[5]) + '\n')
-        f.write(str(i[5]) + '\n')
+        for j in (i[1:5] + i[6:]):
+            f.write(str(j) + '\n')
+        #f.write(str(i[0]) + '\n')
+        #f.write(str(i[1]) + '\n')
+        #f.write(str(i[2]) + '\n')
+        #f.write(str(i[3]) + '\n')
+        #f.write(str(i[4]) + '\n')
+        #f.write(str(i[4]) + '\n')
+        #f.write(str(i[5]) + '\n')
+        #f.write(str(i[5]) + '\n')
     try:
         avg = []
         for j in range(10):

poc/vnni/run.py

@@ -163,18 +163,24 @@ def callback(op):
                     reduction.append(axis)
                 else:
                     simple.append(axis)
+            sch[op].pragma(stencil[0], 'tensorize', pragma)
+
+            if tune.parallel_only:
+                if str(op) != str(output):
+                    sch[op].reorder(*(simple + unroll + reduction + stencil))
+                else:
+                    sch[op].reorder(*([fusion] + unroll + simple + reduction + stencil))
+                return
+
             for i in unroll:
                 sch[op].unroll(i)
-            sch[op].pragma(stencil[0], 'tensorize', pragma)
             #if simple:
             #    unroll = [simple[-1]] + unroll
             #    simple = simple[:-1]
             if str(op) != str(output):
-                #sch[op].reorder(*(simple + reduction + unroll + stencil))
-                sch[op].reorder(*(simple + unroll + reduction + stencil))
+                sch[op].reorder(*(simple + reduction + unroll + stencil))
             else:
-                #sch[op].reorder(*([fusion] + simple + reduction + unroll + stencil))
-                sch[op].reorder(*([fusion] + unroll + simple + reduction + stencil))
+                sch[op].reorder(*([fusion] + simple + reduction + unroll + stencil))
 
     traverse_inline(sch, output, callback)
 

python/tensorizer/intrinsics/cpu.py

@@ -21,6 +21,7 @@ def _ceil_div(a, b):
 
 def analyze_tiling(op, pattern, max_unroll=32, max_parallel=3000):
 
+    print(tvm.arith._ffi_api)
     info = list(tvm.arith._ffi_api.MatchTensorizer(op, pattern))
     assert info
     loops = {}

python/tensorizer/intrinsics/looptiler.py

@@ -13,6 +13,8 @@
 splitk = None
 x86 = {}
 
+parallel_only = None
+
 HOME = os.getenv("HOME")
 
 def load_x86():