Enable `align <n>` to be used in intrinsic definitions.

[hliao]

Summary

• This enable specifying the alignment on the return value as well as parameters.

[hliao]

To prepare the refactoring on D80364, intrinsics interested should be specified with the alignment on the return pointer. With this patch, amdgcn.dispatch.ptr is defined as follows

def int_amdgcn_dispatch_ptr :
  Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 4>], [],
  [Align<-1, 4>, IntrNoMem, IntrSpeculatable]>;

[arsenm]

Should also add this to at least one of the amdgpu intrinsics, and add some tests for it. In particular I would like to see a test that shows an explicit alignment on a call site will override the intrinsic's alignment

[arsenm]

In D80422#2050474, @hliao wrote:

To prepare the refactoring on D80364, intrinsics interested should be specified with the alignment on the return pointer. With this patch, amdgcn.dispatch.ptr is defined as follows

def int_amdgcn_dispatch_ptr :
  Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 4>], [],
  [Align<-1, 4>, IntrNoMem, IntrSpeculatable]>;

Referring to the return index by a negative parameter index is a bit weird. How does this work for multiple return values? Is it [-2, -1]? This needs documenting somewhere

[hliao]

In D80422#2050863, @arsenm wrote:

In D80422#2050474, @hliao wrote:

To prepare the refactoring on D80364, intrinsics interested should be specified with the alignment on the return pointer. With this patch, amdgcn.dispatch.ptr is defined as follows

def int_amdgcn_dispatch_ptr :
  Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 4>], [],
  [Align<-1, 4>, IntrNoMem, IntrSpeculatable]>;

Referring to the return index by a negative parameter index is a bit weird. How does this work for multiple return values? Is it [-2, -1]? This needs documenting somewhere

There's no such thing. Multiple return values are always aggregated into a single struct and there's only one return value. Different from parameter indices, the return index for the attribute list is just 0 only.

351 class AttributeList {                                                                                                                                                                                                                                     
352 public:                                                                                                                                                                                                                                                   
353   enum AttrIndex : unsigned {                                                                                                                                                                                                                             
354     ReturnIndex = 0U,                                                                                                                                                                                                                                     
355     FunctionIndex = ~0U,                                                                                                                                                                                                                                  
356     FirstArgIndex = 1,                                                                                                                                                                                                                                    
357   };

[hliao]

Add tests & example changes in relevant AMDGPU intrinsics.

[hliao]

Revise the comment in the intrinsic emitter.

[hliao]

Follow part of clang-tidy the suggestion on coding style.

[arsenm]

Usually the attributes of a call are the union of (call site attributes, declaration attribute), but in this case it's a bit different since the call site will override it since it happens to be checked first. This is probably fine, but I'm also wondering if defining this as a minimum alignment might be better

[jdoerfert]

In D80422#2055938, @arsenm wrote:

Usually the attributes of a call are the union of (call site attributes, declaration attribute), but in this case it's a bit different since the call site will override it since it happens to be checked first. This is probably fine, but I'm also wondering if defining this as a minimum alignment might be better

Not sure I understand. If you query attributes you'll get the first that is found (I think), so if the call site provides one you get that, otherwise the function one, if present. I haven't seen anything here that would be "weird". Wrt. minimum alignment, I guess that is by definition the case. align(X) anywhere means minimum alignment and the Attributor would for example bump call site alignment to the function alignment if that is known to be higher.

I added some minor comments below. From my perspective this is fine, assuming I did understand @arsenm properly. If so, feel free to accept :)

llvm/include/llvm/IR/Intrinsics.td

86

Nit: maybe uint64_t for the align to match other places, though it will work fine as long as int is 32 bit (or more).

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

173

I think -1 is in principle the right thing here *but* we should try to make it a variable with a proper name, e.g. ReturnIndex.

llvm/lib/IR/Attributes.cpp

1186

Nit: Message for the assert please.

llvm/utils/TableGen/CodeGenIntrinsics.h

156–170

Every time I see a use of tuple I ask myself the same question: Wouldn't a POD-like struct with some properly named members be better? I feel std::get<1>(T) is somewhat less helpful than T->Index. I don't need you to change it here but it's something to think about. In particular it might help to remove magic numbers below, e.g., 0 for the value.

[xbolva00]

llvm/utils/TableGen/CodeGenIntrinsics.h

156–170

+1

[hliao]

• Introduce RetIndex and ArgIndex<n> to specify attribute indices.
• Minor refinement following other review comments.

[hliao]

llvm/include/llvm/IR/Intrinsics.td

86

int in td is represented as int64_t internally.

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

173

introduce RetIndex and ArgIndex<n> to specifiy that attribute indices.

[hliao]

remove <tuple> header.

[arsenm]

LGTM. I would split the ArgIndex refactor into a precommit but up to you

llvm/include/llvm/IR/Attributes.h

public:
   // AttributeList Construction and Mutation
   //===--------------------------------------------------------------------===//
 
   /// Return an AttributeList with the specified parameters in it.
   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,
                            const AttrBuilder &B);
 
   /// Add an attribute to the attribute set at the given index.
   /// Returns a new list because attribute lists are immutable.
   LLVM_NODISCARD AttributeList addAttribute(LLVMContext &C, unsigned Index,
                                             Attribute::AttrKind Kind) const;

llvm/include/llvm/IR/Intrinsics.td

 }
 
 // NoAlias - The specified argument pointer is not aliasing other "noalias" pointer
 // arguments of the intrinsic wrt. the intrinsic scope.
 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 {
   int ArgNo = idx.Value;
 }
 
 // ImmArg - The specified argument must be an immediate.
 class ImmArg<AttrIndex idx> : IntrinsicProperty {

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

 //===----------------------------------------------------------------------===//
 
 defm int_amdgcn_workitem_id : AMDGPUReadPreloadRegisterIntrinsic_xyz;
 defm int_amdgcn_workgroup_id : AMDGPUReadPreloadRegisterIntrinsic_xyz_named
                                <"__builtin_amdgcn_workgroup_id">;
 
 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">,
   Intrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable]>;
 
 def int_amdgcn_dispatch_id :
   GCCBuiltin<"__builtin_amdgcn_dispatch_id">,
   Intrinsic<[llvm_i64_ty], [], [IntrNoMem, IntrSpeculatable]>;
 
 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.
 // FIXME: Should be mangled for wave size.
 def int_amdgcn_init_exec : Intrinsic<[],
   [llvm_i64_ty],      // 64-bit literal constant
   [IntrConvergent, ImmArg<ArgIndex<0>>]>;
 

llvm/lib/IR/Attributes.cpp

AttributeList AttributeList::get(LLVMContext &C, unsigned Index,
                                  ArrayRef<Attribute::AttrKind> Kinds) {
   SmallVector<std::pair<unsigned, Attribute>, 8> Attrs;
   for (const auto K : Kinds)
     Attrs.emplace_back(Index, Attribute::get(C, K));
   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;
   for (const auto &K : Kinds)
     Attrs.emplace_back(Index, Attribute::get(C, K));
   return get(C, Attrs);
 }
 
 AttributeList AttributeList::get(LLVMContext &C,

llvm/test/CodeGen/AMDGPU/reqd-work-group-size.ll

define amdgpu_kernel void @all_local_size(i64 addrspace(1)* nocapture readnone %out) #0 !reqd_work_group_size !0 {
   store volatile i64 %tmp34.i7, i64 addrspace(1)* %out, align 4
   ret void
 }
 
 ; TODO: Should be able to handle this, but not much reason to.
 ; 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()
   %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
 }
 
+; 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,
 ; CHECK: store i32 %group.size.xy
 define amdgpu_kernel void @load_group_size_xy_i32(i32 addrspace(1)* %out) #0 !reqd_work_group_size !0 {
   %dispatch.ptr = tail call i8 addrspace(4)* @llvm.amdgcn.dispatch.ptr()
   %gep.group.size.x = getelementptr inbounds i8, i8 addrspace(4)* %dispatch.ptr, i64 4
   %gep.group.size.x.bc = bitcast i8 addrspace(4)* %gep.group.size.x to i32 addrspace(4)*

llvm/utils/TableGen/CodeGenIntrinsics.h

struct CodeGenIntrinsic {
 
   enum ArgAttrKind {
     NoCapture,
     NoAlias,
     Returned,
     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);
     }
   };
 
   std::vector<ArgAttribute> ArgumentAttributes;
 
   bool hasProperty(enum SDNP Prop) const {
     return Properties & (1 << Prop);
   }
 
   /// Returns true if the parameter at \p ParamIdx is a pointer type. Returns
   /// false if the parameter is not a pointer, or \p ParamIdx is greater than
   /// the size of \p IS.ParamVTs.

llvm/utils/TableGen/CodeGenTarget.cpp

for (unsigned i = 0, e = PropList->size(); i != e; ++i) {
     else if (Property->getName() == "IntrCold")
       isCold = true;
     else if (Property->getName() == "IntrSpeculatable")
       isSpeculatable = true;
     else if (Property->getName() == "IntrHasSideEffects")
       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!");
   }
 
   // Also record the SDPatternOperator Properties.
   Properties = parseSDPatternOperatorProperties(R);
 
   // Sort the argument attributes for later benefit.
   llvm::sort(ArgumentAttributes);
 }
 
 bool CodeGenIntrinsic::isParamAPointer(unsigned ParamIdx) const {
   if (ParamIdx >= IS.ParamVTs.size())
     return false;
   MVT ParamType = MVT(IS.ParamVTs[ParamIdx]);
   return ParamType == MVT::iPTR || ParamType == MVT::iPTRAny;
 }
 
 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);
 }

llvm/utils/TableGen/IntrinsicEmitter.cpp

for (UniqAttrMapTy::const_iterator I = UniqAttributes.begin(),
     unsigned ai = 0, ae = intrinsic.ArgumentAttributes.size();
     if (ae) {
       while (ai != ae) {
         unsigned attrIdx = intrinsic.ArgumentAttributes[ai].Index;
 
         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:
             if (addComma)
               OS << ",";
             OS << "Attribute::NoCapture";
             addComma = true;
             break;
if (ae) {
             addComma = true;
             break;
           case CodeGenIntrinsic::ImmArg:
             if (addComma)
               OS << ',';
             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";
       }
     }
 
     if (!intrinsic.canThrow ||
         (intrinsic.ModRef != CodeGenIntrinsic::ReadWriteMem && !intrinsic.hasSideEffects) ||
         intrinsic.isNoReturn || intrinsic.isNoSync || intrinsic.isWillReturn ||
         intrinsic.isCold || intrinsic.isNoDuplicate || intrinsic.isConvergent ||
         intrinsic.isSpeculatable) {