[clang] 03fef62 - [AMDGPU] Relax __builtin_amdgcn_update_dpp sema check (#113341)

[via cfe-commits]

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')}}
 }