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 @@ -396,6 +396,9 @@ static AttributeList get(LLVMContext &C, ArrayRef Attrs); 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, 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 @@ -73,6 +73,11 @@ int ArgNo = argNo; } +class Align : IntrinsicProperty { + int ArgNo = argNo; + 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<-1, 4>, IntrNoMem, IntrSpeculatable]>; def int_amdgcn_queue_ptr : GCCBuiltin<"__builtin_amdgcn_queue_ptr">, Intrinsic<[LLVMQualPointerType], [], - [IntrNoMem, IntrSpeculatable]>; + [Align<-1, 4>, IntrNoMem, IntrSpeculatable]>; def int_amdgcn_kernarg_segment_ptr : GCCBuiltin<"__builtin_amdgcn_kernarg_segment_ptr">, Intrinsic<[LLVMQualPointerType], [], - [IntrNoMem, IntrSpeculatable]>; + [Align<-1, 4>, IntrNoMem, IntrSpeculatable]>; def int_amdgcn_implicitarg_ptr : GCCBuiltin<"__builtin_amdgcn_implicitarg_ptr">, Intrinsic<[LLVMQualPointerType], [], - [IntrNoMem, IntrSpeculatable]>; + [Align<-1, 4>, 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<-1, 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 --- a/llvm/lib/IR/Attributes.cpp +++ b/llvm/lib/IR/Attributes.cpp @@ -1174,6 +1174,17 @@ return get(C, Attrs); } +AttributeList AttributeList::get(LLVMContext &C, unsigned Index, + ArrayRef Kinds, + ArrayRef Values) { + assert(Kinds.size() == Values.size()); + 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; 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 @@ -16,6 +16,7 @@ #include "SDNodeProperties.h" #include "llvm/Support/MachineValueType.h" #include +#include #include namespace llvm { @@ -149,10 +150,11 @@ ReadOnly, WriteOnly, ReadNone, - ImmArg + ImmArg, + Alignment }; - std::vector> ArgumentAttributes; + std::vector> ArgumentAttributes; bool hasProperty(enum SDNP Prop) const { return Properties & (1 << Prop); 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 @@ -25,6 +25,7 @@ #include "llvm/TableGen/Record.h" #include "llvm/TableGen/TableGenBackend.h" #include +#include using namespace llvm; cl::OptionCategory AsmParserCat("Options for -gen-asm-parser"); @@ -785,25 +786,29 @@ hasSideEffects = true; else if (Property->isSubClassOf("NoCapture")) { unsigned ArgNo = Property->getValueAsInt("ArgNo"); - ArgumentAttributes.push_back(std::make_pair(ArgNo, NoCapture)); + ArgumentAttributes.push_back(std::make_tuple(ArgNo, NoCapture, 0)); } else if (Property->isSubClassOf("NoAlias")) { unsigned ArgNo = Property->getValueAsInt("ArgNo"); - ArgumentAttributes.push_back(std::make_pair(ArgNo, NoAlias)); + ArgumentAttributes.push_back(std::make_tuple(ArgNo, NoAlias, 0)); } else if (Property->isSubClassOf("Returned")) { unsigned ArgNo = Property->getValueAsInt("ArgNo"); - ArgumentAttributes.push_back(std::make_pair(ArgNo, Returned)); + ArgumentAttributes.push_back(std::make_tuple(ArgNo, Returned, 0)); } else if (Property->isSubClassOf("ReadOnly")) { unsigned ArgNo = Property->getValueAsInt("ArgNo"); - ArgumentAttributes.push_back(std::make_pair(ArgNo, ReadOnly)); + ArgumentAttributes.push_back(std::make_tuple(ArgNo, ReadOnly, 0)); } else if (Property->isSubClassOf("WriteOnly")) { unsigned ArgNo = Property->getValueAsInt("ArgNo"); - ArgumentAttributes.push_back(std::make_pair(ArgNo, WriteOnly)); + ArgumentAttributes.push_back(std::make_tuple(ArgNo, WriteOnly, 0)); } else if (Property->isSubClassOf("ReadNone")) { unsigned ArgNo = Property->getValueAsInt("ArgNo"); - ArgumentAttributes.push_back(std::make_pair(ArgNo, ReadNone)); + ArgumentAttributes.push_back(std::make_tuple(ArgNo, ReadNone, 0)); } else if (Property->isSubClassOf("ImmArg")) { unsigned ArgNo = Property->getValueAsInt("ArgNo"); - ArgumentAttributes.push_back(std::make_pair(ArgNo, ImmArg)); + ArgumentAttributes.push_back(std::make_tuple(ArgNo, ImmArg, 0)); + } else if (Property->isSubClassOf("Align")) { + unsigned ArgNo = Property->getValueAsInt("ArgNo"); + unsigned Align = Property->getValueAsInt("Align"); + ArgumentAttributes.push_back(std::make_tuple(ArgNo, Alignment, Align)); } else llvm_unreachable("Unknown property!"); } @@ -823,7 +828,7 @@ } bool CodeGenIntrinsic::isParamImmArg(unsigned ParamIdx) const { - std::pair Val = {ParamIdx, ImmArg}; + std::tuple Val = {ParamIdx, 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 @@ -22,6 +22,7 @@ #include "llvm/TableGen/StringToOffsetTable.h" #include "llvm/TableGen/TableGenBackend.h" #include +#include using namespace llvm; cl::OptionCategory GenIntrinsicCat("Options for -gen-intrinsic-enums"); @@ -661,14 +662,16 @@ unsigned ai = 0, ae = intrinsic.ArgumentAttributes.size(); if (ae) { while (ai != ae) { - unsigned argNo = intrinsic.ArgumentAttributes[ai].first; + unsigned argNo = std::get<0>(intrinsic.ArgumentAttributes[ai]); unsigned attrIdx = argNo + 1; // Must match AttributeList::FirstArgIndex OS << " const Attribute::AttrKind AttrParam" << attrIdx << "[]= {"; bool addComma = false; + bool allZeroValues = true; + SmallVector Values; do { - switch (intrinsic.ArgumentAttributes[ai].second) { + switch (std::get<1>(intrinsic.ArgumentAttributes[ai])) { case CodeGenIntrinsic::NoCapture: if (addComma) OS << ","; @@ -711,13 +714,39 @@ OS << "Attribute::ImmArg"; addComma = true; break; + case CodeGenIntrinsic::Alignment: + if (addComma) + OS << ','; + OS << "Attribute::Alignment"; + addComma = true; + break; } + uint64_t V = std::get<2>(intrinsic.ArgumentAttributes[ai]); + Values.push_back(V); + allZeroValues &= (V == 0); ++ai; - } while (ai != ae && intrinsic.ArgumentAttributes[ai].first == argNo); + } while (ai != ae && + std::get<0>(intrinsic.ArgumentAttributes[ai]) == argNo); OS << "};\n"; + + if (!allZeroValues) { + 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 (!allZeroValues) + OS << ", AttrValParam" << attrIdx; + OS << ");\n"; } }