[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

Author: rampitec

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;

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{{$}}

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