You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@tvm.apache.org by GitBox <gi...@apache.org> on 2022/05/05 20:03:56 UTC
[GitHub] [tvm] Lunderberg commented on a diff in pull request #11181: [LLVM] Fix a possible tbaa issue
Lunderberg commented on code in PR #11181:
URL: https://github.com/apache/tvm/pull/11181#discussion_r866246595
##########
src/target/llvm/codegen_llvm.cc:
##########
@@ -472,23 +472,47 @@ llvm::Type* CodeGenLLVM::GetLLVMType(const PrimExpr& expr) const {
//
// This trick comes from Halide's CodeGen_LLVM
//
-void CodeGenLLVM::AddAliasInfo(llvm::Instruction* inst, const VarNode* buffer, PrimExpr index) {
- if (alias_var_set_.count(buffer) != 0) {
+void CodeGenLLVM::AddAliasInfo(llvm::Instruction* inst, const VarNode* buffer_var, PrimExpr index,
+ DataType access_dtype) {
+ if (alias_var_set_.count(buffer_var) != 0) {
// Mark all possibly aliased pointer as same type.
llvm::MDNode* meta = md_tbaa_alias_set_;
inst->setMetadata("tbaa", md_builder_->createTBAAStructTagNode(meta, meta, 0));
return;
}
+ // Extract the underlying element bit width of the allocated buffer.
+ // fallback to byte type if no type annotation present.
+ int64_t buffer_elem_bits = 8;
+ int64_t access_elem_bits = access_dtype.bits() * access_dtype.lanes();
+ if (buffer_var->type_annotation.defined()) {
+ Type elem_ty = Downcast<PointerType>(buffer_var->type_annotation)->element_type;
+ if (auto* ptype = elem_ty.as<PrimTypeNode>()) {
+ if (!ptype->dtype.is_void()) {
+ buffer_elem_bits = ptype->dtype.bits() * ptype->dtype.lanes();
+ }
+ }
+ }
+
int64_t base = 0, width = 0;
arith::PVar<IntImm> pbase, pstride;
arith::PVar<int> planes;
// create meta-data for alias analysis
// Use a group of binary tree ranges of memory banks.
if (index.defined()) {
Review Comment:
Tangentially-related cleanup: I think we can remove the check on `index.defined()`. `AddAliasInfo` is only called from `BufferAccessHelper`, which provides a defined index.
##########
src/target/llvm/codegen_llvm.cc:
##########
@@ -472,23 +472,47 @@ llvm::Type* CodeGenLLVM::GetLLVMType(const PrimExpr& expr) const {
//
// This trick comes from Halide's CodeGen_LLVM
//
-void CodeGenLLVM::AddAliasInfo(llvm::Instruction* inst, const VarNode* buffer, PrimExpr index) {
- if (alias_var_set_.count(buffer) != 0) {
+void CodeGenLLVM::AddAliasInfo(llvm::Instruction* inst, const VarNode* buffer_var, PrimExpr index,
+ DataType access_dtype) {
+ if (alias_var_set_.count(buffer_var) != 0) {
// Mark all possibly aliased pointer as same type.
llvm::MDNode* meta = md_tbaa_alias_set_;
inst->setMetadata("tbaa", md_builder_->createTBAAStructTagNode(meta, meta, 0));
return;
}
+ // Extract the underlying element bit width of the allocated buffer.
+ // fallback to byte type if no type annotation present.
+ int64_t buffer_elem_bits = 8;
+ int64_t access_elem_bits = access_dtype.bits() * access_dtype.lanes();
+ if (buffer_var->type_annotation.defined()) {
+ Type elem_ty = Downcast<PointerType>(buffer_var->type_annotation)->element_type;
+ if (auto* ptype = elem_ty.as<PrimTypeNode>()) {
+ if (!ptype->dtype.is_void()) {
+ buffer_elem_bits = ptype->dtype.bits() * ptype->dtype.lanes();
+ }
+ }
+ }
+
int64_t base = 0, width = 0;
arith::PVar<IntImm> pbase, pstride;
arith::PVar<int> planes;
// create meta-data for alias analysis
// Use a group of binary tree ranges of memory banks.
if (index.defined()) {
+ int64_t xwith = 0;
if (arith::ramp(pbase, pstride, planes).Match(index)) {
base = pbase.Eval()->value;
- int64_t xwith = planes.Eval() * pstride.Eval()->value;
+ xwith = planes.Eval() * pstride.Eval()->value;
+ } else if (auto* ptr = index.as<tir::IntImmNode>()) {
+ base = ptr->value;
+ xwith = 1;
+ }
+ if (buffer_elem_bits != access_elem_bits) {
+ base = base * access_elem_bits / buffer_elem_bits;
Review Comment:
Would this cause false positives for aliasing of a buffer whose access type is smaller than the allocation type? I'm picturing something like the following:
```python
@T.prim_func
def func():
A = T.alloc_buffer(32, dtype='int32')
A_bytes = T.buffer_decl(128, dtype='int8', data=A.data)
A_bytes[0] = 42
A_bytes[3] = 42
```
By scaling the alias information to the size of the original allocation, both `A_bytes[0]` and `A_bytes[3]` are treated as access of `A[0]`. This would treat it as an alias even though they are accessing different addresses.
##########
src/target/llvm/codegen_llvm.cc:
##########
@@ -472,23 +472,47 @@ llvm::Type* CodeGenLLVM::GetLLVMType(const PrimExpr& expr) const {
//
// This trick comes from Halide's CodeGen_LLVM
//
-void CodeGenLLVM::AddAliasInfo(llvm::Instruction* inst, const VarNode* buffer, PrimExpr index) {
- if (alias_var_set_.count(buffer) != 0) {
+void CodeGenLLVM::AddAliasInfo(llvm::Instruction* inst, const VarNode* buffer_var, PrimExpr index,
+ DataType access_dtype) {
+ if (alias_var_set_.count(buffer_var) != 0) {
// Mark all possibly aliased pointer as same type.
llvm::MDNode* meta = md_tbaa_alias_set_;
inst->setMetadata("tbaa", md_builder_->createTBAAStructTagNode(meta, meta, 0));
return;
}
+ // Extract the underlying element bit width of the allocated buffer.
+ // fallback to byte type if no type annotation present.
+ int64_t buffer_elem_bits = 8;
Review Comment:
I don't think we need the size from the type annotation. The type annotation on the `Var` would only include the buffer's type as allocated, and may not be correlated with the type used for accessing it. When accessing the buffer in `CodeGenLLVM::CreateBufferPtr`, if the allocation type and access type differ, the buffer var is cast to the access type. So the bytes being accessed by a load/store should only depend on the access type and the access index.
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: commits-unsubscribe@tvm.apache.org
For queries about this service, please contact Infrastructure at:
users@infra.apache.org