TVM_0.20

[TVM_0.20] 04 Runtime and Device Interaction Part2

Runtime 상세 파트 2

Introduction to Module Serialization

TVM 런타임 모듈 배포 시, CPU/GPU에 상관 없이 하나의 Dynamic share library만 배포하면 된다. 이는 TVM의 module serialization 메커니즘 덕택이다. 이번장에서는 TVM module serialization 표준 포맷과 상세 구현에 대해 다룬다.

Serialization

Serialization을 위한 기본 API는 tvm.module.Module의 export_library 이다. 이 API에서 내부적으로 하기 스텝을 수행한다.

  1. 모든 DSO(Dynamic sahred object(*.so)?) module을 수집(LLVM 또는 C module)
  2. DSO 모듈이 있으면 save 함수를 호출하여 이를 파일에 저장
  3. CUDA나 OpenCL 같은 imported moudle을 가지고 있는지 확인하고 imported moudle이 있을 경우 devc.o 또는 dev.cc 을 생성한다. 그후 module serialization을 위해 _PackImportsToLLVM 또는 _PackImportsToC을 호출한다.
    • imported moudle의 type은 제한이 없다.
    • imported moudle의 binary blob data를 하나의 dynamic shared library에 임베딩할 수 있다.
  4. dynamic shared library를 만들기 위해 ‘_cc.create_shared’을 invoke하는 fcompile을 호출한다.
    (Note1) C파일의 경우 컴파일 후 DSO module에 링킹한다.
    (Note1) _PackImportsToLLVM 또는 _PackImportsToC의 사용 여부는 TVM에 LLVM을 활성화 하였는지에 달려있다. 두 함수는 같은 목적을 가진다.

Under the Hood of Serialization and Format Standard

serialization 작업은 _PackImportsToLLVM 또는 _PackImportsToC에서 이루어지며 두 함수는 serialization 작업을 위해 SerializeModule을 호출한다. TVM은 SerializeModule에서 ModuleSerializer라는 헬퍼 클래스가 구성되어 있다. 이 헬퍼클래스는 module을 인자로 받아 모듈 index생성 같은 초기화 작업을 진행한다. 그 후 module의 SerializeModule을 사용할 수 있다.

  • (역자주) Helper class : 해당 애플리케이션이나 해당 클래스의 주요 목적이 아닌 일부 기능을 제공하는 데 사용되는 클래

ModuleSerializer 헬퍼 클래스

다음의 코드는 ModuleSerializer를 구성하기 위해 사용된다.

1
2
3
4
5
6
7
8

explicit ModuleSerializer(runtime::Module mod) : mod_(mod) {
  Init();
}
private:
void Init() {
  CreateModuleIndex();
  CreateImportTree();
}
  1. CreateModuleIndex()는 DFS를 이용하여 module import relationship을 조사하고 index를 생성한다.
    • root module의 index는 0이다. (본문 예시 참조)
  2. 인덱스 생성 후 CreateImportTree()는 import tree를 생성한다.
    • import tree는 exported library를 로드할 때 module import relationship을 복원하는데 사용된다.
    • TVM은 import tree를 저장하기 위해 CSR format을 사용한다. 각 Row는 parent index를, child index들은 Row의 children index를 나태낸다. import_tree_row_ptr_ 과 import_tree_child_indices_에서 이를 확인할 수 있다.

SerializeModule

헬퍼 클래스에 의한 초기화 후 SerializeModule을 이용하여 직렬화를 수행 할 수 있다. SerializeModule 함수에서 serialization format을 하기와 같이 사용한다.

1
2
3
4
5
6
7
8

binary_blob_size
binary_blob_type_key
binary_blob_logic
binary_blob_type_key
binary_blob_logic
...
_import_tree
_import_tree_logic
  • binary_blob_size : 직렬화 단계에서 수행할 blob의 갯수 (본문의 예제에서는 3 : LLVM module, CUDA module, _import_tree)
  • binary_blob_type_key : 모듈 blob의 type (LLVM / C module은 _lib, CUDA는 cuda, module->type_key()로 값을 얻을 수 있음)
  • binary_blob_logic : blob의 로직 핸들링 (대부분의 blob은 SaveToBinary함수(CUDA나 OpenCL), 그러나 LLVM/C module은 이것이 DSO모듈임을 나타내기 위해 _lib이라고만 적어놓음)
    • [Note1] virture function인 SaveToBinary를 구현해야하는지 여부는 모듈이 어떻게 사용되는지에 따라 다름
    • [Note2] dynamic shared library를 로드할 때 필요한 데이타가 있다면 SaveToBinary를 구현해야함, 예를 들어 CUDA의 경우 dynamic shared library를 로드할 때 GPU Driver에게 바이너리 데이터를 전달 해야하므로 상기 가상함수를 구현해야함
    • [Note3] DSO 같은 host module은 dynamic shared library를 로드 시 별도의 정보가 필요 없으므로 가상함수를 구현할 필요는 없다. 그러나 미래에 메타 데이터 기록이 필요할 경우가 있다면 가상 함수를 구현하면 된다.
  • _import_tree : import_tree를 저장(루트 모듈이 DSO 하나만 있고, 다른 모듈(import된 모듈)이 없다면 _import_tree는 필요 없음=기록할 트리가 없으므로)
  • import_tree_logic : 그냥 스트림안에 import_tree_row_ptr_ 과 import_tree_child_indices_를 적어 놓는다.
  • 마지막으로 상기 내용을 runtime::symbol::tvm_ffi_library_bin 심볼에 패킹 함

역자주

전체 직렬화 흐름 (export_library)의 흐름은 다음과 같다

  1. DSO 모듈 수집(예: LLVM, C 소스 모듈)
  2. DSO 모듈 저장(.o 또는 .cc로 저장)
  3. 임포트된 모듈(CUDA, OpenCL 등) 확인 및 직렬화(_PackImportsToLLVM 또는 _PackImportsToC 함수로 처리)
  4. 최종 컴파일 및 공유 라이브러리 생성(fcompile → _cc.create_shared() 사용)
 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20

#CPU + CUDA 모듈 직렬화 예시

# Step 1: 연산 정의
n = te.var("n")
A = te.placeholder((n,), name="A")
B = te.compute((n,), lambda i: A[i] + 1, name="B")

# Step 2: 스케줄 정의
s_cpu = te.create_schedule(B.op)
s_cuda = te.create_schedule(B.op)

# Step 3: CPU, CUDA 모듈 각각 빌드
cpu_mod = tvm.build(s_cpu, [A, B], target="llvm", name="cpu_add")
cuda_mod = tvm.build(s_cuda, [A, B], target="cuda", name="cuda_add")

# Step 4: CPU 모듈에 CUDA 모듈을 import
cpu_mod.import_module(cuda_mod)

# Step 5: 하나의 라이브러리로 저장
cpu_mod.export_library("fused_module.so")

Deserialization

Deserialization의 시작 API는 tvm.runtime.load이다. 이 함수는 _LoadFromFile을 호출한다.(이는 Module::LoadFromFile 이다) 예를 들어 배포 바이너리 deploy.so는 dso_library.cc의 module.loadfile_so을 호출한다. 하기 코드는 주요 구조를 나타낸다

 1
 2
 3
 4
 5
 6
 7
 8
 9
10

// Load the imported modules
const char* library_bin = reinterpret_cast<const char*>(
   lib->GetSymbol(runtime::symbol::tvm_ffi_library_bin));
Module root_mod;
if (library_bin != nullptr) {
   root_mod = ProcessLibraryBin(library_bin, lib);
} else {
   // Only have one single DSO Module
   root_mod = Module(n);
}

deserialization동안 runtime::symbol::tvm_ffi_library_bin를 검사한다.(이전에 blob 정보를 여기다 packing하였다). 만약 runtime::symbol::tvm_ffi_library_bin이 있다면 ProcessLibraryBin을 호출한다. 이 함수의 구조는 다음과 같다.

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17

READ(blob_size)
READ(blob_type_key)
for (size_t i = 0; i < blob_size; i++) {
    if (blob_type_key == "_lib") {
      // construct dso module using lib
    } else if (blob_type_key == "_import_tree") {
      // READ(_import_tree_row_ptr)
      // READ(_import_tree_child_indices)
    } else {
      // call module.loadbinary_blob_type_key, such as module.loadbinary_cuda
      // to restore.
    }
}
// Using _import_tree_row_ptr and _import_tree_child_indices to
// restore module import relationship. The first module is the
// root module according to our invariance as said before.
return root_module;

이후 root부터 심볼 탐색이 가능하도록 ctx_address를 root_module로 세팅함

Device/Target Interactions

본 문서는 TVM framework가 특정 device API와 어떻게 상호작용하는지 알기를 원하는 개발자나, 새로운 API 또는 새로운 하드웨어에 대한 지원을 구현하기 위한 사람을 위한 것이다. 새로운 런타임 환경을 위해 구현해야 하는 세 가지 주요 측면이 있다.

  1. Device API class : 특정 디바이스 또는 이것과 상호작용하기 위한 API를 제공한다. 이것은 디바이스 파라메터(메모리 사이즈, 스레드 갯수 등) 쿼링 또는 간단한 동작(호스트- 디바이스간 메모리 복사 등) 수행을 위한 공통의 인터페이스를 정의한다.
  2. Target Class : 함수가 실행될 장치에 대한 설명을 담고 있다. 이것은 target code generators과 optimization pass에서 사용된다.
  3. Target code generation : IRModule에서 1개이상의 PackedFunc으로 구성된 Module을 구성한다.(Module/PackedFunc는 part1 참조)

DeviceAPI

DeviceAPI는 특정 하드웨어 장치 API에 대한 handle을 나타냄. (CUDADeviceAPI 은 CUDA 프레임워크를 통한 모든 상호작용을 핸들링함)
대부분의 DeviceAPI는 접근해야할 디바이스를 지정하기 위해 device_id를 인자로 받음. Python애서는 tvm.runtime.device() 함수를 이용하여 device_id에 접근 가능함 (예를 들어 tvm.runtime.device('cuda',0)은 CUDA API를 통해 물리적 디바이스 0에 접근 가능함).

  • Attribute queries : GetAttr는 쿼리 된 다른 device-specific parameter들에 접근을 허락함(장치 이름, 스레드 갯수 등) 쿼리된 파라메터는 device_api.h 내 enum DeviceAttrKind에 정의 되어 있음. 모든 쿼리 가능한 파라메가 모든 장치에서 지원되는 것은 아닙니다. 파라메터가 쿼리 되지 않았거나(Vulkan의 kMaxClockRate) 파라메터가 해당되지 않는 경우(CPU의 kWarpSize) 해당 쿼리는 nullptr를 리턴함
  • Setting active device : SetDevice는 특정 디바이스를 활성화 상태로 세팅함. 타켓 코드 제너레이터에 의해 생성된 PackedFunc이 디바이스 실행을 요청할 경우, 이것은 활상화 디바이스에서 실행되어야 함.
  • Memory management : 디바이스에서 메모리 할당 및 해제를 위한 유틸리티
    • Allocate data space : AllocDataSpace나 FreeDataSpace는 디바이스에서 메모리 공간을 할당 또는 해제함. 할당된 메모리는 연산자에게 입력 및 출력으로 제공될 수 있으며, 연산자 그래프의 주요 데이터 흐름을 구성. 이것은 호스트와 할당된 메모리 간의 데이터 전송이 가능해야 함. 리턴 value는 void*. 몇몇 구현은 메모리 주소를 리턴하는 반면 이것이 필수는 아니며, void*는 해당 주소를 생성한 장치 백엔드에서만 해석할 수 있는 opaque handle일 수 있다. void*는 CopyDataFromTo와 같은 다른 백엔드 전용 함수의 인수로 사용된다.
    • Allocate work space : AllocWorkspace와 FreeWorkspace는 디바이스에서 메모리 공간을 할당 또는 해제함. data space와는 다르게 이것은 operator definition 내부에서 intermediate values 저장을 위해 사용되며 host device와 데이터 전송을 필요로 하지 않는다. DeviceAPI subclass가 이 방법을 구현하지 않는다면 해당 DeviceAPI는 이에 대응되는 DataSpace 함수들을 호출한다.
    • Copy data : CopyDataFromTo는 데이터를 한곳에서 다른 곳으로 복사한다. 이 복사의 형태는 dev_from나 dev_to 파라메터에 의해 결정된다. 이의 구현은 CPU와 장치간 메모리 복사, 버퍼에서 단일 장치로 데이터 복사를 지원해야 한다. src/dst 위치가 CPU에 있는 경우 해당 void*은 memcpy로 전달할 수 있는 CPU 주소를 나타냄. src/dst 위치가 디바이스에 있는 경우 해당 void*은 이전에 AllocDataSpace 또는 AllocWorkspace에서 생성된 주소이다.
    • 이러한 복사는 특정 TVMStreamHandle에서 실행되도록 대기열에 저장된다. 그러나 CopyDataFromTo 호출이 완료된 후에도 CPU 버퍼가 유효하거나 액세스할 수 있다고 가정해서는 안 됩니다.
  • Execution stream management : 명령을 실행하는데 사용되는 병렬 실행 스트림인 TVMStreamHandle을 핸들링하는데 사용되는 유틸
    • Create stream : CreateStream과 FreeStream은 실행 스트림의 핸들을 할당/해제하여야 한다. 만약 디바이스가 command queue 사이즈가 1이면 CreateStream은 nullptr을 반환한다.
    • Set active stream : SetStream은 스트림을 활성화해야 한다. 타켓 코드 제너레이터에 의해 생성된 PackedFunc이 디바이스 실행을 요청할 경우, 그 작업은 활성화 된 스트림에 제출 되어야 한다.
    • Synchronize to CPU : StreamSync은 실행 스트림을 CPU와 동기화 해야한다. StreamSync 함수 호출은, 그 호출 이전에 제출된 모든 메모리 전송과 연산이 완료된 후에야 반환되어야 한다.
    • Synchronize between streams : SyncStreamFromTo 은 source/destination 스트림 간 synchronization barrier를 도입해야 한다. 즉 SyncStreamFromTo 명령을 통해 source 스트림이 완료될 때 까지 dest 스트림 진행을 멈출 수 있다.
    • (역자주) source 스트림은 먼저 실행된 스트림이고 dest은 나중에 실행된 스트림이다. 두 스트림간 데이터 동기화, 의존성 해결을 위해 한 스트림이 완료된 후 다음 스트림을 실행 할 수 있음

신규 DeviceAPI 등록 절차

새로운 디바이스가 TVM에서 사용되려면 신규 DeviceAPI를 등록해야 함

  1. 새로운 DeviceAPI를 인스턴스화하고 포인터를 반환하는 함수를 만듭니다
1
2
3
4

FooDeviceAPI* FooDeviceAPI::Global() {
  static FooDeviceAPI inst;
  return &inst;
}
  1. TVM registry에 함수를 등록합니다.
1
2

TVM_FFI_REGISTER_GLOBAL("device_api.foo").set_body_typed(FooDeviceAPI::Global);
// device_api.foo는 TVM registry에 등록된 이름
  1. 새로운 DeviceAPI를 위한 엔트리를 base.h 내부의 enum TVMDeviceExtType에 추가한다. 그 값은 DLDeviceType::kDLExtDev보다 크고 DeviceAPIManager::kMaxDeviceAPI 보다 작은 사용되지 않은 값이어야 한다.
  2. device_api.h 내부 DeviceName에 enum값을 string으로 변환하기 위한 case 문을 추가해라. 이 string 값은 TVM_FFI_REGISTER_GLOBAL에 주어진 이름과 같아야 한다.
  3. 새로운 enum값을 위해 엔트리들을 tvm.runtime.Device의 DEVICE_TYPE_TO_NAME과 DEVICE_NAME_TO_TYPE dictionaries에 추가하라.

역자해석

  1. DeviceAPI 클래스 구현
 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45

// foo_device_api.cc

#include <tvm/runtime/device_api.h>
#include <tvm/runtime/registry.h>
#include <iostream>

namespace tvm {
namespace runtime {

class FooDeviceAPI : public DeviceAPI {
 public:
  void SetDevice(DLDevice dev) override {
    std::cout << "Set Foo device: " << dev.device_id << std::endl;
  }

  void* AllocDataSpace(DLDevice dev, size_t nbytes, size_t alignment,
                       DLDataType type_hint) override {
    std::cout << "Alloc Foo device memory\n";
    return malloc(nbytes);
  }

  void FreeDataSpace(DLDevice dev, void* ptr) override {
    std::cout << "Free Foo device memory\n";
    free(ptr);
  }

  void CopyDataFromTo(const void* from, void* to, size_t size,
                      const DLDevice& dev_from, const DLDevice& dev_to,
                      DLDataType type_hint, TVMStreamHandle stream) override {
    std::cout << "Copy data on Foo device\n";
    memcpy(to, from, size);
  }

  void StreamSync(DLDevice dev, TVMStreamHandle stream) override {
    std::cout << "Stream sync on Foo device\n";
  }

  static FooDeviceAPI* Global() {
    static FooDeviceAPI inst;
    return &inst;
  }
};

}  // namespace runtime
}  // namespace tvm
  1. 레지스트리 등록
1
2

TVM_REGISTER_GLOBAL("device_api.foo")
.set_body_typed(tvm::runtime::FooDeviceAPI::Global);
  1. TVMDeviceExtType에 enum 추가 (include/tvm/runtime/c_runtime_api.h)
1
2
3
4
5

typedef enum {
  // 기존 값들 ...
  kDLExtDev = 12,   // 시작점
  kDLFoo = 13,      // 새로운 Foo 디바이스
} TVMDeviceExtType;
  1. DeviceName()에 enum -> 문자열 변환 추가 (device_api.h)
 1
 2
 3
 4
 5
 6
 7
 8
 9
10

inline const char* DeviceName(DLDeviceType type) {
  switch (type) {
    // 기존 장치들
    case kDLCPU: return "cpu";
    case kDLCUDA: return "cuda";
    // 새 장치 추가
    case kDLFoo: return "foo";
    default: return "unknown";
  }
}
  1. Python 딕셔너리 등록 (python/tvm/runtime/device.py)
1
2

DEVICE_TYPE_TO_NAME[13] = "foo"
DEVICE_NAME_TO_TYPE["foo"] = 13
  1. 등록 후 Python 사용
1
2
3
4

import tvm

dev = tvm.runtime.device("foo", 0)
print(dev)  # <tvm.runtime.Device foo(0)>
  • note) TVM은 DeviceAPIManager를 통해 장치별 구현을 자동으로 선택합니다.
  • note) void*로 반환된 핸들은 CopyDataFromTo, StreamSync 등 TVM 백엔드가 해석하므로, 메모리 주소일 필요는 없습니다.
  • note) DeviceAPI는 최소 SetDevice, AllocDataSpace, FreeDataSpace, CopyDataFromTo, StreamSync만 구현해도 동작 가능하지만, 고급 기능을 위해 CreateStream, SyncStreamFromTo, AllocWorkspace 등을 추가 구현할 수 있습니다.

Target Definition

Target object는 physical device와 그것의 hardware/driver제한사항 및 capabilities에 대한 properties의 룩업테이블이다. Target은 optimization과 code generation 단계 모두에서 접근이 가능하다. 같은 Target 클래스가 모든 런타임 대상(runtime target)에서 사용되지만, 각 런타임 대상마다 고유한 옵션을 추가해야 할 수 있습니다.

신규 TARGET_KIND 등록

target_kind.cc내부에 새로운 TVM_REGISTER_TARGET_KIND를 사용해 새로운 타겟 선언

  • 새로운 타켓의 string name과 타겟이 실행될 장치의 TVMDeviceExtType과 DLDeviceType enum 값을 전달 일반적으로 일반적으로 타겟 이름과 디바이스 이름은 일치하지만(“cuda” 타겟은 “kDLCUDA”) 하나의 물리 디바이스에서 서로 다른 코드 생성 타겟을 사용할 수 있기 때문에 예외도 있다. (“llvm"과 “c” 타겟은 둘 다 kDLCPU 디바이스 유형에서 실행)

타겟 옵션 설정

특정 target kind를 위한 옵션은 add_attr_option 함수를 통해 추가할 수 있다.(디폴트 값 설정 가능) Target parser도 set_target_parser로 추가 가능

  • 다른 파라메터에 동적으로 기반하거나 장치 속성에서 쿼리된 파라메터를 처리가능

Target 객체 생성

상기 argument definition은 타겟의 string description을 unpack할 수 있는 parser를 정의한다. 이것은 C++의 Target::Target(const String&) constructor에 의해 완료된다(JSON-formatted string을 수용할 수 있으며 tvm.target.Target 파이썬 object를 통해 호출됨).

  • 예를 들어 tvm.target.Target(’{“kind”: “cuda”, “max_num_threads”: 1024}’)은 max_num_threads를 오버라이딩하는 cuda 객체를 생성

Target 속성 접근

code generator에서 타겟 속성은 c++의 arget->GetAttr<T>(param_name) 또는 파이선의 target.attrs dictionary를 통해 접근 할 수 있다.

역자 주

  1. FooDeviceAPI를 위한 “foo” 타겟을 등록한다고 하면 다음처럼 구현
1
2
3
4
5

TVM_REGISTER_TARGET_KIND("foo", kDLFoo)
    .add_attr_option<Integer>("max_threads_per_block")
    .add_attr_option<Integer>("shared_memory_per_block")
    .add_attr_option<Array<String>>("libs")
    .set_default_keys({"foo", "gpu"});
  1. 등록된 타겟을 python에서 다음과 같이 사용가능하다.
 1
 2
 3
 4
 5
 6
 7
 8
 9
10

from tvm import target

t = target.Target({
    "kind": "foo",
    "max_threads_per_block": 256,
    "shared_memory_per_block": 49152
})

print(t.kind.name)  # foo
print(t.attrs["max_threads_per_block"])  # 256

Target Code Generators

code generator는 최적화된 IRModule을 입력받아 executable representation으로 변환한다. 각 code generator는 TVM framework에서 사용되기 위해서 등록되어야 하며 이는 target.build.foo 라는 이름의 함수를 등록함으로써 이루어진다(foo는 앞서 TVM_REGISTER_TARGET_KIND에서 사용한 타겟 이름과 같아야 함)

1
2

tvm::runtime::Module GeneratorFooCode(IRModule mod, Target target);
TVM_FFI_REGISTER_GLOBAL("target.build.foo").set_body_typed(GeneratorFooCode);

code generator는 2개의 인수를 취한다.

  • IRModule : 컴파일할 중간 표현 모듈
  • Target : 해당 코드가 실행될 디바이스에 대한 정보 컴파일 환경과 코드의 동작 환경이 다를 수 있으므로 코드 생성 시점에 실제 하드웨어 디바이스에 접근하거나 쿼리해서 파라메터를 받아 오지 말고 Target에 저장된 파라메터를 사용해야 함. input IRModule 내부 각 함수는 output runtime::Module에서 이름으로 접근가능해야 한다.

역자주

  1. code generator 예시 : GeneratorFooCode
 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38

#include <tvm/ir/module.h>
#include <tvm/runtime/module.h>
#include <tvm/runtime/registry.h>
#include <tvm/target/codegen.h>
#include <tvm/tir/stmt_functor.h>
#include <tvm/tir/function.h>
#include <tvm/codegen/codegen_c.h>  // CodeGenC 사용

using namespace tvm;
using namespace tvm::runtime;
using namespace tvm::tir;

runtime::Module GeneratorFooCode(IRModule mod, Target target) {
  // 1. C 코드 생성기 초기화
  auto cg = std::make_shared<CodeGenC>();
  cg->Init(target);  // 타겟 정보 전달

  // 2. IRModule 내 함수들을 순회하며 추가
  for (const auto& kv : mod->functions) {
    if (auto* n = kv.second.as<PrimFuncNode>()) {
      PrimFunc f = GetRef<PrimFunc>(n);
      cg->AddFunction(f);
    }
  }

  // 3. 생성된 C 코드 문자열 얻기
  std::string c_code = cg->Finish();

  // 4. 코드를 로그로 출력 (디버깅용)
  std::cout << "[FooCodeGen] Generated C Code:\n" << c_code << "\n";

  // 5. DSO 모듈로 빌드하는 대신, 그냥 소스 코드 모듈로 래핑해서 반환
  return codegen::CreateCSourceModule(c_code, "", Array<String>{}, "foo");
}

// 등록
TVM_REGISTER_GLOBAL("target.build.foo")
    .set_body_typed(GeneratorFooCode);
  1. 상기 코드에 대한 사용 예시
 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14

import tvm
from tvm import relay

x = relay.var("x", shape=(4,), dtype="float32")
y = relay.add(x, x)
func = relay.Function([x], y)
mod = tvm.IRModule.from_expr(func)

# foo 타겟으로 빌드
with tvm.transform.PassContext(opt_level=3):
    lib = relay.build(mod, target="foo")

# 생성된 C 코드를 확인
print(lib.get_source("foo"))
TVM
Licensed under CC BY-NC-SA 4.0
comments powered by Disqus
© 2022 - 2025 MW Devlog
Built with Hugo
Theme Stack designed by Jimmy