Fix formatting

Created using spr 1.3.5

Author: alexey-bataev

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

flang/include/flang/Optimizer/Transforms/CufOpConversion.h renamed to flang/include/flang/Optimizer/Transforms/CUFOpConversion.h

@@ -1,4 +1,4 @@
-//===------- Optimizer/Transforms/CufOpConversion.h -------------*- C++ -*-===//
+//===------- Optimizer/Transforms/CUFOpConversion.h -------------*- C++ -*-===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -19,7 +19,7 @@ class LLVMTypeConverter;
 namespace mlir {
 class DataLayout;
 class SymbolTable;
-}
+} // namespace mlir
 
 namespace cuf {
 

flang/include/flang/Optimizer/Transforms/Passes.h

@@ -40,7 +40,7 @@ namespace fir {
 #define GEN_PASS_DECL_CHARACTERCONVERSION
 #define GEN_PASS_DECL_CFGCONVERSION
 #define GEN_PASS_DECL_CUFADDCONSTRUCTOR
-#define GEN_PASS_DECL_CUFIMPLICITDEVICEGLOBAL
+#define GEN_PASS_DECL_CUFDEVICEGLOBAL
 #define GEN_PASS_DECL_CUFOPCONVERSION
 #define GEN_PASS_DECL_EXTERNALNAMECONVERSION
 #define GEN_PASS_DECL_MEMREFDATAFLOWOPT

flang/include/flang/Optimizer/Transforms/Passes.td

@@ -421,15 +421,15 @@ def AssumedRankOpConversion : Pass<"fir-assumed-rank-op", "mlir::ModuleOp"> {
   ];
 }
 
-def CufOpConversion : Pass<"cuf-convert", "mlir::ModuleOp"> {
+def CUFOpConversion : Pass<"cuf-convert", "mlir::ModuleOp"> {
   let summary = "Convert some CUF operations to runtime calls";
   let dependentDialects = [
     "fir::FIROpsDialect"
   ];
 }
 
-def CufImplicitDeviceGlobal :
-    Pass<"cuf-implicit-device-global", "mlir::ModuleOp"> {
+def CUFDeviceGlobal :
+    Pass<"cuf-device-global", "mlir::ModuleOp"> {
   let summary = "Flag globals used in device function with data attribute";
   let dependentDialects = [
     "cuf::CUFDialect"

flang/lib/Optimizer/Transforms/CMakeLists.txt

@@ -10,8 +10,8 @@ add_flang_library(FIRTransforms
   ConstantArgumentGlobalisation.cpp
   ControlFlowConverter.cpp
   CUFAddConstructor.cpp
-  CufImplicitDeviceGlobal.cpp
-  CufOpConversion.cpp
+  CUFDeviceGlobal.cpp
+  CUFOpConversion.cpp
   ArrayValueCopy.cpp
   ExternalNameConversion.cpp
   MemoryUtils.cpp

flang/lib/Optimizer/Transforms/CufImplicitDeviceGlobal.cpp renamed to flang/lib/Optimizer/Transforms/CUFDeviceGlobal.cpp

@@ -1,4 +1,4 @@
-//===-- CufOpConversion.cpp -----------------------------------------------===//
+//===-- CUFOpConversion.cpp -----------------------------------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -18,7 +18,7 @@
 #include "mlir/Transforms/DialectConversion.h"
 
 namespace fir {
-#define GEN_PASS_DEF_CUFIMPLICITDEVICEGLOBAL
+#define GEN_PASS_DEF_CUFDEVICEGLOBAL
 #include "flang/Optimizer/Transforms/Passes.h.inc"
 } // namespace fir
 
@@ -45,8 +45,7 @@ static void prepareImplicitDeviceGlobals(mlir::func::FuncOp funcOp,
   }
 }
 
-class CufImplicitDeviceGlobal
-    : public fir::impl::CufImplicitDeviceGlobalBase<CufImplicitDeviceGlobal> {
+class CUFDeviceGlobal : public fir::impl::CUFDeviceGlobalBase<CUFDeviceGlobal> {
 public:
   void runOnOperation() override {
     mlir::Operation *op = getOperation();

flang/lib/Optimizer/Transforms/CufOpConversion.cpp renamed to flang/lib/Optimizer/Transforms/CUFOpConversion.cpp

@@ -1,12 +1,12 @@
-//===-- CufOpConversion.cpp -----------------------------------------------===//
+//===-- CUFDeviceGlobal.cpp -----------------------------------------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
 
-#include "flang/Optimizer/Transforms/CufOpConversion.h"
+#include "flang/Optimizer/Transforms/CUFOpConversion.h"
 #include "flang/Common/Fortran.h"
 #include "flang/Optimizer/Builder/Runtime/RTBuilder.h"
 #include "flang/Optimizer/CodeGen/TypeConverter.h"
@@ -619,7 +619,7 @@ struct CufDataTransferOpConversion
   const mlir::SymbolTable &symtab;
 };
 
-class CufOpConversion : public fir::impl::CufOpConversionBase<CufOpConversion> {
+class CUFOpConversion : public fir::impl::CUFOpConversionBase<CUFOpConversion> {
 public:
   void runOnOperation() override {
     auto *ctx = &getContext();

flang/runtime/CUDA/registration.cpp

@@ -14,13 +14,16 @@ namespace Fortran::runtime::cuda {
 
 extern "C" {
 
-extern void **__cudaRegisterFatBinary(void *data);
+extern void **__cudaRegisterFatBinary(void *);
+extern void __cudaRegisterFatBinaryEnd(void *);
 extern void __cudaRegisterFunction(void **fatCubinHandle, const char *hostFun,
     char *deviceFun, const char *deviceName, int thread_limit, uint3 *tid,
     uint3 *bid, dim3 *bDim, dim3 *gDim, int *wSize);
 
 void *RTDECL(CUFRegisterModule)(void *data) {
-  return __cudaRegisterFatBinary(data);
+  void **fatHandle{__cudaRegisterFatBinary(data)};
+  __cudaRegisterFatBinaryEnd(fatHandle);
+  return fatHandle;
 }
 
 void RTDEF(CUFRegisterFunction)(void **module, const char *fct) {