!71 Create a file to hold the global configs

From: @ckey_dou
Reviewed-by: @dylangeng,@anyrenwei
Signed-off-by: @anyrenwei

python/akg/global_configs.py

@@ -0,0 +1,38 @@
+# Copyright 2021 Huawei Technologies Co., Ltd
+#
+# Licensed under the Apache License, Version 2.0 (the "License");
+# you may not use this file except in compliance with the License.
+# You may obtain a copy of the License at
+#
+# http://www.apache.org/licenses/LICENSE-2.0
+#
+# Unless required by applicable law or agreed to in writing, software
+# distributed under the License is distributed on an "AS IS" BASIS,
+# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+# See the License for the specific language governing permissions and
+# limitations under the License.
+
+"""global_configs"""
+import akg.tvm
+
+CUDA_META_PATH = './cuda_meta/'
+ASCEND_META_PATH = './kernel_meta/'
+
+DUMP_IR_FLAG = 'MS_AKG_DUMP_IR'
+DUMP_CODE_FLAG = 'MS_AKG_DUMP_CODE'
+
+@akg.tvm.register_func
+def get_cuda_meta_path():
+  return CUDA_META_PATH
+
+@akg.tvm.register_func
+def get_ascend_meta_path():
+  return ASCEND_META_PATH
+
+@akg.tvm.register_func
+def get_dump_ir_flag():
+  return DUMP_IR_FLAG
+
+@akg.tvm.register_func
+def get_dump_code_flag():
+  return DUMP_CODE_FLAG

python/akg/ms/message.py

@@ -28,6 +28,8 @@ from akg import composite
 from . import cce
 from . import gpu
 from . import op_build
+from akg.global_configs import DUMP_IR_FLAG
+from akg.global_configs import DUMP_CODE_FLAG
 
 
 @vc_util.check_input_type(str)
@@ -82,8 +84,8 @@ def _compilewithjson_to_module(json_str):
                 if kernel_info['attr']:
                     for ext_arg in kernel_info['attr']:
                         op_attrs.append(ext_arg['value'])
-                dump_ir = os.getenv('MS_AKG_DUMP_IR') == "on"
-                dump_code = os.getenv('MS_AKG_DUMP_CODE') == "on"
+                dump_ir = os.getenv(DUMP_IR_FLAG) == "on"
+                dump_code = os.getenv(DUMP_CODE_FLAG) == "on"
                 utils.op_build(op_func, input_shapes, input_types, op_attrs, kernel_info['op'], dump_ir=dump_ir,
                                dump_code=dump_code)
                 return True

python/akg/ms/op_build.py

@@ -1,6 +1,6 @@
 #!/usr/bin/env python3
 # coding: utf-8
-# Copyright 2019 Huawei Technologies Co., Ltd
+# Copyright 2019-2021 Huawei Technologies Co., Ltd
 #
 # Licensed under the Apache License, Version 2.0 (the "License");
 # you may not use this file except in compliance with the License.
@@ -27,13 +27,10 @@ from akg.utils import kernel_exec as utils
 from akg.ms import save_gpu_param as gpu_utils
 from akg.utils import validation_check as vc_util
 from akg.tvm import _api_internal
-
+from akg.global_configs import CUDA_META_PATH
+from akg.global_configs import DUMP_IR_FLAG
 
 BINDS = "binds"
-MS_AKG_DUMP_IR = "MS_AKG_DUMP_IR"
-MS_AKG_DUMP_CODE = "MS_AKG_DUMP_CODE"
-MS_DAVINCI_KERNEL_PATH = "./kernel_meta/"
-
 
 @vc_util.check_input_type(list, (list, tuple), (list, tuple), (types.FunctionType, type(None)), str, str, dict)
 def op_build_to_func(opnames, computes, args, custom_schedule, device, kernel_name, attrs):
@@ -43,7 +40,7 @@ def op_build_to_func(opnames, computes, args, custom_schedule, device, kernel_na
         return None
 
     polyhedral = True
-    dump_ir = os.getenv(MS_AKG_DUMP_IR) == "on"
+    dump_ir = os.getenv(DUMP_IR_FLAG) == "on"
 
     try:
         tmp_outputs = [x.op for x in computes]
@@ -79,7 +76,7 @@ def op_build(opnames, computes, args, custom_schedule, device, kernel_name, attr
         return None
 
     if device == "cuda":
-        kernel_meta_path = "./cuda_meta_" + str(os.getpid()) + "/"
+        kernel_meta_path = CUDA_META_PATH 
         cuda_path = os.path.realpath(kernel_meta_path)
         if not os.path.isdir(cuda_path):
             os.makedirs(cuda_path)

python/akg/utils/dump_ascend_meta.py

@@ -1,6 +1,6 @@
 #!/usr/bin/env python3
 # coding: utf-8
-# Copyright 2020 Huawei Technologies Co., Ltd
+# Copyright 2020-2021 Huawei Technologies Co., Ltd
 #
 # Licensed under the Apache License, Version 2.0 (the "License");
 # you may not use this file except in compliance with the License.
@@ -21,7 +21,7 @@ import json
 import hashlib
 import logging
 import akg.tvm
-
+from akg.global_configs import ASCEND_META_PATH
 
 def write_code(js_dict, fname):
     if os.path.exists(fname):
@@ -55,7 +55,7 @@ def tvm_callback_cce_postproc(code, block_dim=1):
 
     # sha256
     buf_size = 64 * 1024  # once read 64kb
-    kernel_file_name = "kernel_meta/" + bin_file_name + bin_file_suffix
+    kernel_file_name = ASCEND_META_PATH + bin_file_name + bin_file_suffix
     sha256 = hashlib.sha256()
     with open(kernel_file_name, 'rb') as kf:
         while True:
@@ -66,9 +66,9 @@ def tvm_callback_cce_postproc(code, block_dim=1):
     title_dict["sha256"] = sha256.hexdigest()
 
     load_dict = {}
-    if not os.path.exists("kernel_meta"):
+    if not os.path.exists(ASCEND_META_PATH):
         try:
-            os.mkdir("kernel_meta")
+            os.mkdir(ASCEND_META_PATH)
         except OSError as err:
             # 17, OSError: [Errno 17] File exists
             if err.errno == 17:
@@ -76,7 +76,7 @@ def tvm_callback_cce_postproc(code, block_dim=1):
             else:
                 raise err
     else:
-        fname = "kernel_meta/" + kernel_name + "wk.json"
+        fname = ASCEND_META_PATH + kernel_name + "wk.json"
         if os.path.exists(fname):
             with open(fname, "r") as f:
                 load_dict = json.load(f)
@@ -85,7 +85,7 @@ def tvm_callback_cce_postproc(code, block_dim=1):
     final_dict = title_dict.copy()
     final_dict.update(load_dict)
 
-    json_file = "kernel_meta/" + kernel_name + ".json"
+    json_file = ASCEND_META_PATH + kernel_name + ".json"
     write_code(final_dict, json_file)
 
     return code

src/codegen/build_module.cc

@@ -737,10 +737,16 @@ void CreateCode(const std::string &code, const std::string &kernel_name, const s
   std::string file_path;
   std::string file_suffix;
   if (target_name.find("cce") != std::string::npos) {
-    file_path = std::string(kMsAscendKernelPath);
+    const auto *f = air::runtime::Registry::Get("get_ascend_meta_path");
+    CHECK(f != nullptr) << "Function get_ascend_meta_path is not registed";
+
+    file_path = (*f)().operator std::string();
     file_suffix = ".cce";
   } else if (target_name.find("cuda") != std::string::npos) {
-    file_path = std::string(kMsGpuKernelPath) + "_" + std::to_string(getpid()) + "/";
+    const auto *f = air::runtime::Registry::Get("get_cuda_meta_path");
+    CHECK(f != nullptr) << "Function get_cuda_meta_path is not registed";
+
+    file_path = (*f)().operator std::string();
     file_suffix = ".cu";
   }
 

src/composite/util.h

@@ -19,8 +19,6 @@
 #include "picojson.h"
 
 namespace akg {
-constexpr auto kMsAscendKernelPath = "./kernel_meta/";
-constexpr auto kMsGpuKernelPath = "./cuda_meta";
 constexpr auto BLOCK_IDX_X = "blockIdx.x";
 constexpr auto BLOCK_IDX_Y = "blockIdx.y";
 constexpr auto BLOCK_IDX_Z = "blockIdx.z";