You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@tvm.apache.org by tq...@apache.org on 2025/04/14 15:38:33 UTC
(tvm) branch refactor-s1 updated (4bd6ba1bcb -> afe55be216)
This is an automated email from the ASF dual-hosted git repository.
tqchen pushed a change to branch refactor-s1
in repository https://gitbox.apache.org/repos/asf/tvm.git
discard 4bd6ba1bcb add handle add by byteoffset
new 9ee70e95dc add handle add by byteoffset
new f6f5ba817d [CODEGEN] TVMFFIAny support in struct get/set/alloca
new afe55be216 [CODEGEN] handle_byte_offset
This update added new revisions after undoing existing revisions.
That is to say, some revisions that were in the old version of the
branch are not in the new version. This situation occurs
when a user --force pushes a change and generates a repository
containing something like this:
* -- * -- B -- O -- O -- O (4bd6ba1bcb)
\
N -- N -- N refs/heads/refactor-s1 (afe55be216)
You should already have received notification emails for all of the O
revisions, and so the following emails describe only the N revisions
from the common base, B.
Any revisions marked "omit" are not gone; other references still
refer to them. Any revisions marked "discard" are gone forever.
The 3 revisions listed above as "new" are entirely new to this
repository and will be described in separate emails. The revisions
listed as "add" were already present in the repository and have only
been added to this reference.
Summary of changes:
python/tvm/tir/op.py | 76 +++++++++++++++++++++++++++++--------
src/target/llvm/codegen_cpu.cc | 41 ++++++++++++++++++++
src/target/llvm/codegen_cpu.h | 3 +-
src/target/llvm/codegen_llvm.cc | 4 ++
src/target/llvm/codegen_llvm.h | 4 ++
src/target/source/codegen_c.cc | 29 ++++++++++++++
src/target/source/codegen_c_host.cc | 2 +
7 files changed, 141 insertions(+), 18 deletions(-)
(tvm) 02/03: [CODEGEN] TVMFFIAny support in struct get/set/alloca
Posted by tq...@apache.org.
This is an automated email from the ASF dual-hosted git repository.
tqchen pushed a commit to branch refactor-s1
in repository https://gitbox.apache.org/repos/asf/tvm.git
commit f6f5ba817d6c7366653a52e3fb64d18155303ede
Author: tqchen <ti...@gmail.com>
AuthorDate: Mon Apr 14 11:27:01 2025 -0400
[CODEGEN] TVMFFIAny support in struct get/set/alloca
---
src/target/llvm/codegen_cpu.cc | 41 +++++++++++++++++++++++++++++++++++++
src/target/llvm/codegen_cpu.h | 3 +--
src/target/llvm/codegen_llvm.h | 4 ++++
src/target/source/codegen_c.cc | 22 ++++++++++++++++++++
src/target/source/codegen_c_host.cc | 2 ++
5 files changed, 70 insertions(+), 2 deletions(-)
diff --git a/src/target/llvm/codegen_cpu.cc b/src/target/llvm/codegen_cpu.cc
index 5e30695517..ded7b91566 100644
--- a/src/target/llvm/codegen_cpu.cc
+++ b/src/target/llvm/codegen_cpu.cc
@@ -100,6 +100,8 @@ void CodeGenCPU::Init(const std::string& module_name, LLVMTarget* llvm_target,
// Defined in include/tvm/runtime/c_runtime_api.h:
// typedef union { ... } TVMValue;
t_tvm_value_ = llvm::StructType::create({t_float64_});
+ // Defined in include/tvm/ffi/c_api.h:
+ t_tvm_ffi_any_ = llvm::StructType::create({t_int32_, t_int32_, t_float64_});
// Defined in include/tvm/runtime/c_backend_api.h:
// typedef struct { void* sync_handle; int32_t num_task; } TVMParallelGroupEnv;
t_tvm_parallel_group_env_ = llvm::StructType::create({llvmGetPointerTo(t_int32_, 0), t_int32_});
@@ -345,6 +347,33 @@ CodeGenLLVM::TypedPointer CodeGenCPU::CreateStructRefPtr(DataType t, llvm::Value
builder_->CreateInBoundsGEP(t_tvm_array_, buf, {index, ConstInt32(1), ConstInt32(0)});
return TypedPointer(member_type, member_addr);
}
+ case builtin::kTVMFFIAnyTypeIndex: {
+ buf = builder_->CreatePointerCast(buf, llvmGetPointerTo(t_tvm_ffi_any_, 0));
+ buf = builder_->CreateInBoundsGEP(t_int32_, buf, {index, ConstInt32(0)});
+ return TypedPointer(t_int32_, buf);
+ }
+ case builtin::kTVMFFIAnyUnionValue: {
+ ICHECK_EQ(t.lanes(), 1);
+ buf = builder_->CreatePointerCast(buf, llvmGetPointerTo(t_tvm_ffi_any_, 0));
+ // field 2 is the union value
+ buf = builder_->CreateInBoundsGEP(t_tvm_ffi_any_, buf, {index, ConstInt32(2)});
+ if (t.is_bool()) {
+ // it should be safe to set the pointer to the first byte of the union value
+ buf = builder_->CreatePointerCast(buf, llvmGetPointerTo(DTypeToLLVMType(t), 0));
+ return TypedPointer(t_int8_, buf);
+ } else if (t.is_int() && t.bits() == 64) {
+ buf = builder_->CreatePointerCast(buf, llvmGetPointerTo(t_int64_, 0));
+ return TypedPointer(t_int64_, buf);
+ } else if (t.is_float() && t.bits() == 64) {
+ buf = builder_->CreatePointerCast(buf, llvmGetPointerTo(t_float64_, 0));
+ return TypedPointer(t_float64_, buf);
+ } else if (t.is_handle()) {
+ builder_->CreatePointerCast(buf, llvmGetPointerTo(t_void_p_, 0));
+ return TypedPointer(t_void_p_, buf);
+ } else {
+ LOG(DEBUG) << "DataType " << t << " cannot be stored into a TVMFFIAny's value field";
+ }
+ }
case builtin::kTVMValueContent: {
ICHECK_EQ(t.lanes(), 1);
if (t.is_bool()) {
@@ -1034,6 +1063,16 @@ llvm::Value* CodeGenCPU::CreateIntrinsic(const CallNode* op) {
if (value->getType()->isPointerTy()) {
value = builder_->CreatePointerCast(value, ref.type);
}
+
+ if (kind == builtin::kTVMFFIAnyUnionValue) {
+ // when we set any union value, we need to be careful to
+ // clear off the union value to zero if the set size is less than 64 bits
+ if (data_layout_->getTypeAllocSize(ref.type) != 8) {
+ llvm::Value* i64_addr = builder_->CreatePointerCast(
+ ref.addr, llvmGetPointerTo(t_int64_, 0));
+ builder_->CreateStore(ConstInt64(0), i64_addr);
+ }
+ }
builder_->CreateStore(value, ref.addr);
return ConstInt32(0);
} else if (op->op.same_as(builtin::tvm_stack_alloca())) {
@@ -1049,6 +1088,8 @@ llvm::Value* CodeGenCPU::CreateIntrinsic(const CallNode* op) {
return builder_->CreateAlloca(t_tvm_value_, num);
} else if (type == "arg_tcode") {
return builder_->CreateAlloca(t_int_, num);
+ } else if (type == "tvm_ffi_any") {
+ return builder_->CreateAlloca(t_tvm_ffi_any_, num);
} else if (type == "array") {
return builder_->CreateAlloca(t_tvm_array_, num);
} else {
diff --git a/src/target/llvm/codegen_cpu.h b/src/target/llvm/codegen_cpu.h
index 182dae81ce..5187b3e3f2 100644
--- a/src/target/llvm/codegen_cpu.h
+++ b/src/target/llvm/codegen_cpu.h
@@ -88,12 +88,11 @@ class CodeGenCPU : public CodeGenLLVM {
llvm::StructType* t_tvm_type_{nullptr};
llvm::StructType* t_tvm_array_{nullptr};
llvm::StructType* t_tvm_value_{nullptr};
+ llvm::StructType* t_tvm_ffi_any_{nullptr};
llvm::StructType* t_tvm_parallel_group_env_{nullptr};
-
llvm::FunctionType* ftype_tvm_backend_packed_c_func_{nullptr};
llvm::StructType* t_tvm_crt_func_registry_{nullptr};
llvm::StructType* t_tvm_crt_module_{nullptr};
-
llvm::FunctionType* ftype_tvm_parallel_lambda_{nullptr};
llvm::FunctionType* ftype_tvm_func_call_{nullptr};
llvm::FunctionType* ftype_tvm_get_func_from_env_{nullptr};
diff --git a/src/target/llvm/codegen_llvm.h b/src/target/llvm/codegen_llvm.h
index 302a0d97b3..c9cbc16b11 100644
--- a/src/target/llvm/codegen_llvm.h
+++ b/src/target/llvm/codegen_llvm.h
@@ -194,6 +194,10 @@ class CodeGenLLVM : public ExprFunctor<llvm::Value*(const PrimExpr&)>,
llvm::Constant* ConstInt32(int64_t value) const {
return llvm::ConstantInt::getSigned(t_int32_, value);
}
+ // Short hande code to get a constant int 64
+ llvm::Constant* ConstInt64(int64_t value) const {
+ return llvm::ConstantInt::getSigned(t_int64_, value);
+ }
// override codegen
llvm::Value* VisitExpr_(const VarNode* op) override;
llvm::Value* VisitExpr_(const CastNode* op) override;
diff --git a/src/target/source/codegen_c.cc b/src/target/source/codegen_c.cc
index a67cb80b91..2d05414f80 100644
--- a/src/target/source/codegen_c.cc
+++ b/src/target/source/codegen_c.cc
@@ -315,6 +315,28 @@ std::string CodeGenC::GetStructRef(DataType t, const PrimExpr& buffer, const Pri
}
os << ')';
return os.str();
+ } else if (kind == builtin::kTVMFFIAnyTypeIndex) {
+ std::ostringstream os;
+ os << "(((TVMFFIAny*)";
+ this->PrintExpr(buffer, os);
+ os << ")[" << index << "].type_index)";
+ return os.str();
+ } else if (kind == builtin::kTVMFFIAnyUnionValue) {
+ std::ostringstream os;
+ os << "(((TVMFFIAny*)";
+ this->PrintExpr(buffer, os);
+ os << ")[" << index << "].";
+ if (t.is_handle()) {
+ os << "v_ptr";
+ } else if (t.is_float()) {
+ os << "v_float64";
+ } else if (t.is_int()) {
+ os << "v_int64";
+ } else {
+ LOG(FATAL) << "Do not know how to handle type" << t;
+ }
+ os << ")";
+ return os.str();
} else {
ICHECK_LT(kind, builtin::kTVMValueKindBound_);
std::ostringstream os;
diff --git a/src/target/source/codegen_c_host.cc b/src/target/source/codegen_c_host.cc
index b00763a8a3..95c596fbd4 100644
--- a/src/target/source/codegen_c_host.cc
+++ b/src/target/source/codegen_c_host.cc
@@ -335,6 +335,8 @@ void CodeGenCHost::VisitExpr_(const CallNode* op, std::ostream& os) { // NOLINT
size = (num->value * sizeof(tvm_index_t) + unit - 1) / unit;
} else if (type == "arg_value") {
size = (num->value * sizeof(TVMValue) + unit - 1) / unit;
+ } else if (type == "tvm_ffi_any") {
+ size = (num->value * sizeof(TVMFFIAny) + unit - 1) / unit;
} else if (type == "arg_tcode") {
size = (num->value * sizeof(int) + unit - 1) / unit;
} else if (type == "array") {
(tvm) 01/03: add handle add by byteoffset
Posted by tq...@apache.org.
This is an automated email from the ASF dual-hosted git repository.
tqchen pushed a commit to branch refactor-s1
in repository https://gitbox.apache.org/repos/asf/tvm.git
commit 9ee70e95dc88c3a29d4cc61e37667610a6f47b87
Author: tqchen <ti...@gmail.com>
AuthorDate: Mon Apr 14 10:50:26 2025 -0400
add handle add by byteoffset
---
include/tvm/tir/builtin.h | 11 ++++
python/tvm/script/ir_builder/tir/ir.py | 2 +
python/tvm/tir/__init__.py | 2 +-
python/tvm/tir/op.py | 91 ++++++++++++++++++++++++++++------
src/tir/op/builtin.cc | 4 ++
5 files changed, 95 insertions(+), 15 deletions(-)
diff --git a/include/tvm/tir/builtin.h b/include/tvm/tir/builtin.h
index 2579250de2..822763d0b2 100644
--- a/include/tvm/tir/builtin.h
+++ b/include/tvm/tir/builtin.h
@@ -263,6 +263,15 @@ TVM_DLL const Op& tvm_context_id();
*/
TVM_DLL const Op& tvm_tuple();
+/*!
+ * \brief See pesudo code
+ *
+ * void* handle_add_byte_offset(void* handle, int offset) {
+ * return reinterpret_cast<v*>(reinterpret_cast<char*>(handle) + offset);
+ * }
+ */
+TVM_DLL const Op& handle_add_byte_offset();
+
/*!
* \brief See pesudo code
*
@@ -990,6 +999,8 @@ enum TVMStructFieldKind : int {
kArrKindBound_,
// TVMValue field
kTVMValueContent,
+ kTVMFFIAnyTypeIndex,
+ kTVMFFIAnyUnionValue,
kTVMValueKindBound_
};
} // namespace builtin
diff --git a/python/tvm/script/ir_builder/tir/ir.py b/python/tvm/script/ir_builder/tir/ir.py
index 3e835e8d9d..3814f2df88 100644
--- a/python/tvm/script/ir_builder/tir/ir.py
+++ b/python/tvm/script/ir_builder/tir/ir.py
@@ -1881,6 +1881,7 @@ call_cpacked = _op_wrapper(_tir_op.call_cpacked)
call_packed_lowered = _op_wrapper(_tir_op.call_packed_lowered)
call_cpacked_lowered = _op_wrapper(_tir_op.call_cpacked_lowered)
tvm_tuple = _op_wrapper(_tir_op.tvm_tuple)
+handle_add_byte_offset = _op_wrapper(_tir_op.handle_add_byte_offset)
tvm_struct_set = _op_wrapper(_tir_op.tvm_struct_set)
tvm_struct_get = _tir_op.tvm_struct_get
tvm_thread_invariant = _op_wrapper(_tir_op.tvm_thread_invariant)
@@ -2178,6 +2179,7 @@ __all__ = [
"call_llvm_pure_intrin",
"call_pure_extern",
"tvm_tuple",
+ "handle_add_byte_offset",
"tvm_struct_set",
"tvm_struct_get",
"tvm_thread_invariant",
diff --git a/python/tvm/tir/__init__.py b/python/tvm/tir/__init__.py
index 4f56ec3c15..63d3cb8f31 100644
--- a/python/tvm/tir/__init__.py
+++ b/python/tvm/tir/__init__.py
@@ -50,7 +50,7 @@ from .op import call_packed, call_cpacked, call_intrin, call_pure_extern, call_e
from .op import call_llvm_intrin, call_llvm_pure_intrin, ret, all, any, min_value, max_value, trace
from .op import tvm_check_return
from .op import tvm_stack_alloca, tvm_stack_make_shape, tvm_stack_make_array
-from .op import tvm_tuple, tvm_struct_get, tvm_struct_set
+from .op import tvm_tuple, handle_add_byte_offset, tvm_struct_get, tvm_struct_set
from .op import address_of, lookup_param, assume, undef
from .op import tvm_thread_allreduce, type_annotation, tvm_access_ptr, tvm_throw_last_error
from .op import (
diff --git a/python/tvm/tir/op.py b/python/tvm/tir/op.py
index 53c92fff86..2e0d2a64c9 100644
--- a/python/tvm/tir/op.py
+++ b/python/tvm/tir/op.py
@@ -32,7 +32,11 @@ from .expr import Call, CommReducer, IntImm, PrimExprWithOp, Var
def _pack_buffer(buf, span=None):
"""Build intrinsics that packs the buffer."""
shape = Call("handle", "tir.tvm_stack_make_shape", buf.shape, span)
- strides = Call("handle", "tir.tvm_stack_make_shape", buf.strides, span) if buf.strides else 0
+ strides = (
+ Call("handle", "tir.tvm_stack_make_shape", buf.strides, span)
+ if buf.strides
+ else 0
+ )
pack_args = [
buf.data,
shape,
@@ -334,7 +338,9 @@ def tvm_check_return(expected, return_unexpected, nested_call):
call : PrimExpr
The call expression.
"""
- return call_intrin("int32", "tir.tvm_check_return", expected, return_unexpected, nested_call)
+ return call_intrin(
+ "int32", "tir.tvm_check_return", expected, return_unexpected, nested_call
+ )
def tvm_stack_alloca(dtype_str, num):
@@ -401,7 +407,14 @@ def tvm_stack_make_array(data, shape, strides, ndim, arr_dtype, elem_offset):
The call expression.
"""
return call_intrin(
- "handle", "tir.tvm_stack_make_array", data, shape, strides, ndim, arr_dtype, elem_offset
+ "handle",
+ "tir.tvm_stack_make_array",
+ data,
+ shape,
+ strides,
+ ndim,
+ arr_dtype,
+ elem_offset,
)
@@ -495,6 +508,25 @@ def tvm_tuple(*value):
return call_intrin("handle", "tir.tvm_tuple", *value)
+def handle_add_byte_offset(handle, offset):
+ """Add offset to handle
+
+ Parameters
+ ----------
+ handle : Expr
+ The handle.
+
+ offset : int
+ The offset.
+
+ Returns
+ -------
+ call : PrimExpr
+ The call expression.
+ """
+ return call_intrin("handle", "tir.handle_add_byte_offset", handle, offset)
+
+
def tvm_struct_get(arr, index, field, dtype):
"""Get struct field value in array
@@ -653,7 +685,9 @@ def tvm_warp_shuffle(mask, value, warp_id, width, warp_size):
call : PrimExpr
The call expression.
"""
- return call_intrin(value.dtype, "tir.tvm_warp_shuffle", mask, value, warp_id, width, warp_size)
+ return call_intrin(
+ value.dtype, "tir.tvm_warp_shuffle", mask, value, warp_id, width, warp_size
+ )
def tvm_warp_shuffle_up(mask, value, offset, width, warp_size):
@@ -762,7 +796,9 @@ def tvm_access_ptr(ptype, data, offset, extent, rw_mask):
call : PrimExpr
The call expression.
"""
- return call_intrin("handle", "tir.tvm_access_ptr", ptype, data, offset, extent, rw_mask)
+ return call_intrin(
+ "handle", "tir.tvm_access_ptr", ptype, data, offset, extent, rw_mask
+ )
def tvm_throw_last_error():
@@ -1298,7 +1334,9 @@ def mma_fill(dtype, local_size, local_ptr, offset):
)
-def ptx_ldmatrix(dtype, trans, num, type, local_ptr, local_offset, smem_ptr, smem_offset):
+def ptx_ldmatrix(
+ dtype, trans, num, type, local_ptr, local_offset, smem_ptr, smem_offset
+):
"""TVM intrinsic for ptx load matrix from shared memory
https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#warp-level-matrix-instructions-ldmatrix
@@ -1376,7 +1414,13 @@ def ptx_cp_async(dtype, shared_ptr, shared_offset, global_ptr, global_offset, by
The call expression.
"""
return call_intrin(
- dtype, "tir.ptx_cp_async", shared_ptr, shared_offset, global_ptr, global_offset, bytes
+ dtype,
+ "tir.ptx_cp_async",
+ shared_ptr,
+ shared_offset,
+ global_ptr,
+ global_offset,
+ bytes,
)
@@ -1489,7 +1533,9 @@ def ptx_init_barrier_thread_count(barrier_id, thread_count):
call : PrimExpr
The call expression.
"""
- return call_intrin("", "tir.ptx_init_barrier_thread_count", barrier_id, thread_count)
+ return call_intrin(
+ "", "tir.ptx_init_barrier_thread_count", barrier_id, thread_count
+ )
def ptx_arrive_barrier(barrier_id):
@@ -1595,7 +1641,9 @@ def make_filled_simdgroup_matrix(
call : PrimExpr
The call expression.
"""
- return call_intrin("handle", "tir.make_filled_simdgroup_matrix", d, index, value, col, row)
+ return call_intrin(
+ "handle", "tir.make_filled_simdgroup_matrix", d, index, value, col, row
+ )
def simdgroup_load(
@@ -1691,7 +1739,15 @@ def simdgroup_store(
The call expression.
"""
return call_intrin(
- "handle", "tir.simdgroup_store", d, index, ptr, stride, col, row, transpose_matrix
+ "handle",
+ "tir.simdgroup_store",
+ d,
+ index,
+ ptr,
+ stride,
+ col,
+ row,
+ transpose_matrix,
)
@@ -3348,10 +3404,13 @@ def comm_reducer(fcombine, fidentity, name="reduce"):
if where is None:
where = tir.convert(True)
if init is None:
- outputs = tuple(tvm.tir.Reduce(combiner, expr, axis, where, i, []) for i in range(size))
+ outputs = tuple(
+ tvm.tir.Reduce(combiner, expr, axis, where, i, []) for i in range(size)
+ )
else:
outputs = tuple(
- tvm.tir.Reduce(combiner, expr, axis, where, i, init) for i in range(size)
+ tvm.tir.Reduce(combiner, expr, axis, where, i, init)
+ for i in range(size)
)
return outputs[0] if size == 1 else outputs
@@ -3407,7 +3466,9 @@ def comm_reducer(fcombine, fidentity, name="reduce"):
return reducer
-def TVMBackendAllocWorkspace(device_type, device_id, nbytes, dtype_code_hint, dtype_bits_hint):
+def TVMBackendAllocWorkspace(
+ device_type, device_id, nbytes, dtype_code_hint, dtype_bits_hint
+):
"""Backend function to allocate temporal workspace
Parameters
@@ -3462,7 +3523,9 @@ def TVMBackendFreeWorkspace(device_type, device_id, ptr):
call : PrimExpr
The call expression.
"""
- return call_intrin("int32", "tir.TVMBackendFreeWorkspace", device_type, device_id, ptr)
+ return call_intrin(
+ "int32", "tir.TVMBackendFreeWorkspace", device_type, device_id, ptr
+ )
def anylist_getitem(list_handle, index):
diff --git a/src/tir/op/builtin.cc b/src/tir/op/builtin.cc
index 39a09d9922..f0894acecf 100644
--- a/src/tir/op/builtin.cc
+++ b/src/tir/op/builtin.cc
@@ -163,6 +163,10 @@ TIR_DEFINE_BUILTIN_FUNC(tvm_context_id)
TIR_DEFINE_BUILTIN_FUNC(tvm_tuple).set_attr<TCallEffectKind>("TCallEffectKind",
Integer(CallEffectKind::kEmbedInfo));
+TIR_DEFINE_BUILTIN_FUNC(handle_add_byte_offset)
+ .set_num_inputs(2)
+ .set_attr<TCallEffectKind>("TCallEffectKind", Integer(CallEffectKind::kPure));
+
TIR_DEFINE_BUILTIN_FUNC(tvm_struct_get)
.set_num_inputs(3)
.set_attr<TCallEffectKind>("TCallEffectKind", Integer(CallEffectKind::kReadState))
(tvm) 03/03: [CODEGEN] handle_byte_offset
Posted by tq...@apache.org.
This is an automated email from the ASF dual-hosted git repository.
tqchen pushed a commit to branch refactor-s1
in repository https://gitbox.apache.org/repos/asf/tvm.git
commit afe55be2164fac1d05372f8e11ff60b3e98bb709
Author: tqchen <ti...@gmail.com>
AuthorDate: Mon Apr 14 11:35:45 2025 -0400
[CODEGEN] handle_byte_offset
---
src/target/llvm/codegen_cpu.cc | 4 ++--
src/target/llvm/codegen_llvm.cc | 4 ++++
src/target/source/codegen_c.cc | 11 +++++++++--
3 files changed, 15 insertions(+), 4 deletions(-)
diff --git a/src/target/llvm/codegen_cpu.cc b/src/target/llvm/codegen_cpu.cc
index ded7b91566..0bc40a306f 100644
--- a/src/target/llvm/codegen_cpu.cc
+++ b/src/target/llvm/codegen_cpu.cc
@@ -1068,8 +1068,8 @@ llvm::Value* CodeGenCPU::CreateIntrinsic(const CallNode* op) {
// when we set any union value, we need to be careful to
// clear off the union value to zero if the set size is less than 64 bits
if (data_layout_->getTypeAllocSize(ref.type) != 8) {
- llvm::Value* i64_addr = builder_->CreatePointerCast(
- ref.addr, llvmGetPointerTo(t_int64_, 0));
+ llvm::Value* i64_addr =
+ builder_->CreatePointerCast(ref.addr, llvmGetPointerTo(t_int64_, 0));
builder_->CreateStore(ConstInt64(0), i64_addr);
}
}
diff --git a/src/target/llvm/codegen_llvm.cc b/src/target/llvm/codegen_llvm.cc
index 5a2fb86311..c2ea454064 100644
--- a/src/target/llvm/codegen_llvm.cc
+++ b/src/target/llvm/codegen_llvm.cc
@@ -1418,6 +1418,10 @@ llvm::Value* CodeGenLLVM::CreateIntrinsic(const CallNode* op) {
return llvm::Constant::getNullValue(t_void_p_);
} else if (op->op.same_as(builtin::isnullptr())) {
return builder_->CreateIsNull(MakeValue(op->args[0]));
+ } else if (op->op.same_as(builtin::handle_add_byte_offset())) {
+ llvm::Value* ptr = MakeValue(op->args[0]);
+ llvm::Value* offset = MakeValue(op->args[1]);
+ return builder_->CreateInBoundsGEP(t_int8_, ptr, offset);
} else if (op->op.same_as(builtin::large_uint_imm())) {
ICHECK_EQ(op->args.size(), 2U);
uint64_t low = static_cast<uint64_t>(Downcast<IntImm>(op->args[0])->value);
diff --git a/src/target/source/codegen_c.cc b/src/target/source/codegen_c.cc
index 2d05414f80..db5eaac50b 100644
--- a/src/target/source/codegen_c.cc
+++ b/src/target/source/codegen_c.cc
@@ -321,8 +321,8 @@ std::string CodeGenC::GetStructRef(DataType t, const PrimExpr& buffer, const Pri
this->PrintExpr(buffer, os);
os << ")[" << index << "].type_index)";
return os.str();
- } else if (kind == builtin::kTVMFFIAnyUnionValue) {
- std::ostringstream os;
+ } else if (kind == builtin::kTVMFFIAnyUnionValue) {
+ std::ostringstream os;
os << "(((TVMFFIAny*)";
this->PrintExpr(buffer, os);
os << ")[" << index << "].";
@@ -693,6 +693,13 @@ void CodeGenC::VisitExpr_(const CallNode* op, std::ostream& os) { // NOLINT(*)
os << "(";
this->PrintExpr(op->args[0], os);
os << " == NULL)";
+ } else if (op->op.same_as(builtin::handle_add_byte_offset())) {
+ ICHECK_EQ(op->args.size(), 2U);
+ os << "((void*)((char*)";
+ this->PrintExpr(op->args[0], os);
+ os << " + ";
+ this->PrintExpr(op->args[1], os);
+ os << "))";
} else if (op->op.same_as(builtin::reinterpret())) {
auto target_dtype = op->dtype;
auto source_dtype = op->args[0]->dtype;