diff --git a/akg b/akg index 97dc7e96c2..1e6b226a04 160000 --- a/akg +++ b/akg @@ -1 +1 @@ -Subproject commit 97dc7e96c2ffedf2e6e38310a903ffa205a6e656 +Subproject commit 1e6b226a0417d23d2d0a2333d5e80f13fe9e8d0f diff --git a/mindspore/_extends/parallel_compile/akg_compiler/akg_process.py b/mindspore/_extends/parallel_compile/akg_compiler/akg_process.py index 16ddbcb587..c6487c9f17 100644 --- a/mindspore/_extends/parallel_compile/akg_compiler/akg_process.py +++ b/mindspore/_extends/parallel_compile/akg_compiler/akg_process.py @@ -21,6 +21,7 @@ from multiprocessing import Pool, cpu_count from mindspore import log as logger from mindspore._extends.parallel_compile.akg_compiler.get_file_path import get_akg_path + def copy_json(pid_path, ppid_path): """ copy json from pid_path to ppid_path @@ -32,7 +33,7 @@ def copy_json(pid_path, ppid_path): shutil.move(os.path.join(pid_path, json_file), ppid_path) -def _compile_akg_task_gpu(*json_strs): +def _compile_akg_task_gpu(json_strs, attrs): """ compile func called in single process @@ -45,9 +46,9 @@ def _compile_akg_task_gpu(*json_strs): func = getattr(p.ms, "compilewithjson") for json_str in json_strs: - res = func(json_str) + res = func(json_str, attrs) if not res: - raise ValueError("Compile error, args: {}!".format(json_str)) + raise ValueError("Compile error, args: {}! build attrs: {}".format(json_str, attrs)) pid_path = os.path.realpath("./cuda_meta_" + str(os.getpid())) if os.path.exists(pid_path): @@ -55,23 +56,25 @@ def _compile_akg_task_gpu(*json_strs): shutil.rmtree(pid_path) -def _compile_akg_task_ascend(*json_strs): +def _compile_akg_task_ascend(json_strs, attrs): """ compile func called in single process Parameters: json_strs: list. List contains multiple kernel infos, suitable for json compile api. """ + if attrs is None: + attrs = "{}" akg_compiler = os.path.join(os.path.split( os.path.realpath(__file__))[0], "compiler.py") for json_str in json_strs: try: - subprocess.run([sys.executable, akg_compiler, json_str], text=True, check=True) + subprocess.run([sys.executable, akg_compiler, json_str, attrs], text=True, check=True) except BaseException as e: - logger.error(e, "Failed, args: {}!".format(json_str)) + logger.error(e, "Failed, args: {}! build attrs: {}".format(json_str, attrs)) -def create_akg_parallel_process(process_num, wait_time, platform=""): +def create_akg_parallel_process(process_num, wait_time, platform): """ create AkgParallelCompiler object @@ -84,7 +87,7 @@ def create_akg_parallel_process(process_num, wait_time, platform=""): class AkgProcess: """akg kernel parallel process""" - def __init__(self, process_num, wait_time, platform=""): + def __init__(self, process_num, wait_time, platform): """ Args: process_num: int. processes number @@ -103,7 +106,7 @@ class AkgProcess: self.platform = platform self.argc = 0 - def compile(self): + def compile(self, attrs=None): """ compile kernel by multi processes Return: @@ -111,13 +114,14 @@ class AkgProcess: """ if self.argc == 0: raise ValueError("json must be not null") + args = [(arg, attrs) for arg in self.args] if self.platform == "GPU": with Pool(processes=self.process_num) as pool: - res = pool.starmap_async(_compile_akg_task_gpu, self.args) + res = pool.starmap_async(_compile_akg_task_gpu, args) res.get(timeout=self.wait_time) elif self.platform == "ASCEND": with Pool(processes=self.process_num) as pool: - res = pool.starmap_async(_compile_akg_task_ascend, self.args) + res = pool.starmap_async(_compile_akg_task_ascend, args) res.get(timeout=self.wait_time) else: raise ValueError("The value of 'platform' must be 'GPU' or 'ASCEND'.") diff --git a/mindspore/_extends/parallel_compile/akg_compiler/compiler.py b/mindspore/_extends/parallel_compile/akg_compiler/compiler.py index 34f5d85ebb..55bdded469 100644 --- a/mindspore/_extends/parallel_compile/akg_compiler/compiler.py +++ b/mindspore/_extends/parallel_compile/akg_compiler/compiler.py @@ -16,7 +16,7 @@ import sys -def run_compiler(op_json): +def run_compiler(op_json, attrs=None): """ Run AKG compiler to compile op with subprocess, if this process of compilation failed, an exception will be raised @@ -31,10 +31,13 @@ def run_compiler(op_json): sys.path.insert(0, get_akg_path()) p = __import__("akg", globals(), locals(), ['ms'], 0) func = getattr(p.ms, "compilewithjson") - res = func(op_json) + res = func(op_json, attrs) if not res: raise ValueError("Compile error") if __name__ == "__main__": - run_compiler(sys.argv[1]) + if len(sys.argv) > 2: + run_compiler(sys.argv[1], sys.argv[2]) + else: + run_compiler(sys.argv[1]) diff --git a/mindspore/_extends/remote/kernel_build_server.py b/mindspore/_extends/remote/kernel_build_server.py index 965032d372..da042e95a2 100644 --- a/mindspore/_extends/remote/kernel_build_server.py +++ b/mindspore/_extends/remote/kernel_build_server.py @@ -128,27 +128,24 @@ class Messager: class AkgBuilder(): """Akg building wrapper""" - def __init__(self): - pass + def __init__(self, platform): + self.platform = platform + self.attrs = None - def create(self, process_num, waitime, platform=""): + def create(self, process_num, waitime): """ Create akg processor""" - - self.akg_processor = create_akg_parallel_process(process_num, waitime, platform) + self.akg_processor = create_akg_parallel_process(process_num, waitime, self.platform) def accept_json(self, json): """ Accept json""" - return self.akg_processor.accept_json(json) def compile(self): """Compile""" + return self.akg_processor.compile(self.attrs) - return self.akg_processor.compile() - - def handle(self, messager, arg, platform=""): + def handle(self, messager, arg): """Handle message about akg""" - if arg == 'AKG/PID': messager.send_res(os.getpid()) elif arg == 'AKG/START': @@ -156,7 +153,11 @@ class AkgBuilder(): process_num_str = messager.get_message() messager.send_ack() wait_time_str = messager.get_message() - self.create(int(process_num_str), int(wait_time_str), platform) + messager.send_ack() + self.create(int(process_num_str), int(wait_time_str)) + elif arg == 'AKG/ATTR': + messager.send_ack() + self.attrs = messager.get_message() messager.send_ack() elif arg == 'AKG/DATA': messager.send_ack() @@ -176,7 +177,7 @@ class AkgBuilder(): messager.send_ack() json = messager.get_message() try: - akg_compile_single(json) + akg_compile_single(json, self.attrs) except ValueError: messager.send_ack(False) messager.exit() diff --git a/mindspore/_extends/remote/kernel_build_server_ascend.py b/mindspore/_extends/remote/kernel_build_server_ascend.py index db3284ec84..b22d73d45a 100644 --- a/mindspore/_extends/remote/kernel_build_server_ascend.py +++ b/mindspore/_extends/remote/kernel_build_server_ascend.py @@ -63,7 +63,7 @@ class AscendMessager(Messager): super().__init__(fdin, fdout) get_logger().info("[TRACE] Ascend Messager init...") self.tbe_builder = TbeBuilder() - self.akg_builder = AkgBuilder() + self.akg_builder = AkgBuilder("ASCEND") def tbe_handle(self, arg): """ @@ -119,7 +119,7 @@ class AscendMessager(Messager): if arg.startswith('TBE'): self.tbe_handle(arg) elif arg.startswith('AKG'): - self.akg_builder.handle(self, arg, "ASCEND") + self.akg_builder.handle(self, arg) elif arg == 'FORMAT': self.send_ack() json = self.get_message() diff --git a/mindspore/_extends/remote/kernel_build_server_gpu.py b/mindspore/_extends/remote/kernel_build_server_gpu.py index 28c1efaa3a..c951508f22 100644 --- a/mindspore/_extends/remote/kernel_build_server_gpu.py +++ b/mindspore/_extends/remote/kernel_build_server_gpu.py @@ -27,7 +27,7 @@ class GpuMessager(Messager): def __init__(self, fdin, fdout): super().__init__(fdin, fdout) get_logger().info("[TRACE] GPU Messager init...") - self.akg_builder = AkgBuilder() + self.akg_builder = AkgBuilder("GPU") def handle(self): """ @@ -36,7 +36,7 @@ class GpuMessager(Messager): """ arg = self.get_message() if "AKG" in arg: - self.akg_builder.handle(self, arg, "GPU") + self.akg_builder.handle(self, arg) else: self.send_ack(False) self.exit() @@ -45,6 +45,7 @@ class GpuMessager(Messager): get_logger().info("[TRACE] GPU Messager Exit...") exit() + if __name__ == '__main__': warnings.simplefilter("ignore") if len(sys.argv) != 3: diff --git a/mindspore/ccsrc/backend/kernel_compiler/akg/akg_kernel_build.cc b/mindspore/ccsrc/backend/kernel_compiler/akg/akg_kernel_build.cc index 660e4a308e..8b047f153a 100644 --- a/mindspore/ccsrc/backend/kernel_compiler/akg/akg_kernel_build.cc +++ b/mindspore/ccsrc/backend/kernel_compiler/akg/akg_kernel_build.cc @@ -23,8 +23,10 @@ #include #include #include +#include "nlohmann/json.hpp" #include "ir/dtype.h" #include "ir/func_graph.h" +#include "utils/context/graph_kernel_flags.h" #include "backend/kernel_compiler/common_utils.h" #include "backend/kernel_compiler/akg/akg_kernel_json_generator.h" #include "backend/kernel_compiler/akg/akg_kernel_attrs_process.h" @@ -108,6 +110,11 @@ bool AkgKernelBuilder::AkgOpParallelBuild(const std::vector &build MS_LOG(ERROR) << "Akg start failed."; return false; } + auto attrs = CollectBuildAttrs(); + if (!attrs.empty() && !client->AkgSendAttr(attrs)) { + MS_LOG(ERROR) << "Akg send attr failed."; + return false; + } if (!client->AkgSendData(jsons)) { MS_LOG(ERROR) << "Akg send data failed."; return false; @@ -134,7 +141,9 @@ bool AkgKernelBuilder::AkgKernelParallelBuild(const std::vector &anf std::vector json_and_node; for (const auto &anf_node : anf_nodes) { MS_EXCEPTION_IF_NULL(anf_node); - AkgKernelJsonGenerator akg_kernel_json_generator; + DumpOption option; + option.get_compute_capability = true; + AkgKernelJsonGenerator akg_kernel_json_generator(option); auto cnode = anf_node->cast(); MS_EXCEPTION_IF_NULL(cnode); if (AnfAlgo::IsGraphKernel(cnode)) { @@ -146,14 +155,13 @@ bool AkgKernelBuilder::AkgKernelParallelBuild(const std::vector &anf func_graph->set_manager(mng); } std::vector node_list, input_list, output_list; - MS_LOG(INFO) << "Akg start compile composite op[" << anf_node->fullname_with_scope() << "]"; GetValidKernelNodes(func_graph, &node_list, &input_list, &output_list); if (!akg_kernel_json_generator.CollectFusedJson(node_list, input_list, output_list)) { - MS_EXCEPTION(UnknownError) << "Akg build failed composite op[" << anf_node->fullname_with_scope() << "]."; + MS_EXCEPTION(UnknownError) << "Collect op info failed. op[" << anf_node->fullname_with_scope() << "]."; } } else { if (!akg_kernel_json_generator.CollectJson(anf_node)) { - MS_EXCEPTION(UnknownError) << "Akg build failed basic op[" << anf_node->fullname_with_scope() << "]."; + MS_EXCEPTION(UnknownError) << "Collect op info failed. op[" << anf_node->fullname_with_scope() << "]."; } } json_and_node.push_back({akg_kernel_json_generator, anf_node}); @@ -167,6 +175,7 @@ bool AkgKernelBuilder::AkgKernelParallelBuild(const std::vector &anf struct timeval start_time, end_time; (void)gettimeofday(&start_time, nullptr); + MS_LOG(INFO) << "Akg start parallel build. kernel count: " << json_and_node.size(); bool res = AkgOpParallelBuild(json_and_node); if (!res) { MS_LOG(ERROR) << "Akg build kernel failed."; @@ -179,5 +188,17 @@ bool AkgKernelBuilder::AkgKernelParallelBuild(const std::vector &anf MS_LOG(INFO) << "Akg kernel build time: " << cost << " us."; return true; } + +std::string AkgKernelBuilder::CollectBuildAttrs() { + auto &flags = context::GraphKernelFlags::GetInstance(); + nlohmann::json attrs; + if (flags.online_tuning > 0) { + attrs["online_tuning"] = flags.online_tuning; + } + if (!flags.repository_path.empty()) { + attrs["repository_path"] = flags.repository_path; + } + return attrs.empty() ? "" : attrs.dump(); +} } // namespace kernel } // namespace mindspore diff --git a/mindspore/ccsrc/backend/kernel_compiler/akg/akg_kernel_build.h b/mindspore/ccsrc/backend/kernel_compiler/akg/akg_kernel_build.h index 81870dd0d4..c0012ece6f 100644 --- a/mindspore/ccsrc/backend/kernel_compiler/akg/akg_kernel_build.h +++ b/mindspore/ccsrc/backend/kernel_compiler/akg/akg_kernel_build.h @@ -49,6 +49,7 @@ class AkgKernelBuilder { bool HandleRepeatNodes(); bool AkgOpParallelBuild(const std::vector &build_args); std::vector repeat_nodes_; + std::string CollectBuildAttrs(); }; } // namespace kernel } // namespace mindspore diff --git a/mindspore/ccsrc/backend/kernel_compiler/akg/akg_kernel_json_generator.cc b/mindspore/ccsrc/backend/kernel_compiler/akg/akg_kernel_json_generator.cc index 41300aac50..39b2445bbf 100644 --- a/mindspore/ccsrc/backend/kernel_compiler/akg/akg_kernel_json_generator.cc +++ b/mindspore/ccsrc/backend/kernel_compiler/akg/akg_kernel_json_generator.cc @@ -22,6 +22,9 @@ #include #include #include +#if ENABLE_GPU +#include +#endif #include "backend/kernel_compiler/akg/akg_kernel_attrs_process.h" #include "backend/kernel_compiler/common_utils.h" #include "backend/kernel_compiler/oplib/oplib.h" @@ -549,6 +552,9 @@ bool AkgKernelJsonGenerator::CollectJson(const AnfNodePtr &anf_node, nlohmann::j (*kernel_json)[kJsonKeyPlatform] = "AKG"; (*kernel_json)[kJsonKeyProcess] = GetStrProcessorFromContext(); // GetProcessorStr(anf_node); (*kernel_json)[kJsonKeyComposite] = false; + if (dump_option_.get_compute_capability) { + (*kernel_json)[kJsonKeyComputeCapability] = ComputeCapability::Get(); + } if (!GetIOSize(*kernel_json, &input_size_list_, &output_size_list_)) { MS_LOG(ERROR) << "Cal mem size failed."; @@ -638,6 +644,9 @@ bool AkgKernelJsonGenerator::CollectFusedJson(const std::vector &anf (*kernel_json)[kJsonKeyProcess] = GetStrProcessorFromContext(); (*kernel_json)[kJsonKeyComposite] = true; (*kernel_json)[kJsonKeyCompositeGraph] = fg->ToString(); + if (dump_option_.get_compute_capability) { + (*kernel_json)[kJsonKeyComputeCapability] = ComputeCapability::Get(); + } GenStitchJson(anf_nodes, &node_json_map, kernel_json); @@ -837,5 +846,23 @@ bool AkgKernelJsonGenerator::CollectFusedJson(const std::vector &anf kernel_json_ = nlohmann::json(); return CollectFusedJson(anf_nodes, input_list, output_list, &kernel_json_); } + +void ComputeCapability::GetComputeCapability() { +#if ENABLE_GPU + int a, b; + auto ret = cuDeviceGetAttribute(&a, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, 0); + if (ret != CUDA_SUCCESS) { + MS_LOG(WARNING) << "Get CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR fail, ret=" << ret; + return; + } + ret = cuDeviceGetAttribute(&b, CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, 0); + if (ret != CUDA_SUCCESS) { + MS_LOG(WARNING) << "Get CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR fail, ret=" << ret; + return; + } + this->compute_capability_ = std::to_string(a) + "." + std::to_string(b); +#endif + return; +} } // namespace kernel } // namespace mindspore diff --git a/mindspore/ccsrc/backend/kernel_compiler/akg/akg_kernel_json_generator.h b/mindspore/ccsrc/backend/kernel_compiler/akg/akg_kernel_json_generator.h index 193ac54dcb..319f23028d 100644 --- a/mindspore/ccsrc/backend/kernel_compiler/akg/akg_kernel_json_generator.h +++ b/mindspore/ccsrc/backend/kernel_compiler/akg/akg_kernel_json_generator.h @@ -58,6 +58,7 @@ constexpr auto kJsonKeyRecomputeOps = "recompute_ops"; constexpr auto kJsonKeyBufferStitch = "buffer_stitch"; constexpr auto kJsonKeyStitchOp = "stitch_op"; constexpr auto kJsonKeyStitchAtomicOp = "stitch_atomic_op"; +constexpr auto kJsonKeyComputeCapability = "compute_capability"; constexpr auto kAttrInputNames = "input_names"; @@ -66,6 +67,23 @@ struct DumpOption { bool is_before_select_kernel = false; bool save_ptr_address = false; bool extract_opinfo_from_anfnode = false; + bool get_compute_capability = false; +}; + +class ComputeCapability { + public: + static const std::string &Get() { + static std::unique_ptr instance = nullptr; + if (instance == nullptr) { + instance = std::make_unique(); + instance->GetComputeCapability(); + } + return instance->compute_capability_; + } + + private: + void GetComputeCapability(); + std::string compute_capability_; }; class AkgKernelJsonGenerator { diff --git a/mindspore/ccsrc/backend/kernel_compiler/akg/gpu/akg_gpu_kernel_build.cc b/mindspore/ccsrc/backend/kernel_compiler/akg/gpu/akg_gpu_kernel_build.cc index 5f07c0641b..d39e75e291 100644 --- a/mindspore/ccsrc/backend/kernel_compiler/akg/gpu/akg_gpu_kernel_build.cc +++ b/mindspore/ccsrc/backend/kernel_compiler/akg/gpu/akg_gpu_kernel_build.cc @@ -27,7 +27,6 @@ namespace mindspore { namespace kernel { constexpr int32_t ARGS_SIZE = 1; -constexpr auto kCompileWithJsonFunc = "compilewithjson"; KernelPackPtr AkgGpuKernelBuilder::AkgSearchCache(const std::string &kernel_name) { return SearchCache(kernel_name, kProcessorCuda); diff --git a/mindspore/ccsrc/backend/session/kernel_build_client.cc b/mindspore/ccsrc/backend/session/kernel_build_client.cc index a95c0df7bc..097c4564a5 100644 --- a/mindspore/ccsrc/backend/session/kernel_build_client.cc +++ b/mindspore/ccsrc/backend/session/kernel_build_client.cc @@ -50,6 +50,20 @@ bool KernelBuildClient::AkgStart(int process_num, int wait_time) { return true; } +bool KernelBuildClient::AkgSendAttr(const std::string &attr) { + auto res = SendRequest(kAkgAttr); + if (res != kAck) { + MS_LOG(ERROR) << "AKG/ATTR failed, res: " << res; + return false; + } + res = SendRequest(attr); + if (res != kAck) { + MS_LOG(ERROR) << "AKG/ATTR.. responds failed, res: " << res << ", when sending [" << attr << "]"; + return false; + } + return true; +} + bool KernelBuildClient::AkgSendData(const std::vector &jsons) { auto res = SendRequest(kAkgData); if (res != kAck) { diff --git a/mindspore/ccsrc/backend/session/kernel_build_client.h b/mindspore/ccsrc/backend/session/kernel_build_client.h index 44cfe43a34..5101d035c1 100644 --- a/mindspore/ccsrc/backend/session/kernel_build_client.h +++ b/mindspore/ccsrc/backend/session/kernel_build_client.h @@ -54,6 +54,7 @@ class KernelBuildClient { constexpr inline static auto kFinish = "FINISH"; constexpr inline static auto kAkgStart = "AKG/START"; constexpr inline static auto kAkgData = "AKG/DATA"; + constexpr inline static auto kAkgAttr = "AKG/ATTR"; constexpr inline static auto kAkgWait = "AKG/WAIT"; // Receive the response from server constexpr inline static auto kAck = "ACK"; @@ -129,6 +130,7 @@ class KernelBuildClient { // Run AKG building. bool AkgStart(int process_num, int wait_time); + bool AkgSendAttr(const std::string &attr); bool AkgSendData(const std::vector &jsons); bool AkgWait(); diff --git a/mindspore/ccsrc/utils/context/graph_kernel_flags.cc b/mindspore/ccsrc/utils/context/graph_kernel_flags.cc index 8928168631..97cdf4aab2 100644 --- a/mindspore/ccsrc/utils/context/graph_kernel_flags.cc +++ b/mindspore/ccsrc/utils/context/graph_kernel_flags.cc @@ -185,8 +185,10 @@ void GraphKernelFlags::RegisterFlags(std::map *flag_ma reg.AddFlag("enable_parallel_fusion", &enable_parallel_fusion, opt_level == OptLevel_3); // Integer flags - reg.AddFlag("auto_tune", &auto_tune); - reg.AddFlag("cluster_limit", &cluster_limit); + reg.AddFlag("online_tuning", &online_tuning); + + // String flags + reg.AddFlag("repository_path", &repository_path); // String list flags reg.AddFlag("enable_expand_ops", &enable_expand_ops); @@ -208,8 +210,9 @@ std::string GraphKernelFlags::DumpAllFlags() const { json["enable_parallel_fusion"] = enable_parallel_fusion; json["opt_level"] = opt_level; - json["auto_tune"] = auto_tune; - json["cluster_limit"] = cluster_limit; + json["online_tuning"] = online_tuning; + + json["repository_path"] = repository_path; json["enable_expand_ops"] = enable_expand_ops; json["enable_expand_ops_only"] = enable_expand_ops_only; diff --git a/mindspore/ccsrc/utils/context/graph_kernel_flags.h b/mindspore/ccsrc/utils/context/graph_kernel_flags.h index d7a1f02501..81dfbdd9fa 100644 --- a/mindspore/ccsrc/utils/context/graph_kernel_flags.h +++ b/mindspore/ccsrc/utils/context/graph_kernel_flags.h @@ -91,14 +91,16 @@ class GraphKernelFlags { unsigned int opt_level; // defaults 0 or 2 /** - * auto_tune, unsupported now. + * Online tuning level, value from 0 to 3. + * 0: Disable online tuning + * 1-3: The higher level, the larger tuning space, and the more time it takes. */ - unsigned int auto_tune{0}; + unsigned int online_tuning{0}; /** - * cluster_limit, unsupported now. + * AKG's operator repository file path. */ - unsigned int cluster_limit{0}; + std::string repository_path; /** * Additional expanding operators (case sensitive). diff --git a/tests/ut/cpp/stub/runtime/cuda.cc b/tests/ut/cpp/stub/runtime/cuda.cc index a9595ec6e7..9b368d4f4b 100644 --- a/tests/ut/cpp/stub/runtime/cuda.cc +++ b/tests/ut/cpp/stub/runtime/cuda.cc @@ -17,8 +17,10 @@ CUresult cuModuleLoadData(CUmodule *module, const void *image) { return CUDA_SUCCESS; } -CUresult cuModuleLoadDataEx(CUmodule *module, const void *image, unsigned int numOptions, - CUjit_option *options, void **optionValues) { return CUDA_SUCCESS; } +CUresult cuModuleLoadDataEx(CUmodule *module, const void *image, unsigned int numOptions, CUjit_option *options, + void **optionValues) { + return CUDA_SUCCESS; +} CUresult cuModuleGetFunction(CUfunction *hfunc, CUmodule hmod, const char *name) { return CUDA_SUCCESS; } @@ -31,3 +33,8 @@ CUresult cuLaunchKernel(CUfunction f, unsigned int gridDimX, unsigned int gridDi CUresult cuModuleUnload(CUmodule hmod) { return CUDA_SUCCESS; } CUresult cuGetErrorName(CUresult error, const char **pStr) { return CUDA_SUCCESS; } + +CUresult cuDeviceGetAttribute(int *pi, CUdevice_attribute attrib, int dev) { + *pi = 0; + return CUDA_SUCCESS; +} diff --git a/tests/ut/cpp/stub/runtime/cuda.h b/tests/ut/cpp/stub/runtime/cuda.h index dc4d326acc..318b6dc4f2 100644 --- a/tests/ut/cpp/stub/runtime/cuda.h +++ b/tests/ut/cpp/stub/runtime/cuda.h @@ -22,31 +22,35 @@ typedef enum cudaError_enum { CUDA_ERROR_DEINITIALIZED = 2, } CUresult; -typedef enum CUjit_option_enum -{ - CU_JIT_MAX_REGISTERS = 0, - CU_JIT_THREADS_PER_BLOCK, - CU_JIT_WALL_TIME, - CU_JIT_INFO_LOG_BUFFER, - CU_JIT_INFO_LOG_BUFFER_SIZE_BYTES, - CU_JIT_ERROR_LOG_BUFFER, - CU_JIT_ERROR_LOG_BUFFER_SIZE_BYTES, - CU_JIT_OPTIMIZATION_LEVEL, - CU_JIT_TARGET_FROM_CUCONTEXT, - CU_JIT_TARGET, - CU_JIT_FALLBACK_STRATEGY, - CU_JIT_GENERATE_DEBUG_INFO, - CU_JIT_LOG_VERBOSE, - CU_JIT_GENERATE_LINE_INFO, - CU_JIT_CACHE_MODE, - CU_JIT_NEW_SM3X_OPT, - CU_JIT_FAST_COMPILE, - CU_JIT_GLOBAL_SYMBOL_NAMES, - CU_JIT_GLOBAL_SYMBOL_ADDRESSES, - CU_JIT_GLOBAL_SYMBOL_COUNT, - CU_JIT_NUM_OPTIONS +typedef enum CUjit_option_enum { + CU_JIT_MAX_REGISTERS = 0, + CU_JIT_THREADS_PER_BLOCK, + CU_JIT_WALL_TIME, + CU_JIT_INFO_LOG_BUFFER, + CU_JIT_INFO_LOG_BUFFER_SIZE_BYTES, + CU_JIT_ERROR_LOG_BUFFER, + CU_JIT_ERROR_LOG_BUFFER_SIZE_BYTES, + CU_JIT_OPTIMIZATION_LEVEL, + CU_JIT_TARGET_FROM_CUCONTEXT, + CU_JIT_TARGET, + CU_JIT_FALLBACK_STRATEGY, + CU_JIT_GENERATE_DEBUG_INFO, + CU_JIT_LOG_VERBOSE, + CU_JIT_GENERATE_LINE_INFO, + CU_JIT_CACHE_MODE, + CU_JIT_NEW_SM3X_OPT, + CU_JIT_FAST_COMPILE, + CU_JIT_GLOBAL_SYMBOL_NAMES, + CU_JIT_GLOBAL_SYMBOL_ADDRESSES, + CU_JIT_GLOBAL_SYMBOL_COUNT, + CU_JIT_NUM_OPTIONS } CUjit_option; +typedef enum CUdevice_attribute_enum { + CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MAJOR, + CU_DEVICE_ATTRIBUTE_COMPUTE_CAPABILITY_MINOR, +} CUdevice_attribute; + struct CUctx_st { int arch; }; @@ -65,13 +69,14 @@ typedef struct CUmod_st *CUmodule; typedef struct CUfunc_st *CUfunction; typedef struct CUstream_st *CUstream; - CUresult cuModuleLoadData(CUmodule *module, const void *image); -CUresult cuModuleLoadDataEx(CUmodule *module, const void *image, unsigned int numOptions, CUjit_option *options, void **optionValues); +CUresult cuModuleLoadDataEx(CUmodule *module, const void *image, unsigned int numOptions, CUjit_option *options, + void **optionValues); CUresult cuModuleGetFunction(CUfunction *hfunc, CUmodule hmod, const char *name); CUresult cuLaunchKernel(CUfunction f, unsigned int gridDimX, unsigned int gridDimY, unsigned int gridDimZ, unsigned int blockDimX, unsigned int blockDimY, unsigned int blockDimZ, unsigned int sharedMemBytes, CUstream hStream, void **kernelParams, void **extra); CUresult cuModuleUnload(CUmodule hmod); CUresult cuGetErrorName(CUresult error, const char **pStr); +CUresult cuDeviceGetAttribute(int *pi, CUdevice_attribute attrib, int dev); #endif // TESTS_UT_STUB_RUNTIME_INCLUDE_CUDA_H_