hmmmmm

Author: Jian Weng

models/arm/compile.py

@@ -95,6 +95,8 @@ def tracer(module, info, is_before):
 
     print('Model Load!')
     import tensorizer
+    from tensorizer import tune
+    tune.enable = True
     with tvm.transform.PassContext(config={'tir.add_lower_pass': [(1, tensorizer.rewrite)]},
                                    trace=tracer, opt_level=3):
             graph, lib, params = relay.build_module.build(

models/cuda/compile.py

@@ -62,7 +62,6 @@ def load_model(symbol_file, param_file, logger=None):
 
 
 def compile_via_tvm(sym, arg_params, aux_params, symbol_file, data_shape):
-    tune = False
 
     input_shape = [1] + list(data_shape)
     input_dict = {'data': input_shape}
@@ -91,6 +90,8 @@ def tracer(module, info, is_before):
             print('Executes: ', info.name, (time.time() - timing) * 1000)
 
     import tensorizer
+    from tensorizer import tune
+    tune.enable = True
     with tvm.transform.PassContext(config={'tir.add_lower_pass': [(1, tensorizer.rewrite)]},
                                    trace=tracer, opt_level=3, disabled_pass=['FoldScaleAxis']):
     #with tvm.transform.PassContext(trace=tracer, opt_level=4):

poc/tensorcore/conv-tc.cu

@@ -118,17 +118,20 @@ int main() {
                                          yDesc, y) );
   cudaDeviceSynchronize();
   begin_roi();
-  checkCudnnErr( cudnnConvolutionForward(handle, (void*)(&alpha),
-                                         xDesc, x,
-                                         wDesc, w, convDesc, algo,
-                                         workSpace, workSpaceSize,
-                                         (void*)(&beta),
-                                         yDesc, y) );
-  checkCudaErr(cudaDeviceSynchronize());
-  auto elps = end_roi();
+  for (int i = 0; i < 10; ++i) {
+    checkCudnnErr( cudnnConvolutionForward(handle, (void*)(&alpha),
+                                           xDesc, x,
+                                           wDesc, w, convDesc, algo,
+                                           workSpace, workSpaceSize,
+                                           (void*)(&beta),
+                                           yDesc, y) );
+    checkCudaErr(cudaDeviceSynchronize());
+  }
+  double elps = end_roi();
+  elps /= 10.;
   std::cout << "Exec: " << elps << "us" << std::endl;
   std::cout << ((double) dimY[0] * dimY[1] * dimY[2] * dimY[3] * dimB[1] * dimB[2] * dimB[3]) / elps / 1000.
             << " GFLOP/s" << std::endl;
 
   return 0;
-}
+}

poc/tensorcore/input

@@ -14,26 +14,26 @@
 (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)]
+
 for i in workloads:
+    res = []
     for dtype in [0, 2]:
-        for tc in range(2):
+        for tc in range(2 if dtype == 2 else 1):
             exec_time = []
             for algo in range(8):
                 with open('input', 'w') as f:
-                    f.write('1\n')
+                    #f.write('1\n')
                     for j in i:
                         f.write(str(j) + '\n')
                     f.write(str(dtype) + '\n')
                     f.write(str(algo) + '\n')
                     f.write(str(tc) + '\n')
                 try:
-                    avg = []
-                    for j in range(10):
-                        res = subprocess.check_output('./tensorcore.nvcc < input', shell=True).decode('utf-8')
-                        res = int(res.split('\n')[0].lstrip('Exec: ').rstrip('us'))
-                        avg.append(res)
-                    exec_time.append(sum(avg) / len(avg))
+                    proc = subprocess.check_output('./tensorcore.nvcc < input', shell=True).decode('utf-8')
+                    proc = float(proc.split('\n')[0].lstrip('Exec: ').rstrip('us'))
+                    exec_time.append(proc)
                 except:
                     pass
-                    #print(i, algo, tc, 'fails')
-            print(i, ['fp32', None, 'fp16'][dtype], tc, min(exec_time))
+            res.append(min(exec_time))
+    print(i, ['fp32', None, 'fp16'][dtype], tc, res)