diff --git a/llvm/include/llvm/IR/Attributes.h b/llvm/include/llvm/IR/Attributes.h --- a/llvm/include/llvm/IR/Attributes.h +++ b/llvm/include/llvm/IR/Attributes.h @@ -397,6 +397,9 @@ static AttributeList get(LLVMContext &C, unsigned Index, ArrayRef Kinds); static AttributeList get(LLVMContext &C, unsigned Index, + ArrayRef Kinds, + ArrayRef Values); + static AttributeList get(LLVMContext &C, unsigned Index, ArrayRef Kind); static AttributeList get(LLVMContext &C, unsigned Index, const AttrBuilder &B); diff --git a/llvm/include/llvm/IR/Intrinsics.td b/llvm/include/llvm/IR/Intrinsics.td --- a/llvm/include/llvm/IR/Intrinsics.td +++ b/llvm/include/llvm/IR/Intrinsics.td @@ -81,6 +81,11 @@ int ArgNo = idx.Value; } +class Align : IntrinsicProperty { + int ArgNo = idx.Value; + int Align = align; +} + // Returned - The specified argument is always the return value of the // intrinsic. class Returned : IntrinsicProperty { diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td --- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td +++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td @@ -142,22 +142,22 @@ def int_amdgcn_dispatch_ptr : Intrinsic<[LLVMQualPointerType], [], - [IntrNoMem, IntrSpeculatable]>; + [Align, IntrNoMem, IntrSpeculatable]>; def int_amdgcn_queue_ptr : GCCBuiltin<"__builtin_amdgcn_queue_ptr">, Intrinsic<[LLVMQualPointerType], [], - [IntrNoMem, IntrSpeculatable]>; + [Align, IntrNoMem, IntrSpeculatable]>; def int_amdgcn_kernarg_segment_ptr : GCCBuiltin<"__builtin_amdgcn_kernarg_segment_ptr">, Intrinsic<[LLVMQualPointerType], [], - [IntrNoMem, IntrSpeculatable]>; + [Align, IntrNoMem, IntrSpeculatable]>; def int_amdgcn_implicitarg_ptr : GCCBuiltin<"__builtin_amdgcn_implicitarg_ptr">, Intrinsic<[LLVMQualPointerType], [], - [IntrNoMem, IntrSpeculatable]>; + [Align, IntrNoMem, IntrSpeculatable]>; def int_amdgcn_groupstaticsize : GCCBuiltin<"__builtin_amdgcn_groupstaticsize">, @@ -170,7 +170,7 @@ def int_amdgcn_implicit_buffer_ptr : GCCBuiltin<"__builtin_amdgcn_implicit_buffer_ptr">, Intrinsic<[LLVMQualPointerType], [], - [IntrNoMem, IntrSpeculatable]>; + [Align, 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 --- a/llvm/lib/IR/Attributes.cpp +++ b/llvm/lib/IR/Attributes.cpp @@ -1175,6 +1175,17 @@ } AttributeList AttributeList::get(LLVMContext &C, unsigned Index, + ArrayRef Kinds, + ArrayRef Values) { + assert(Kinds.size() == Values.size() && "Mismatched attribute values."); + SmallVector, 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 Kinds) { SmallVector, 8> Attrs; for (const auto &K : Kinds) diff --git a/llvm/test/CodeGen/AMDGPU/reqd-work-group-size.ll b/llvm/test/CodeGen/AMDGPU/reqd-work-group-size.ll --- a/llvm/test/CodeGen/AMDGPU/reqd-work-group-size.ll +++ b/llvm/test/CodeGen/AMDGPU/reqd-work-group-size.ll @@ -390,7 +390,7 @@ ; 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 @@ 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 --- a/llvm/utils/TableGen/CodeGenIntrinsics.h +++ b/llvm/utils/TableGen/CodeGenIntrinsics.h @@ -149,18 +149,21 @@ 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 --- a/llvm/utils/TableGen/CodeGenTarget.cpp +++ b/llvm/utils/TableGen/CodeGenTarget.cpp @@ -795,25 +795,29 @@ 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::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 --- a/llvm/utils/TableGen/IntrinsicEmitter.cpp +++ b/llvm/utils/TableGen/IntrinsicEmitter.cpp @@ -668,6 +668,8 @@ OS << " const Attribute::AttrKind AttrParam" << attrIdx << "[]= {"; bool addComma = false; + bool AllValuesAreZero = true; + SmallVector Values; do { switch (intrinsic.ArgumentAttributes[ai].Kind) { case CodeGenIntrinsic::NoCapture: @@ -712,13 +714,39 @@ 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"; } }