()
| 30 | |
| 31 | @pytest.mark.gpu |
| 32 | def test_cuda_multi_lib(): |
| 33 | pytest.importorskip("cloudpickle") |
| 34 | |
| 35 | # test combining two system lib together |
| 36 | # each contains a fatbin component in cuda |
| 37 | dev = tvm.cuda(0) |
| 38 | for device in ["llvm", "cuda"]: |
| 39 | if not tvm.testing.device_enabled(device): |
| 40 | print(f"skip because {device} is not enabled...") |
| 41 | return |
| 42 | |
| 43 | @tvm.script.ir_module |
| 44 | class ModA: |
| 45 | I.module_attrs({"system_lib_prefix": "modA_"}) |
| 46 | |
| 47 | @T.prim_func(s_tir=True) |
| 48 | def my_inplace_update(x: T.Buffer((12), "float32")) -> None: |
| 49 | T.func_attr({"global_symbol": "modA_my_inplace_update"}) |
| 50 | for bx in T.thread_binding(T.int64(1), thread="blockIdx.x"): |
| 51 | for tx in T.thread_binding(T.int64(12), thread="threadIdx.x"): |
| 52 | x[tx] = x[tx] + 1 |
| 53 | |
| 54 | @tvm.script.ir_module |
| 55 | class ModB: |
| 56 | I.module_attrs({"system_lib_prefix": "modB_"}) |
| 57 | |
| 58 | @T.prim_func(s_tir=True) |
| 59 | def my_inplace_update(x: T.Buffer((12), "float32")) -> None: |
| 60 | T.func_attr({"global_symbol": "modB_my_inplace_update"}) |
| 61 | for bx in T.thread_binding(T.int64(1), thread="blockIdx.x"): |
| 62 | for tx in T.thread_binding(T.int64(12), thread="threadIdx.x"): |
| 63 | x[tx] = x[tx] + 2 |
| 64 | |
| 65 | temp = utils.tempdir() |
| 66 | target = tvm.target.Target("cuda", host="llvm") |
| 67 | libA = tvm.compile(ModA, target=target) |
| 68 | libB = tvm.compile(ModB, target=target) |
| 69 | |
| 70 | pathA = temp.relpath("libA.tar") |
| 71 | pathB = temp.relpath("libB.tar") |
| 72 | pathAll = temp.relpath("libAll.a") |
| 73 | |
| 74 | path_dso = temp.relpath("mylib.so") |
| 75 | libA.export_library(pathA, fcompile=tar.tar) |
| 76 | libB.export_library(pathB, fcompile=tar.tar) |
| 77 | cc.create_staticlib(pathAll, [pathA, pathB]) |
| 78 | # package two static libs together |
| 79 | cc.create_shared(path_dso, ["-Wl,--whole-archive", pathAll, "-Wl,--no-whole-archive"]) |
| 80 | |
| 81 | def popen_check(): |
| 82 | # Load dll, will trigger system library registration |
| 83 | ctypes.CDLL(path_dso) |
| 84 | # Load the system wide library |
| 85 | dev = tvm.cuda() |
| 86 | a_np = np.random.uniform(size=12).astype("float32") |
| 87 | a_nd = tvm.runtime.tensor(a_np, dev) |
| 88 | b_nd = tvm.runtime.tensor(a_np, dev) |
| 89 | syslibA = tvm.runtime.system_lib("modA_") |
nothing calls this directly
no test coverage detected
searching dependent graphs…