Merge branch 'main' into fix/111854

Author: a-tarasyuk

clang/test/Format/error-unfound-files.cpp

@@ -0,0 +1,5 @@
+// RUN: rm -f a.c b.c
+
+// RUN: not clang-format a.c b.c 2>&1 | FileCheck %s
+// CHECK: a.c:
+// CHECK-NEXT: b.c:

clang/test/Modules/no-external-type-id.cppm

@@ -23,7 +23,7 @@ export module b;
 import a;
 export int b();
 
-// CHECK: <DECL_FUNCTION {{.*}} op8=4112
+// CHECK: <DECL_FUNCTION {{.*}} op8=4120
 // CHECK: <TYPE_FUNCTION_PROTO
 
 //--- a.v1.cppm

clang/tools/clang-format/ClangFormat.cpp

@@ -410,7 +410,7 @@ static bool format(StringRef FileName, bool ErrorOnIncompleteFormat = false) {
   const bool IsSTDIN = FileName == "-";
   if (!OutputXML && Inplace && IsSTDIN) {
     errs() << "error: cannot use -i when reading from stdin.\n";
-    return false;
+    return true;
   }
   // On Windows, overwriting a file with an open file mapping doesn't work,
   // so read the whole file into memory when formatting in-place.
@@ -419,7 +419,7 @@ static bool format(StringRef FileName, bool ErrorOnIncompleteFormat = false) {
           ? MemoryBuffer::getFileAsStream(FileName)
           : MemoryBuffer::getFileOrSTDIN(FileName, /*IsText=*/true);
   if (std::error_code EC = CodeOrErr.getError()) {
-    errs() << EC.message() << "\n";
+    errs() << FileName << ": " << EC.message() << "\n";
     return true;
   }
   std::unique_ptr<llvm::MemoryBuffer> Code = std::move(CodeOrErr.get());

compiler-rt/lib/builtins/CMakeLists.txt

@@ -104,6 +104,7 @@ set(GENERIC_SOURCES
   divti3.c
   extendsfdf2.c
   extendhfsf2.c
+  extendhfxf2.c
   ffsdi2.c
   ffssi2.c
   ffsti2.c

compiler-rt/lib/builtins/extendhfxf2.c

@@ -0,0 +1,18 @@
+//===-- lib/extendhfxf2.c - half -> long double conversion --------*- C -*-===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#define SRC_HALF
+#define DST_DOUBLE
+#include "fp_extend_impl.inc"
+
+// Use a forwarding definition and noinline to implement a poor man's alias,
+// as there isn't a good cross-platform way of defining one.
+// Long double are expected to be as precise as double.
+COMPILER_RT_ABI NOINLINE long double __extendhfxf2(src_t a) {
+  return (long double)__extendXfYf2__(a);
+}

compiler-rt/lib/builtins/macho_embedded/common.txt

@@ -60,6 +60,7 @@ divsf3
 divsi3
 extendsfdf2
 extendhfsf2
+extendhfxf2
 ffssi2
 fixdfsi
 fixsfsi

compiler-rt/test/builtins/Unit/extendhfxf2_test.c

@@ -0,0 +1,71 @@
+// RUN: %clang_builtins %s %librt -o %t && %run %t
+// REQUIRES: librt_has_extendhfxf2
+
+#include <limits.h>
+#include <math.h> // for isnan, isinf
+#include <stdio.h>
+
+#if __LDBL_MANT_DIG__ >= 64 && defined(COMPILER_RT_HAS_FLOAT16)
+
+long double __extendhfxf2(_Float16 f);
+
+int test_extendhfxf2(_Float16 a, long double expected) {
+  long double x = __extendhfxf2(a);
+  __uint16_t *b = (void *)&a;
+  int ret = !((isnan(x) && isnan(expected)) || x == expected);
+  if (ret) {
+    printf("error in test__extendhfxf2(%#.4x) = %.20Lf, "
+           "expected %.20Lf\n",
+           *b, x, expected);
+  }
+  return ret;
+}
+
+char assumption_1[sizeof(_Float16) * CHAR_BIT == 16] = {0};
+
+int main() {
+  // Small positive value
+  if (test_extendhfxf2(0.09997558593750000000f, 0.09997558593750000000L))
+    return 1;
+
+  // Small negative value
+  if (test_extendhfxf2(-0.09997558593750000000f, -0.09997558593750000000L))
+    return 1;
+
+  // Zero
+  if (test_extendhfxf2(0.0f, 0.0L))
+    return 1;
+
+  // Smallest positive non-zero value
+  if (test_extendhfxf2(0x1p-16f, 0x1p-16L))
+    return 1;
+
+  // Smallest negative non-zero value
+  if (test_extendhfxf2(-0x1p-16f, -0x1p-16L))
+    return 1;
+
+  // Positive infinity
+  if (test_extendhfxf2(__builtin_huge_valf16(), __builtin_huge_valf64x()))
+    return 1;
+
+  // Negative infinity
+  if (test_extendhfxf2(-__builtin_huge_valf16(),
+                       (long double)-__builtin_huge_valf64x()))
+    return 1;
+
+  // NaN
+  if (test_extendhfxf2(__builtin_nanf16(""),
+                       (long double)__builtin_nanf64x("")))
+    return 1;
+
+  return 0;
+}
+
+#else
+
+int main() {
+  printf("skipped\n");
+  return 0;
+}
+
+#endif

flang/include/flang/Optimizer/Transforms/CUFCommon.h

@@ -0,0 +1,25 @@
+//===-- CUFCommon.h -------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef FORTRAN_OPTIMIZER_TRANSFORMS_CUFCOMMON_H_
+#define FORTRAN_OPTIMIZER_TRANSFORMS_CUFCOMMON_H_
+
+#include "mlir/Dialect/GPU/IR/GPUDialect.h"
+#include "mlir/IR/BuiltinOps.h"
+
+static constexpr llvm::StringRef cudaDeviceModuleName = "cuda_device_mod";
+
+namespace cuf {
+
+/// Retrieve or create the CUDA Fortran GPU module in the given \p mod.
+mlir::gpu::GPUModuleOp getOrCreateGPUModule(mlir::ModuleOp mod,
+                                            mlir::SymbolTable &symTab);
+
+} // namespace cuf
+
+#endif // FORTRAN_OPTIMIZER_TRANSFORMS_CUFCOMMON_H_

flang/lib/Optimizer/Transforms/CMakeLists.txt

@@ -9,6 +9,7 @@ add_flang_library(FIRTransforms
   CompilerGeneratedNames.cpp
   ConstantArgumentGlobalisation.cpp
   ControlFlowConverter.cpp
+  CUFCommon.cpp
   CUFAddConstructor.cpp
   CUFDeviceGlobal.cpp
   CUFOpConversion.cpp

flang/lib/Optimizer/Transforms/CUFAddConstructor.cpp

@@ -11,6 +11,7 @@
 #include "flang/Optimizer/Dialect/FIRAttr.h"
 #include "flang/Optimizer/Dialect/FIRDialect.h"
 #include "flang/Optimizer/Dialect/FIROpsSupport.h"
+#include "flang/Optimizer/Transforms/CUFCommon.h"
 #include "flang/Runtime/entry-names.h"
 #include "mlir/Dialect/GPU/IR/GPUDialect.h"
 #include "mlir/Dialect/LLVMIR/LLVMDialect.h"
@@ -24,8 +25,6 @@ namespace fir {
 
 namespace {
 
-static constexpr llvm::StringRef cudaModName{"cuda_device_mod"};
-
 static constexpr llvm::StringRef cudaFortranCtorName{
     "__cudaFortranConstructor"};
 
@@ -60,15 +59,15 @@ struct CUFAddConstructor
     builder.create<mlir::LLVM::CallOp>(loc, funcTy, cufRegisterAllocatorRef);
 
     // Register kernels
-    auto gpuMod = symTab.lookup<mlir::gpu::GPUModuleOp>(cudaModName);
+    auto gpuMod = symTab.lookup<mlir::gpu::GPUModuleOp>(cudaDeviceModuleName);
     if (gpuMod) {
       auto llvmPtrTy = mlir::LLVM::LLVMPointerType::get(ctx);
       auto registeredMod = builder.create<cuf::RegisterModuleOp>(
           loc, llvmPtrTy, mlir::SymbolRefAttr::get(ctx, gpuMod.getName()));
       for (auto func : gpuMod.getOps<mlir::gpu::GPUFuncOp>()) {
         if (func.isKernel()) {
           auto kernelName = mlir::SymbolRefAttr::get(
-              builder.getStringAttr(cudaModName),
+              builder.getStringAttr(cudaDeviceModuleName),
               {mlir::SymbolRefAttr::get(builder.getContext(), func.getName())});
           builder.create<cuf::RegisterKernelOp>(loc, kernelName, registeredMod);
         }