maxjhandsome
/
llvm-project
forked from OSchip/llvm-project
1
0
Fork 0
Code Issues Pull Requests Packages Releases Wiki Activity

Make __builtin_amdgcn_dispatch_ptr dereferenceable and align at 4 

Differential Revision: https://reviews.llvm.org/D75028
This commit is contained in:
Yaxun (Sam) Liu 2020-02-20 19:02:53 -05:00
parent eee22ec3c3
commit a57d9652a0
4 changed files with 18 additions and 4 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

15
clang/lib/CodeGen/CGBuiltin.cpp
View File

@ -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:

4
clang/test/CodeGenCUDA/builtins-amdgcn.cu
View File

@ -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;

2
clang/test/CodeGenOpenCL/builtins-amdgcn.cl
View File

@ -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();

1
llvm/include/llvm/IR/IntrinsicsAMDGPU.td
View File

@ -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]>;