[AMDGPU] Relax __builtin_amdgcn_update_dpp sema check

[rampitec]

Recent change applied too strict check for old and src operands match. These shall be compatible, but not necessarily exactly the same.

[llvmbot]

@llvm/pr-subscribers-backend-amdgpu

@llvm/pr-subscribers-clang

Author: Stanislav Mekhanoshin (rampitec)

Changes

Recent change applied too strict check for old and src operands match. These shall be compatible, but not necessarily exactly the same.

---

Full diff: https://github.com/llvm/llvm-project/pull/113341.diff

3 Files Affected:

• (modified) clang/lib/Sema/SemaAMDGPU.cpp (+13-7)
• (modified) clang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl (+23)
• (modified) clang/test/SemaOpenCL/builtins-amdgcn-error-gfx9.cl (+1)

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_different_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_different_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 different types ('__private int' vs '__private long')}}
+  *out = __builtin_amdgcn_update_dpp(0.5f, i, 0, 0, 0, false); // expected-error{{arguments are of different types ('float' vs '__private int')}}
 }