Make __builtin_amdgcn_dispatch_ptr dereferenceable and align at 4

[arsenm]

LGTM

[arsenm]

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

144

I don't think you actually need to delete this. It should be fine to leave it

[yaxunl]

llvm/include/llvm/IR/IntrinsicsAMDGPU.td

144

I tried. If I leave this, clang will use the standard handling of target builtin functions for builtin_amdgcn_dispatch_ptr and will not reach EmitAMDGPUBuiltinExpr to handle builtin_amdgcn_dispatch_ptr.

clang/lib/CodeGen/CGBuiltin.cpp

Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
   case AMDGPU::BI__builtin_amdgcn_rsq_clampf:
     return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_rsq_clamp);
   case AMDGPU::BI__builtin_amdgcn_sinf:
   case AMDGPU::BI__builtin_amdgcn_sinh:
     return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_sin);
   case AMDGPU::BI__builtin_amdgcn_cosf:
   case AMDGPU::BI__builtin_amdgcn_cosh:
     return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_cos);
+  case AMDGPU::BI__builtin_amdgcn_dispatch_ptr: {
+    auto *F = CGM.getIntrinsic(Intrinsic::amdgcn_dispatch_ptr);
+    auto *Call = Builder.CreateCall(F);
+    Call->addAttribute(
+        AttributeList::ReturnIndex,
+        Attribute::getWithDereferenceableBytes(Call->getContext(), 64));
+    Call->addAttribute(
+        AttributeList::ReturnIndex,
+        Attribute::getWithAlignment(Call->getContext(), Align(4)));
+    QualType BuiltinRetType = E->getType();
+    auto *RetTy = cast<llvm::PointerType>(ConvertType(BuiltinRetType));
+    if (RetTy == Call->getType())
+      return Call;
+    return Builder.CreateAddrSpaceCast(Call, RetTy);
+  }
   case AMDGPU::BI__builtin_amdgcn_log_clampf:
     return emitUnaryBuiltin(*this, E, Intrinsic::amdgcn_log_clamp);
   case AMDGPU::BI__builtin_amdgcn_ldexp:
   case AMDGPU::BI__builtin_amdgcn_ldexpf:
   case AMDGPU::BI__builtin_amdgcn_ldexph:
     return emitFPIntBuiltin(*this, E, Intrinsic::amdgcn_ldexp);
   case AMDGPU::BI__builtin_amdgcn_frexp_mant:
   case AMDGPU::BI__builtin_amdgcn_frexp_mantf:

clang/test/CodeGenCUDA/builtins-amdgcn.cu

 // RUN: %clang_cc1 -triple amdgcn -fcuda-is-device -emit-llvm %s -o - | FileCheck %s
 #include "Inputs/cuda.h"
 
 // CHECK-LABEL: @_Z16use_dispatch_ptrPi(
-// CHECK: %[[PTR:.*]] = call i8 addrspace(4)* @llvm.amdgcn.dispatch.ptr()
-// CHECK: %{{.*}} = addrspacecast i8 addrspace(4)* %[[PTR]] to i8 addrspace(4)**
+// CHECK: %[[PTR:.*]] = call align 4 dereferenceable(64) i8 addrspace(4)* @llvm.amdgcn.dispatch.ptr()
+// CHECK: %{{.*}} = addrspacecast i8 addrspace(4)* %[[PTR]] to i8*
 __global__ void use_dispatch_ptr(int* out) {
   const int* dispatch_ptr = (const int*)__builtin_amdgcn_dispatch_ptr();
   *out = *dispatch_ptr;
 }
 
 // CHECK-LABEL: @_Z12test_ds_fmaxf(
 // CHECK: call float @llvm.amdgcn.ds.fmax(float addrspace(3)* @_ZZ12test_ds_fmaxfE6shared, float %{{[^,]*}}, i32 0, i32 0, i1 false)
 __global__
 void test_ds_fmax(float src) {
   __shared__ float shared;
   volatile float x = __builtin_amdgcn_ds_fmaxf(&shared, src, 0, 0, false);
 }

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 i8 addrspace(4)* @llvm.amdgcn.dispatch.ptr()
+// CHECK: call align 4 dereferenceable(64) i8 addrspace(4)* @llvm.amdgcn.dispatch.ptr()
 void test_dispatch_ptr(__constant unsigned char ** out)
 {
   *out = __builtin_amdgcn_dispatch_ptr();
 }
 
 // CHECK-LABEL: @test_queue_ptr
 // CHECK: call i8 addrspace(4)* @llvm.amdgcn.queue.ptr()
 void test_queue_ptr(__constant unsigned char ** out)

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 :
-  GCCBuiltin<"__builtin_amdgcn_dispatch_ptr">,
   Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 4>], [],
   [IntrNoMem, IntrSpeculatable]>;
 
 def int_amdgcn_queue_ptr :
   GCCBuiltin<"__builtin_amdgcn_queue_ptr">,
   Intrinsic<[LLVMQualPointerType<llvm_i8_ty, 4>], [],
   [IntrNoMem, IntrSpeculatable]>;