like20
/
mindspore2022
forked from huawei/mindspore2022
0
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

!20174 Send compilation attrs to akg 

Merge pull request !20174 from DeshiChen/0708_compilewithjson
This commit is contained in:
i-robot 2021-07-19 08:13:36 +00:00 committed by Gitee
parent 924863b655 3ab53dd26d
commit fd72de08d9
17 changed files with 179 additions and 71 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
akg

@ -1 +1 @@
Subproject commit 97dc7e96c2ffedf2e6e38310a903ffa205a6e656
Subproject commit 1e6b226a0417d23d2d0a2333d5e80f13fe9e8d0f

26
mindspore/_extends/parallel_compile/akg_compiler/akg_process.py
View File

@ -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'.")

9
mindspore/_extends/parallel_compile/akg_compiler/compiler.py
View File

@ -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])

25
mindspore/_extends/remote/kernel_build_server.py
View File

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

4
mindspore/_extends/remote/kernel_build_server_ascend.py
View File

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

5
mindspore/_extends/remote/kernel_build_server_gpu.py
View File

@ -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:

29
mindspore/ccsrc/backend/kernel_compiler/akg/akg_kernel_build.cc
View File

@ -23,8 +23,10 @@
#include <unordered_set>
#include <utility>
#include <vector>
#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<JsonNodePair> &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<AnfNodePtr> &anf
std::vector<JsonNodePair> 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<CNodePtr>();
MS_EXCEPTION_IF_NULL(cnode);
if (AnfAlgo::IsGraphKernel(cnode)) {
@ -146,14 +155,13 @@ bool AkgKernelBuilder::AkgKernelParallelBuild(const std::vector<AnfNodePtr> &anf
func_graph->set_manager(mng);
}
std::vector<AnfNodePtr> 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<AnfNodePtr> &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<AnfNodePtr> &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

1
mindspore/ccsrc/backend/kernel_compiler/akg/akg_kernel_build.h
View File

@ -49,6 +49,7 @@ class AkgKernelBuilder {
bool HandleRepeatNodes();
bool AkgOpParallelBuild(const std::vector<JsonNodePair> &build_args);
std::vector<JsonNodePair> repeat_nodes_;
std::string CollectBuildAttrs();
};
} // namespace kernel
} // namespace mindspore

27
mindspore/ccsrc/backend/kernel_compiler/akg/akg_kernel_json_generator.cc
View File

@ -22,6 +22,9 @@
#include <set>
#include <sstream>
#include <tuple>
#if ENABLE_GPU
#include <cuda.h>
#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<AnfNodePtr> &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<AnfNodePtr> &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

18
mindspore/ccsrc/backend/kernel_compiler/akg/akg_kernel_json_generator.h
View File

@ -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<ComputeCapability> instance = nullptr;
if (instance == nullptr) {
instance = std::make_unique<ComputeCapability>();
instance->GetComputeCapability();
}
return instance->compute_capability_;
}
private:
void GetComputeCapability();
std::string compute_capability_;
};
class AkgKernelJsonGenerator {

1
mindspore/ccsrc/backend/kernel_compiler/akg/gpu/akg_gpu_kernel_build.cc
View File

@ -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);

14
mindspore/ccsrc/backend/session/kernel_build_client.cc
View File

@ -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<std::string> &jsons) {
auto res = SendRequest(kAkgData);
if (res != kAck) {

2
mindspore/ccsrc/backend/session/kernel_build_client.h
View File

@ -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<std::string> &jsons);
bool AkgWait();

11
mindspore/ccsrc/utils/context/graph_kernel_flags.cc
View File

@ -185,8 +185,10 @@ void GraphKernelFlags::RegisterFlags(std::map<std::string, std::string> *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;

10
mindspore/ccsrc/utils/context/graph_kernel_flags.h
View File

@ -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).

11
tests/ut/cpp/stub/runtime/cuda.cc
View File

@ -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;
}

55
tests/ut/cpp/stub/runtime/cuda.h
View File

@ -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_