Author: Yaxun (Sam) Liu Date: 2020-02-25T13:58:20-05:00 New Revision: a57d9652a0dcc823921f2d4bac29680db5dbef64
URL: https://github.com/llvm/llvm-project/commit/a57d9652a0dcc823921f2d4bac29680db5dbef64 DIFF: https://github.com/llvm/llvm-project/commit/a57d9652a0dcc823921f2d4bac29680db5dbef64.diff LOG: Make __builtin_amdgcn_dispatch_ptr dereferenceable and align at 4 Differential Revision: https://reviews.llvm.org/D75028 Added: Modified: clang/lib/CodeGen/CGBuiltin.cpp clang/test/CodeGenCUDA/builtins-amdgcn.cu clang/test/CodeGenOpenCL/builtins-amdgcn.cl llvm/include/llvm/IR/IntrinsicsAMDGPU.td Removed: ################################################################################ diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp index db58738c2701..47b3abdc5fac 100644 --- a/clang/lib/CodeGen/CGBuiltin.cpp +++ b/clang/lib/CodeGen/CGBuiltin.cpp @@ -13292,6 +13292,21 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID, 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: diff --git a/clang/test/CodeGenCUDA/builtins-amdgcn.cu b/clang/test/CodeGenCUDA/builtins-amdgcn.cu index 409a917b352e..5469e78ea101 100644 --- a/clang/test/CodeGenCUDA/builtins-amdgcn.cu +++ b/clang/test/CodeGenCUDA/builtins-amdgcn.cu @@ -2,8 +2,8 @@ #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; diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl index 6cda2a767d94..85e921cbe12a 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl @@ -461,7 +461,7 @@ void test_read_exec_hi(global uint* out) { } // 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(); diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td index 207b5b55e4bd..32be19109bb4 100644 --- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td +++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td @@ -141,7 +141,6 @@ 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]>; _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits