Author: Stanislav Mekhanoshin Date: 2024-10-22T12:32:08-07:00 New Revision: 03fef62b84469c5dbbed04235c30eb96b6b48369
URL: https://github.com/llvm/llvm-project/commit/03fef62b84469c5dbbed04235c30eb96b6b48369 DIFF: https://github.com/llvm/llvm-project/commit/03fef62b84469c5dbbed04235c30eb96b6b48369.diff LOG: [AMDGPU] Relax __builtin_amdgcn_update_dpp sema check (#113341) Recent change applied too strict check for old and src operands match. These shall be compatible, but not necessarily exactly the same. Fixes: SWDEV-493072 Added: Modified: clang/lib/Sema/SemaAMDGPU.cpp clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl clang/test/SemaOpenCL/builtins-amdgcn-error-gfx9.cl Removed: ################################################################################ diff --git a/clang/lib/Sema/SemaAMDGPU.cpp b/clang/lib/Sema/SemaAMDGPU.cpp index 9e05e8f28b2c08..f59654c14f08fb 100644 --- a/clang/lib/Sema/SemaAMDGPU.cpp +++ b/clang/lib/Sema/SemaAMDGPU.cpp @@ -93,13 +93,19 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID, return true; } } - if (ArgTys[0] != ArgTys[1]) { - SemaRef.Diag(Args[1]->getBeginLoc(), - diag::err_typecheck_call_ diff erent_arg_types) - << ArgTys[0] << ArgTys[1]; - return true; - } - return false; + if (getASTContext().hasSameUnqualifiedType(ArgTys[0], ArgTys[1])) + return false; + if (((ArgTys[0]->isUnsignedIntegerType() && + ArgTys[1]->isSignedIntegerType()) || + (ArgTys[0]->isSignedIntegerType() && + ArgTys[1]->isUnsignedIntegerType())) && + getASTContext().getTypeSize(ArgTys[0]) == + getASTContext().getTypeSize(ArgTys[1])) + return false; + SemaRef.Diag(Args[1]->getBeginLoc(), + diag::err_typecheck_call_ diff erent_arg_types) + << ArgTys[0] << ArgTys[1]; + return true; } default: return false; diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl index 65b54c1d552742..269f20e2f53fe1 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl @@ -218,6 +218,29 @@ void test_update_dpp_half(half *x, global half *p) { *p = __builtin_amdgcn_update_dpp(*x, *x, 0x101, 0xf, 0xf, 0); } +// CHECK-LABEL: @test_update_dpp_int_uint +// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.update.dpp.i32(i32 %arg1, i32 %arg2, i32 0, i32 0, i32 0, i1 false) +void test_update_dpp_int_uint(global int* out, int arg1, unsigned int arg2) +{ + *out = __builtin_amdgcn_update_dpp(arg1, arg2, 0, 0, 0, false); +} + +// CHECK-LABEL: @test_update_dpp_lit_int +// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.update.dpp.i32(i32 5, i32 %arg1, i32 0, i32 0, i32 0, i1 false) +void test_update_dpp_lit_int(global int* out, int arg1) +{ + *out = __builtin_amdgcn_update_dpp(5, arg1, 0, 0, 0, false); +} + +__constant int gi = 5; + +// CHECK-LABEL: @test_update_dpp_const_int +// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.update.dpp.i32(i32 5, i32 %arg1, i32 0, i32 0, i32 0, i1 false) +void test_update_dpp_const_int(global int* out, int arg1) +{ + *out = __builtin_amdgcn_update_dpp(gi, arg1, 0, 0, 0, false); +} + // CHECK-LABEL: @test_ds_fadd // CHECK: atomicrmw fadd ptr addrspace(3) %out, float %src monotonic, align 4{{$}} // CHECK: atomicrmw volatile fadd ptr addrspace(3) %out, float %src monotonic, align 4{{$}} diff --git a/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx9.cl b/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx9.cl index 47b56c703e4c9d..7c07632aeb60b7 100644 --- a/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx9.cl +++ b/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx9.cl @@ -56,4 +56,5 @@ void test_update_dpp(global int* out, int arg1, int arg2, int i, int2 i2, long l *out = __builtin_amdgcn_update_dpp(fc, arg2, 0, 0, 0, false); // expected-error{{used type '__private _Complex float' where integer or floating point type is required}} *out = __builtin_amdgcn_update_dpp(arg1, fc, 0, 0, 0, false); // expected-error{{used type '__private _Complex float' where integer or floating point type is required}} *out = __builtin_amdgcn_update_dpp(i, l, 0, 0, 0, false); // expected-error{{arguments are of diff erent types ('__private int' vs '__private long')}} + *out = __builtin_amdgcn_update_dpp(0.5f, i, 0, 0, 0, false); // expected-error{{arguments are of diff erent types ('float' vs '__private int')}} } _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
