Intrinsics: Allow tablegen to mark parameters with dereferenceable

[arsenm]

Summary

Also increases the alignment of llvm.amdgcn.implicitarg.ptr to 8 to
match clang.

[arsenm]

Revert implicitarg.ptr changes since not-HSA has different alignment for no reason. Also with the size differences between amdhsa and different CO versions we're already wrong for emitting 256 unconditionally

[foad]

I think the tablegen functionality should be a separate patch from the amdgpu changes.

[arsenm]

In D142823#4093357, @foad wrote:

I think the tablegen functionality should be a separate patch from the amdgpu changes.

Maybe, but then it’s untested in the patch which adds it

[foad]

In D142823#4093363, @arsenm wrote:

In D142823#4093357, @foad wrote:

I think the tablegen functionality should be a separate patch from the amdgpu changes.

Maybe, but then it’s untested in the patch which adds it

Not if you add a test. There are some already like test/TableGen/immarg.td and test/TableGen/intrin-side-effects.td.

[arsenm]

Split out amdgpu parts

[nikic]

LGTM

[arsenm]

1717c18664d5880f78db85eb0075a2c1379df2d9

clang/lib/CodeGen/CGBuiltin.cpp

 
 namespace {
 // If \p E is not null pointer, insert address space cast to match return
 // type of \p E if necessary.
 Value *EmitAMDGPUDispatchPtr(CodeGenFunction &CGF,
                              const CallExpr *E = nullptr) {
   auto *F = CGF.CGM.getIntrinsic(Intrinsic::amdgcn_dispatch_ptr);
   auto *Call = CGF.Builder.CreateCall(F);
-  Call->addRetAttr(
-      Attribute::getWithDereferenceableBytes(Call->getContext(), 64));
-  Call->addRetAttr(Attribute::getWithAlignment(Call->getContext(), Align(4)));
   if (!E)
     return Call;
-  QualType BuiltinRetType = E->getType();
-  auto *RetTy = cast<llvm::PointerType>(CGF.ConvertType(BuiltinRetType));
-  if (RetTy == Call->getType())
-    return Call;
-  return CGF.Builder.CreateAddrSpaceCast(Call, RetTy);
+  return CGF.Builder.CreateAddrSpaceCast(Call, CGF.ConvertType(E->getType()));
 }
 
 Value *EmitAMDGPUImplicitArgPtr(CodeGenFunction &CGF) {
   auto *F = CGF.CGM.getIntrinsic(Intrinsic::amdgcn_implicitarg_ptr);
-  auto *Call = CGF.Builder.CreateCall(F);
-  Call->addRetAttr(
-      Attribute::getWithDereferenceableBytes(Call->getContext(), 256));
-  Call->addRetAttr(Attribute::getWithAlignment(Call->getContext(), Align(8)));
-  return Call;
+  return CGF.Builder.CreateCall(F);
 }
 
 // \p Index is 0, 1, and 2 for x, y, and z dimension, respectively.
 Value *EmitAMDGPUWorkGroupSize(CodeGenFunction &CGF, unsigned Index) {
   bool IsCOV_5 = CGF.getTarget().getTargetOpts().CodeObjectVersion ==
                  clang::TargetOptions::COV_5;
   Constant *Offset;
   Value *DP;

clang/test/CodeGenCUDA/amdgpu-workgroup-size.cu

 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \
 // RUN:     -fcuda-is-device -emit-llvm -o - -x hip %s \
-// RUN:     | FileCheck -check-prefix=PRECOV5 %s
+// RUN:     | FileCheck -check-prefixes=PRECOV5,CHECK %s
 
 
 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa \
 // RUN:     -fcuda-is-device -mcode-object-version=5 -emit-llvm -o - -x hip %s \
-// RUN:     | FileCheck -check-prefix=COV5 %s
+// RUN:     | FileCheck -check-prefixes=COV5,CHECK %s
 
 #include "Inputs/cuda.h"
 
 // PRECOV5-LABEL: test_get_workgroup_size
-// PRECOV5: call align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
+// PRECOV5: call ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
 // PRECOV5: getelementptr i8, ptr addrspace(4) %{{.*}}, i32 4
 // PRECOV5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef
 // PRECOV5: getelementptr i8, ptr addrspace(4) %{{.*}}, i32 6
 // PRECOV5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef
 // PRECOV5: getelementptr i8, ptr addrspace(4) %{{.*}}, i32 8
 // PRECOV5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef
 
 // COV5-LABEL: test_get_workgroup_size
-// COV5: call align 8 dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
+// COV5: call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
 // COV5: getelementptr i8, ptr addrspace(4) %{{.*}}, i32 12
 // COV5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef
 // COV5: getelementptr i8, ptr addrspace(4) %{{.*}}, i32 14
 // COV5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef
 // COV5: getelementptr i8, ptr addrspace(4) %{{.*}}, i32 16
 // COV5: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef
 __device__ void test_get_workgroup_size(int d, int *out)
 {
   switch (d) {
   case 0: *out = __builtin_amdgcn_workgroup_size_x(); break;
   case 1: *out = __builtin_amdgcn_workgroup_size_y(); break;
   case 2: *out = __builtin_amdgcn_workgroup_size_z(); break;
   default: *out = 0;
   }
 }
 
+// COV4: declare align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
+// COV5: declare align 8 dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
+
 // CHECK-DAG: [[$WS_RANGE]] = !{i16 1, i16 1025}

clang/test/CodeGenCUDA/builtins-amdgcn.cu

 // CHECK-NEXT:    [[DISPATCH_PTR:%.*]] = alloca ptr, align 8, addrspace(5)
 // CHECK-NEXT:    [[OUT_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT]] to ptr
 // CHECK-NEXT:    [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT_ADDR]] to ptr
 // CHECK-NEXT:    [[DISPATCH_PTR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[DISPATCH_PTR]] to ptr
 // CHECK-NEXT:    [[TMP0:%.*]] = addrspacecast ptr addrspace(1) [[OUT_COERCE:%.*]] to ptr
 // CHECK-NEXT:    store ptr [[TMP0]], ptr [[OUT_ASCAST]], align 8
 // CHECK-NEXT:    [[OUT1:%.*]] = load ptr, ptr [[OUT_ASCAST]], align 8
 // CHECK-NEXT:    store ptr [[OUT1]], ptr [[OUT_ADDR_ASCAST]], align 8
-// CHECK-NEXT:    [[TMP1:%.*]] = call align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
+// CHECK-NEXT:    [[TMP1:%.*]] = call ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
 // CHECK-NEXT:    [[TMP2:%.*]] = addrspacecast ptr addrspace(4) [[TMP1]] to ptr
 // CHECK-NEXT:    store ptr [[TMP2]], ptr [[DISPATCH_PTR_ASCAST]], align 8
 // CHECK-NEXT:    [[TMP3:%.*]] = load ptr, ptr [[DISPATCH_PTR_ASCAST]], align 8
 // CHECK-NEXT:    [[TMP4:%.*]] = load i32, ptr [[TMP3]], align 4
 // CHECK-NEXT:    [[TMP5:%.*]] = load ptr, ptr [[OUT_ADDR_ASCAST]], align 8
 // CHECK-NEXT:    store i32 [[TMP4]], ptr [[TMP5]], align 4
 // CHECK-NEXT:    ret void
 //

 __global__ void test_ds_fmin(float src, float *shared) {
   volatile float x = __builtin_amdgcn_ds_fminf(shared, src, 0, 0, false);
 }
 
 // CHECK-LABEL: @_Z33test_ret_builtin_nondef_addrspacev(
 // CHECK-NEXT:  entry:
 // CHECK-NEXT:    [[X:%.*]] = alloca ptr, align 8, addrspace(5)
 // CHECK-NEXT:    [[X_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[X]] to ptr
-// CHECK-NEXT:    [[TMP0:%.*]] = call align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
+// CHECK-NEXT:    [[TMP0:%.*]] = call ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
 // CHECK-NEXT:    [[TMP1:%.*]] = addrspacecast ptr addrspace(4) [[TMP0]] to ptr
 // CHECK-NEXT:    store ptr [[TMP1]], ptr [[X_ASCAST]], align 8
 // CHECK-NEXT:    ret void
 //
 __device__ void test_ret_builtin_nondef_addrspace() {
   void *x = __builtin_amdgcn_dispatch_ptr();
 }
 

clang/test/CodeGenOpenCL/builtins-amdgcn.cl

 
 // CHECK-LABEL: @test_read_exec_hi(
 // CHECK: call i32 @llvm.read_register.i32(metadata ![[$EXEC_HI:[0-9]+]]) #[[$READ_EXEC_ATTRS]]
 void test_read_exec_hi(global uint* out) {
   *out = __builtin_amdgcn_read_exec_hi();
 }
 
 // CHECK-LABEL: @test_dispatch_ptr
-// CHECK: call align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
+// CHECK: call ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
 void test_dispatch_ptr(__constant unsigned char ** out)
 {
   *out = __builtin_amdgcn_dispatch_ptr();
 }
 
+// CHECK: declare align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
+
+
 // CHECK-LABEL: @test_queue_ptr
 // CHECK: call ptr addrspace(4) @llvm.amdgcn.queue.ptr()
 void test_queue_ptr(__constant unsigned char ** out)
 {
   *out = __builtin_amdgcn_queue_ptr();
 }
 
 // CHECK-LABEL: @test_kernarg_segment_ptr
 // CHECK: call ptr addrspace(4) @llvm.amdgcn.kernarg.segment.ptr()
 void test_kernarg_segment_ptr(__constant unsigned char ** out)
 {
   *out = __builtin_amdgcn_kernarg_segment_ptr();
 }
 
 // CHECK-LABEL: @test_implicitarg_ptr
 // CHECK: call ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
 void test_implicitarg_ptr(__constant unsigned char ** out)
 {
   *out = __builtin_amdgcn_implicitarg_ptr();
 }
 
+// CHECK: declare align 8 dereferenceable(256) ptr addrspace(4) @llvm.amdgcn.implicitarg.ptr()
+
+
 // CHECK-LABEL: @test_get_group_id(
 // CHECK: tail call i32 @llvm.amdgcn.workgroup.id.x()
 // CHECK: tail call i32 @llvm.amdgcn.workgroup.id.y()
 // CHECK: tail call i32 @llvm.amdgcn.workgroup.id.z()
 void test_get_group_id(int d, global int *out)
 {
   switch (d) {
   case 0: *out = __builtin_amdgcn_workgroup_id_x(); break;
void test_get_local_id(int d, global int *out)
   case 0: *out = __builtin_amdgcn_workitem_id_x(); break;
   case 1: *out = __builtin_amdgcn_workitem_id_y(); break;
   case 2: *out = __builtin_amdgcn_workitem_id_z(); break;
   default: *out = 0;
   }
 }
 
 // CHECK-LABEL: @test_get_workgroup_size(
-// CHECK: call align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
+// CHECK: call ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
 // CHECK: getelementptr i8, ptr addrspace(4) %{{.*}}, i64 4
 // CHECK: load i16, ptr addrspace(4) %{{.*}}, align 4, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef
 // CHECK: getelementptr i8, ptr addrspace(4) %{{.*}}, i64 6
 // CHECK: load i16, ptr addrspace(4) %{{.*}}, align 2, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef
 // CHECK: getelementptr i8, ptr addrspace(4) %{{.*}}, i64 8
 // CHECK: load i16, ptr addrspace(4) %{{.*}}, align 4, !range [[$WS_RANGE:![0-9]*]], !invariant.load{{.*}}, !noundef
 void test_get_workgroup_size(int d, global int *out)
 {
   switch (d) {
   case 0: *out = __builtin_amdgcn_workgroup_size_x() + 1; break;
   case 1: *out = __builtin_amdgcn_workgroup_size_y(); break;
   case 2: *out = __builtin_amdgcn_workgroup_size_z(); break;
   default: *out = 0;
   }
 }
 
 // CHECK-LABEL: @test_get_grid_size(
-// CHECK: call align 4 dereferenceable(64) ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
+// CHECK: call ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
 // CHECK: getelementptr i8, ptr addrspace(4) %{{.*}}, i64 12
 // CHECK: load i32, ptr addrspace(4) %{{.*}}, align 4, !invariant.load
 // CHECK: getelementptr i8, ptr addrspace(4) %{{.*}}, i64 16
 // CHECK: load i32, ptr addrspace(4) %{{.*}}, align 4, !invariant.load
 // CHECK: getelementptr i8, ptr addrspace(4) %{{.*}}, i64 20
 // CHECK: load i32, ptr addrspace(4) %{{.*}}, align 4, !invariant.load
 void test_get_grid_size(int d, global int *out)
 {

llvm/include/llvm/IR/Intrinsics.td

class NonNull<AttrIndex idx> : IntrinsicProperty {
   int ArgNo = idx.Value;
 }
 
 class Align<AttrIndex idx, int align> : IntrinsicProperty {
   int ArgNo = idx.Value;
   int Align = align;
 }
 
+class Dereferenceable<AttrIndex idx, int bytes> : IntrinsicProperty {
+  int ArgNo = idx.Value;
+  int Bytes = bytes;
+}
+
 // 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

 // ABI Special Intrinsics
 //===----------------------------------------------------------------------===//
 
 defm int_amdgcn_workitem_id : AMDGPUReadPreloadRegisterIntrinsic_xyz;
 defm int_amdgcn_workgroup_id : AMDGPUReadPreloadRegisterIntrinsic_xyz_named
                                <"__builtin_amdgcn_workgroup_id">;
 
 def int_amdgcn_dispatch_ptr :
+  ClangBuiltin<"__builtin_amdgcn_dispatch_ptr">,
   DefaultAttrsIntrinsic<[LLVMQualPointerType<llvm_i8_ty, 4>], [],
-  [Align<RetIndex, 4>, IntrNoMem, IntrSpeculatable]>;
+  [Align<RetIndex, 4>, Dereferenceable<RetIndex, 64>, IntrNoMem,
+   IntrSpeculatable]>;
 
 def int_amdgcn_queue_ptr :
   ClangBuiltin<"__builtin_amdgcn_queue_ptr">,
   DefaultAttrsIntrinsic<[LLVMQualPointerType<llvm_i8_ty, 4>], [],
   [Align<RetIndex, 4>, IntrNoMem, IntrSpeculatable]>;
 
 def int_amdgcn_kernarg_segment_ptr :
   ClangBuiltin<"__builtin_amdgcn_kernarg_segment_ptr">,
   DefaultAttrsIntrinsic<[LLVMQualPointerType<llvm_i8_ty, 4>], [],
   [Align<RetIndex, 4>, IntrNoMem, IntrSpeculatable]>;
 
 def int_amdgcn_implicitarg_ptr :
   ClangBuiltin<"__builtin_amdgcn_implicitarg_ptr">,
   DefaultAttrsIntrinsic<[LLVMQualPointerType<llvm_i8_ty, 4>], [],
-  [Align<RetIndex, 4>, IntrNoMem, IntrSpeculatable]>;
+  [Align<RetIndex, 8>, Dereferenceable<RetIndex, 256>,
+   IntrNoMem, IntrSpeculatable]>;
 
 def int_amdgcn_groupstaticsize :
   ClangBuiltin<"__builtin_amdgcn_groupstaticsize">,
   DefaultAttrsIntrinsic<[llvm_i32_ty], [], [IntrNoMem, IntrSpeculatable]>;
 
 def int_amdgcn_dispatch_id :
   ClangBuiltin<"__builtin_amdgcn_dispatch_id">,
   DefaultAttrsIntrinsic<[llvm_i64_ty], [], [IntrNoMem, IntrSpeculatable]>;

llvm/utils/TableGen/CodeGenIntrinsics.h

enum ArgAttrKind {
     NoAlias,
     NoUndef,
     NonNull,
     Returned,
     ReadOnly,
     WriteOnly,
     ReadNone,
     ImmArg,
-    Alignment
+    Alignment,
+    Dereferenceable
   };
 
   struct ArgAttribute {
     ArgAttrKind Kind;
     uint64_t Value;
 
     ArgAttribute(ArgAttrKind K, uint64_t V) : Kind(K), Value(V) {}
 

llvm/utils/TableGen/CodeGenTarget.cpp

else if (R->isSubClassOf("NoCapture")) {
     addArgAttribute(ArgNo, ReadNone);
   } else if (R->isSubClassOf("ImmArg")) {
     unsigned ArgNo = R->getValueAsInt("ArgNo");
     addArgAttribute(ArgNo, ImmArg);
   } else if (R->isSubClassOf("Align")) {
     unsigned ArgNo = R->getValueAsInt("ArgNo");
     uint64_t Align = R->getValueAsInt("Align");
     addArgAttribute(ArgNo, Alignment, Align);
+  } else if (R->isSubClassOf("Dereferenceable")) {
+    unsigned ArgNo = R->getValueAsInt("ArgNo");
+    uint64_t Bytes = R->getValueAsInt("Bytes");
+    addArgAttribute(ArgNo, Dereferenceable, Bytes);
   } else
     llvm_unreachable("Unknown property!");
 }
 
 bool CodeGenIntrinsic::isParamAPointer(unsigned ParamIdx) const {
   if (ParamIdx >= IS.ParamVTs.size())
     return false;
   MVT ParamType = MVT(IS.ParamVTs[ParamIdx]);

llvm/utils/TableGen/IntrinsicEmitter.cpp

for (auto &Attrs : Int.ArgumentAttributes) {
           break;
         case CodeGenIntrinsic::ImmArg:
           OS << "      Attribute::get(C, Attribute::ImmArg),\n";
           break;
         case CodeGenIntrinsic::Alignment:
           OS << "      Attribute::get(C, Attribute::Alignment, "
              << Attr.Value << "),\n";
           break;
+        case CodeGenIntrinsic::Dereferenceable:
+          OS << "      Attribute::get(C, Attribute::Dereferenceable, "
+             << Attr.Value << "),\n";
+          break;
         }
       }
       OS << "    });\n";
     }
   }
   OS << "  }\n";
   OS << "}\n\n";