[llvm] fa342b5 - Enable `align <n>` to be used in the intrinsic definition.
Michael Liao via llvm-commits
llvm-commits at lists.llvm.org
Wed May 27 13:38:44 PDT 2020
Author: Michael Liao
Date: 2020-05-27T16:38:18-04:00
New Revision: fa342b5c8054dad4cfd1032ac580d71f0f4943d3
URL: https://github.com/llvm/llvm-project/commit/fa342b5c8054dad4cfd1032ac580d71f0f4943d3
DIFF: https://github.com/llvm/llvm-project/commit/fa342b5c8054dad4cfd1032ac580d71f0f4943d3.diff
LOG: Enable `align <n>` to be used in the intrinsic definition.
- This allow us to specify the (minimal) alignment on an intrinsic's
arguments and, more importantly, the return value.
Differential Revision: https://reviews.llvm.org/D80422
Added:
Modified:
llvm/include/llvm/IR/Attributes.h
llvm/include/llvm/IR/Intrinsics.td
llvm/include/llvm/IR/IntrinsicsAMDGPU.td
llvm/lib/IR/Attributes.cpp
llvm/test/CodeGen/AMDGPU/reqd-work-group-size.ll
llvm/utils/TableGen/CodeGenIntrinsics.h
llvm/utils/TableGen/CodeGenTarget.cpp
llvm/utils/TableGen/IntrinsicEmitter.cpp
Removed:
################################################################################
diff --git a/llvm/include/llvm/IR/Attributes.h b/llvm/include/llvm/IR/Attributes.h
index a76e29132acc..58365aa2b764 100644
--- a/llvm/include/llvm/IR/Attributes.h
+++ b/llvm/include/llvm/IR/Attributes.h
@@ -396,6 +396,9 @@ class AttributeList {
static AttributeList get(LLVMContext &C, ArrayRef<AttributeList> Attrs);
static AttributeList get(LLVMContext &C, unsigned Index,
ArrayRef<Attribute::AttrKind> Kinds);
+ static AttributeList get(LLVMContext &C, unsigned Index,
+ ArrayRef<Attribute::AttrKind> Kinds,
+ ArrayRef<uint64_t> Values);
static AttributeList get(LLVMContext &C, unsigned Index,
ArrayRef<StringRef> Kind);
static AttributeList get(LLVMContext &C, unsigned Index,
diff --git a/llvm/include/llvm/IR/Intrinsics.td b/llvm/include/llvm/IR/Intrinsics.td
index f6df3faba83f..78409df8f816 100644
--- a/llvm/include/llvm/IR/Intrinsics.td
+++ b/llvm/include/llvm/IR/Intrinsics.td
@@ -81,6 +81,11 @@ class NoAlias<AttrIndex idx> : IntrinsicProperty {
int ArgNo = idx.Value;
}
+class Align<AttrIndex idx, int align> : IntrinsicProperty {
+ int ArgNo = idx.Value;
+ int Align = align;
+}
+
// Returned - The specified argument is always the return value of the
// intrinsic.
class Returned<AttrIndex idx> : IntrinsicProperty {
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index 132d6b7360f7..e2d8f3cb1bd6 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -142,22 +142,22 @@ defm int_amdgcn_workgroup_id : AMDGPUReadPreloadRegisterIntrinsic_xyz_named
def int_amdgcn_dispatch_ptr :
Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 4>], [],
- [IntrNoMem, IntrSpeculatable]>;
+ [Align<RetIndex, 4>, IntrNoMem, IntrSpeculatable]>;
def int_amdgcn_queue_ptr :
GCCBuiltin<"__builtin_amdgcn_queue_ptr">,
Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 4>], [],
- [IntrNoMem, IntrSpeculatable]>;
+ [Align<RetIndex, 4>, IntrNoMem, IntrSpeculatable]>;
def int_amdgcn_kernarg_segment_ptr :
GCCBuiltin<"__builtin_amdgcn_kernarg_segment_ptr">,
Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 4>], [],
- [IntrNoMem, IntrSpeculatable]>;
+ [Align<RetIndex, 4>, IntrNoMem, IntrSpeculatable]>;
def int_amdgcn_implicitarg_ptr :
GCCBuiltin<"__builtin_amdgcn_implicitarg_ptr">,
Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 4>], [],
- [IntrNoMem, IntrSpeculatable]>;
+ [Align<RetIndex, 4>, IntrNoMem, IntrSpeculatable]>;
def int_amdgcn_groupstaticsize :
GCCBuiltin<"__builtin_amdgcn_groupstaticsize">,
@@ -170,7 +170,7 @@ def int_amdgcn_dispatch_id :
def int_amdgcn_implicit_buffer_ptr :
GCCBuiltin<"__builtin_amdgcn_implicit_buffer_ptr">,
Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 4>], [],
- [IntrNoMem, IntrSpeculatable]>;
+ [Align<RetIndex, 4>, IntrNoMem, IntrSpeculatable]>;
// Set EXEC to the 64-bit value given.
// This is always moved to the beginning of the basic block.
diff --git a/llvm/lib/IR/Attributes.cpp b/llvm/lib/IR/Attributes.cpp
index 122cfe5d5fca..191668dacc18 100644
--- a/llvm/lib/IR/Attributes.cpp
+++ b/llvm/lib/IR/Attributes.cpp
@@ -1174,6 +1174,17 @@ AttributeList AttributeList::get(LLVMContext &C, unsigned Index,
return get(C, Attrs);
}
+AttributeList AttributeList::get(LLVMContext &C, unsigned Index,
+ ArrayRef<Attribute::AttrKind> Kinds,
+ ArrayRef<uint64_t> Values) {
+ assert(Kinds.size() == Values.size() && "Mismatched attribute values.");
+ SmallVector<std::pair<unsigned, Attribute>, 8> Attrs;
+ auto VI = Values.begin();
+ for (const auto K : Kinds)
+ Attrs.emplace_back(Index, Attribute::get(C, K, *VI++));
+ return get(C, Attrs);
+}
+
AttributeList AttributeList::get(LLVMContext &C, unsigned Index,
ArrayRef<StringRef> Kinds) {
SmallVector<std::pair<unsigned, Attribute>, 8> Attrs;
diff --git a/llvm/test/CodeGen/AMDGPU/reqd-work-group-size.ll b/llvm/test/CodeGen/AMDGPU/reqd-work-group-size.ll
index 1903f89789b1..1e69d4551359 100644
--- a/llvm/test/CodeGen/AMDGPU/reqd-work-group-size.ll
+++ b/llvm/test/CodeGen/AMDGPU/reqd-work-group-size.ll
@@ -390,7 +390,7 @@ define amdgpu_kernel void @all_local_size(i64 addrspace(1)* nocapture readnone %
; CHECK-LABEL: @partial_load_group_size_x(
; CHECK-NEXT: %dispatch.ptr = tail call i8 addrspace(4)* @llvm.amdgcn.dispatch.ptr()
; CHECK-NEXT: %gep.group.size.x = getelementptr inbounds i8, i8 addrspace(4)* %dispatch.ptr, i64 4
-; CHECK-NEXT: %group.size.x.lo = load i8, i8 addrspace(4)* %gep.group.size.x, align 1
+; CHECK-NEXT: %group.size.x.lo = load i8, i8 addrspace(4)* %gep.group.size.x, align 4
; CHECK-NEXT: store i8 %group.size.x.lo, i8 addrspace(1)* %out, align 1
define amdgpu_kernel void @partial_load_group_size_x(i8 addrspace(1)* %out) #0 !reqd_work_group_size !0 {
%dispatch.ptr = tail call i8 addrspace(4)* @llvm.amdgcn.dispatch.ptr()
@@ -400,6 +400,19 @@ define amdgpu_kernel void @partial_load_group_size_x(i8 addrspace(1)* %out) #0 !
ret void
}
+; CHECK-LABEL: @partial_load_group_size_x_explicit_callsite_align(
+; CHECK-NEXT: %dispatch.ptr = tail call align 2 i8 addrspace(4)* @llvm.amdgcn.dispatch.ptr()
+; CHECK-NEXT: %gep.group.size.x = getelementptr inbounds i8, i8 addrspace(4)* %dispatch.ptr, i64 4
+; CHECK-NEXT: %group.size.x.lo = load i8, i8 addrspace(4)* %gep.group.size.x, align 2
+; CHECK-NEXT: store i8 %group.size.x.lo, i8 addrspace(1)* %out, align 1
+define amdgpu_kernel void @partial_load_group_size_x_explicit_callsite_align(i8 addrspace(1)* %out) #0 !reqd_work_group_size !0 {
+ %dispatch.ptr = tail call align 2 i8 addrspace(4)* @llvm.amdgcn.dispatch.ptr()
+ %gep.group.size.x = getelementptr inbounds i8, i8 addrspace(4)* %dispatch.ptr, i64 4
+ %group.size.x.lo = load i8, i8 addrspace(4)* %gep.group.size.x, align 1
+ store i8 %group.size.x.lo, i8 addrspace(1)* %out
+ ret void
+}
+
; TODO: Should be able to handle this
; CHECK-LABEL: @load_group_size_xy_i32(
; CHECK: %group.size.xy = load i32,
diff --git a/llvm/utils/TableGen/CodeGenIntrinsics.h b/llvm/utils/TableGen/CodeGenIntrinsics.h
index 5ebdbf995ebf..6503f39cfd8e 100644
--- a/llvm/utils/TableGen/CodeGenIntrinsics.h
+++ b/llvm/utils/TableGen/CodeGenIntrinsics.h
@@ -149,18 +149,21 @@ struct CodeGenIntrinsic {
ReadOnly,
WriteOnly,
ReadNone,
- ImmArg
+ ImmArg,
+ Alignment
};
struct ArgAttribute {
unsigned Index;
ArgAttrKind Kind;
+ uint64_t Value;
- ArgAttribute(unsigned Idx, ArgAttrKind K)
- : Index(Idx), Kind(K) {}
+ ArgAttribute(unsigned Idx, ArgAttrKind K, uint64_t V)
+ : Index(Idx), Kind(K), Value(V) {}
bool operator<(const ArgAttribute &Other) const {
- return std::tie(Index, Kind) < std::tie(Other.Index, Other.Kind);
+ return std::tie(Index, Kind, Value) <
+ std::tie(Other.Index, Other.Kind, Other.Value);
}
};
diff --git a/llvm/utils/TableGen/CodeGenTarget.cpp b/llvm/utils/TableGen/CodeGenTarget.cpp
index 35d5deecf32c..78fb73287701 100644
--- a/llvm/utils/TableGen/CodeGenTarget.cpp
+++ b/llvm/utils/TableGen/CodeGenTarget.cpp
@@ -795,25 +795,29 @@ CodeGenIntrinsic::CodeGenIntrinsic(Record *R) {
hasSideEffects = true;
else if (Property->isSubClassOf("NoCapture")) {
unsigned ArgNo = Property->getValueAsInt("ArgNo");
- ArgumentAttributes.emplace_back(ArgNo, NoCapture);
+ ArgumentAttributes.emplace_back(ArgNo, NoCapture, 0);
} else if (Property->isSubClassOf("NoAlias")) {
unsigned ArgNo = Property->getValueAsInt("ArgNo");
- ArgumentAttributes.emplace_back(ArgNo, NoAlias);
+ ArgumentAttributes.emplace_back(ArgNo, NoAlias, 0);
} else if (Property->isSubClassOf("Returned")) {
unsigned ArgNo = Property->getValueAsInt("ArgNo");
- ArgumentAttributes.emplace_back(ArgNo, Returned);
+ ArgumentAttributes.emplace_back(ArgNo, Returned, 0);
} else if (Property->isSubClassOf("ReadOnly")) {
unsigned ArgNo = Property->getValueAsInt("ArgNo");
- ArgumentAttributes.emplace_back(ArgNo, ReadOnly);
+ ArgumentAttributes.emplace_back(ArgNo, ReadOnly, 0);
} else if (Property->isSubClassOf("WriteOnly")) {
unsigned ArgNo = Property->getValueAsInt("ArgNo");
- ArgumentAttributes.emplace_back(ArgNo, WriteOnly);
+ ArgumentAttributes.emplace_back(ArgNo, WriteOnly, 0);
} else if (Property->isSubClassOf("ReadNone")) {
unsigned ArgNo = Property->getValueAsInt("ArgNo");
- ArgumentAttributes.emplace_back(ArgNo, ReadNone);
+ ArgumentAttributes.emplace_back(ArgNo, ReadNone, 0);
} else if (Property->isSubClassOf("ImmArg")) {
unsigned ArgNo = Property->getValueAsInt("ArgNo");
- ArgumentAttributes.emplace_back(ArgNo, ImmArg);
+ ArgumentAttributes.emplace_back(ArgNo, ImmArg, 0);
+ } else if (Property->isSubClassOf("Align")) {
+ unsigned ArgNo = Property->getValueAsInt("ArgNo");
+ uint64_t Align = Property->getValueAsInt("Align");
+ ArgumentAttributes.emplace_back(ArgNo, Alignment, Align);
} else
llvm_unreachable("Unknown property!");
}
@@ -834,7 +838,7 @@ bool CodeGenIntrinsic::isParamAPointer(unsigned ParamIdx) const {
bool CodeGenIntrinsic::isParamImmArg(unsigned ParamIdx) const {
// Convert argument index to attribute index starting from `FirstArgIndex`.
- ArgAttribute Val{ParamIdx + 1, ImmArg};
+ ArgAttribute Val{ParamIdx + 1, ImmArg, 0};
return std::binary_search(ArgumentAttributes.begin(),
ArgumentAttributes.end(), Val);
}
diff --git a/llvm/utils/TableGen/IntrinsicEmitter.cpp b/llvm/utils/TableGen/IntrinsicEmitter.cpp
index 0480a838ea6c..ab42f33cf23f 100644
--- a/llvm/utils/TableGen/IntrinsicEmitter.cpp
+++ b/llvm/utils/TableGen/IntrinsicEmitter.cpp
@@ -668,6 +668,8 @@ void IntrinsicEmitter::EmitAttributes(const CodeGenIntrinsicTable &Ints,
OS << " const Attribute::AttrKind AttrParam" << attrIdx << "[]= {";
bool addComma = false;
+ bool AllValuesAreZero = true;
+ SmallVector<uint64_t, 8> Values;
do {
switch (intrinsic.ArgumentAttributes[ai].Kind) {
case CodeGenIntrinsic::NoCapture:
@@ -712,13 +714,39 @@ void IntrinsicEmitter::EmitAttributes(const CodeGenIntrinsicTable &Ints,
OS << "Attribute::ImmArg";
addComma = true;
break;
+ case CodeGenIntrinsic::Alignment:
+ if (addComma)
+ OS << ',';
+ OS << "Attribute::Alignment";
+ addComma = true;
+ break;
}
+ uint64_t V = intrinsic.ArgumentAttributes[ai].Value;
+ Values.push_back(V);
+ AllValuesAreZero &= (V == 0);
++ai;
} while (ai != ae && intrinsic.ArgumentAttributes[ai].Index == attrIdx);
OS << "};\n";
+
+ // Generate attribute value array if not all attribute values are zero.
+ if (!AllValuesAreZero) {
+ OS << " const uint64_t AttrValParam" << attrIdx << "[]= {";
+ addComma = false;
+ for (const auto V : Values) {
+ if (addComma)
+ OS << ',';
+ OS << V;
+ addComma = true;
+ }
+ OS << "};\n";
+ }
+
OS << " AS[" << numAttrs++ << "] = AttributeList::get(C, "
- << attrIdx << ", AttrParam" << attrIdx << ");\n";
+ << attrIdx << ", AttrParam" << attrIdx;
+ if (!AllValuesAreZero)
+ OS << ", AttrValParam" << attrIdx;
+ OS << ");\n";
}
}
More information about the llvm-commits
mailing list