Merge upstream LLVM into amd-gfx12

Author: SC llvm team

clang-tools-extra/clang-tidy/llvm/UseNewMLIROpBuilderCheck.cpp

@@ -111,17 +111,23 @@ EditGenerator rewrite(RangeSelector Call, RangeSelector Builder,
 }
 
 RewriteRuleWith<std::string> useNewMlirOpBuilderCheckRule() {
-  return makeRule(
+  Stencil message = cat("use 'OpType::create(builder, ...)' instead of "
+                        "'builder.create<OpType>(...)'");
+  // Match a create call on an OpBuilder.
+  ast_matchers::internal::Matcher<Stmt> base =
       cxxMemberCallExpr(
           on(expr(hasType(
                       cxxRecordDecl(isSameOrDerivedFrom("::mlir::OpBuilder"))))
                  .bind("builder")),
           callee(cxxMethodDecl(hasTemplateArgument(0, templateArgument()))),
           callee(cxxMethodDecl(hasName("create"))))
-          .bind("call"),
-      rewrite(node("call"), node("builder"), callArgs("call")),
-      cat("use 'OpType::create(builder, ...)' instead of "
-          "'builder.create<OpType>(...)'"));
+          .bind("call");
+  return applyFirst(
+      //  Attempt rewrite given an lvalue builder, else just warn.
+      {makeRule(cxxMemberCallExpr(unless(on(cxxTemporaryObjectExpr())), base),
+                rewrite(node("call"), node("builder"), callArgs("call")),
+                message),
+       makeRule(base, noopEdit(node("call")), message)});
 }
 } // namespace
 

clang-tools-extra/test/clang-tidy/checkers/llvm/use-new-mlir-op-builder.cpp

@@ -69,4 +69,8 @@ void f() {
   // CHECK-MESSAGES: :[[@LINE+2]]:3: warning: use 'OpType::create(builder, ...)' instead of 'builder.create<OpType>(...)' [llvm-use-new-mlir-op-builder]
   // CHECK-FIXES: mlir::ModuleOp::create(ib)
   ib.create<mlir::ModuleOp>(   );
+
+  // CHECK-MESSAGES: :[[@LINE+2]]:3: warning: use 'OpType::create(builder, ...)' instead of 'builder.create<OpType>(...)' [llvm-use-new-mlir-op-builder]
+  // CHECK-FIXES: mlir::OpBuilder().create<mlir::ModuleOp>(builder.getUnknownLoc());
+  mlir::OpBuilder().create<mlir::ModuleOp>(builder.getUnknownLoc());
 }

clang/docs/LanguageExtensions.rst

@@ -68,6 +68,9 @@ It can be used like this:
   ``__has_builtin`` should not be used to detect support for a builtin macro;
   use ``#ifdef`` instead.
 
+  When compiling with target offloading, ``__has_builtin`` only considers the
+  currently active target.
+
 ``__has_constexpr_builtin``
 ---------------------------
 

clang/docs/ReleaseNotes.rst

@@ -37,6 +37,8 @@ Potentially Breaking Changes
 C/C++ Language Potentially Breaking Changes
 -------------------------------------------
 
+- The ``__has_builtin`` function now only considers the currently active target when being used with target offloading.
+
 C++ Specific Potentially Breaking Changes
 -----------------------------------------
 - For C++20 modules, the Reduced BMI mode will be the default option. This may introduce

clang/lib/Driver/Driver.cpp

@@ -910,7 +910,7 @@ getSystemOffloadArchs(Compilation &C, Action::OffloadKind Kind) {
 
   SmallVector<std::string> GPUArchs;
   if (llvm::ErrorOr<std::string> Executable =
-          llvm::sys::findProgramByName(Program)) {
+          llvm::sys::findProgramByName(Program, {C.getDriver().Dir})) {
     llvm::SmallVector<StringRef> Args{*Executable};
     if (Kind == Action::OFK_HIP)
       Args.push_back("--only=amdgpu");
@@ -4919,13 +4919,14 @@ Action *Driver::BuildOffloadingActions(Compilation &C,
   }
 
   // HIP code in device-only non-RDC mode will bundle the output if it invoked
-  // the linker.
+  // the linker or if the user explicitly requested it.
   bool ShouldBundleHIP =
-      HIPNoRDC && offloadDeviceOnly() &&
       Args.hasFlag(options::OPT_gpu_bundle_output,
-                   options::OPT_no_gpu_bundle_output, true) &&
-      !llvm::any_of(OffloadActions,
-                    [](Action *A) { return A->getType() != types::TY_Image; });
+                   options::OPT_no_gpu_bundle_output, false) ||
+      (HIPNoRDC && offloadDeviceOnly() &&
+       llvm::none_of(OffloadActions, [](Action *A) {
+         return A->getType() != types::TY_Image;
+       }));
 
   // All kinds exit now in device-only mode except for non-RDC mode HIP.
   if (offloadDeviceOnly() && !ShouldBundleHIP)

clang/lib/Driver/ToolChains/HIPSPV.cpp

@@ -69,8 +69,17 @@ void HIPSPV::Linker::constructLinkAndEmitSpirvCommand(
 
   // Link LLVM bitcode.
   ArgStringList LinkArgs{};
+
   for (auto Input : Inputs)
     LinkArgs.push_back(Input.getFilename());
+
+  // Add static device libraries using the common helper function.
+  // This handles unbundling archives (.a) containing bitcode bundles.
+  StringRef Arch = getToolChain().getTriple().getArchName();
+  StringRef Target =
+      "generic"; // SPIR-V is generic, no specific target ID like -mcpu
+  tools::AddStaticDeviceLibsLinking(C, *this, JA, Inputs, Args, LinkArgs, Arch,
+                                    Target, /*IsBitCodeSDL=*/true);
   LinkArgs.append({"-o", TempFile});
   const char *LlvmLink =
       Args.MakeArgString(getToolChain().GetProgramPath("llvm-link"));

clang/lib/Lex/PPMacroExpansion.cpp

@@ -1760,7 +1760,8 @@ void Preprocessor::ExpandBuiltinMacro(Token &Tok) {
               Tok, *this, diag::err_feature_check_malformed);
           if (!II)
             return false;
-          else if (II->getBuiltinID() != 0) {
+          unsigned BuiltinID = II->getBuiltinID();
+          if (BuiltinID != 0) {
             switch (II->getBuiltinID()) {
             case Builtin::BI__builtin_cpu_is:
               return getTargetInfo().supportsCpuIs();
@@ -1774,8 +1775,11 @@ void Preprocessor::ExpandBuiltinMacro(Token &Tok) {
               // usual allocation and deallocation functions. Required by libc++
               return 201802;
             default:
+              // __has_builtin should return false for aux builtins.
+              if (getBuiltinInfo().isAuxBuiltinID(BuiltinID))
+                return false;
               return Builtin::evaluateRequiredTargetFeatures(
-                  getBuiltinInfo().getRequiredFeatures(II->getBuiltinID()),
+                  getBuiltinInfo().getRequiredFeatures(BuiltinID),
                   getTargetInfo().getTargetOpts().FeatureMap);
             }
             return true;

clang/test/CodeGenCUDA/bf16.cu

@@ -35,8 +35,8 @@ __device__ __bf16 external_func( __bf16 in);
 // CHECK:        .param .align 2 .b8 _Z9test_callDF16b_param_0[2]
 __device__ __bf16 test_call( __bf16 in) {
 // CHECK:        ld.param.b16    %[[R:rs[0-9]+]], [_Z9test_callDF16b_param_0];
-// CHECK:        st.param.b16    [param0], %[[R]];
 // CHECK:        .param .align 2 .b8 retval0[2];
+// CHECK:        st.param.b16    [param0], %[[R]];
 // CHECK:        call.uni (retval0), _Z13external_funcDF16b, (param0);
 // CHECK:        ld.param.b16    %[[RET:rs[0-9]+]], [retval0];
   return external_func(in);

clang/test/Driver/hip-binding.hip

@@ -61,7 +61,7 @@
 // MULTI-D-ONLY-NEXT: # "amdgcn-amd-amdhsa" - "clang", inputs: ["[[INPUT]]"], output: "[[GFX90a:.+]]"
 // MULTI-D-ONLY-NEXT: # "amdgcn-amd-amdhsa" - "AMDGCN::Linker", inputs: ["[[GFX90a]]"], output: "[[GFX90a_OUT:.+]]"
 //
-// RUN: not %clang -### --target=x86_64-linux-gnu --offload-new-driver -ccc-print-bindings -nogpulib -nogpuinc \
+// RUN: not %clang -### --target=x86_64-linux-gnu --offload-new-driver -ccc-print-bindings -nogpulib -nogpuinc -emit-llvm \
 // RUN:        --no-gpu-bundle-output --offload-arch=gfx90a --offload-arch=gfx908 --offload-device-only -c -o %t %s 2>&1 \
 // RUN: | FileCheck -check-prefix=MULTI-D-ONLY-NO-BUNDLE-O %s
 // MULTI-D-ONLY-NO-BUNDLE-O: error: cannot specify -o when generating multiple output files
@@ -75,6 +75,13 @@
 // MULTI-D-ONLY-O-NEXT: "amdgcn-amd-amdhsa" - "AMDGCN::Linker", inputs: ["[[GFX90A_OBJ]]"], output: "[[GFX90A:.+]]"
 // MULTI-D-ONLY-O-NEXT: "amdgcn-amd-amdhsa" - "AMDGCN::Linker", inputs: ["[[GFX908]]", "[[GFX90A]]"], output: "a.out"
 
+// RUN: %clang -### --target=x86_64-linux-gnu --offload-new-driver -ccc-print-bindings -nogpulib -nogpuinc -emit-llvm \
+// RUN:        --gpu-bundle-output --offload-arch=gfx90a --offload-arch=gfx908 --offload-device-only -c -o a.out %s 2>&1 \
+// RUN: | FileCheck -check-prefix=MULTI-D-ONLY-BC %s
+//      MULTI-D-ONLY-BC: "amdgcn-amd-amdhsa" - "clang", inputs: ["[[INPUT:.+]]"], output: "[[GFX908_BC:.+]]"
+// MULTI-D-ONLY-BC-NEXT: "amdgcn-amd-amdhsa" - "clang", inputs: ["[[INPUT]]"], output: "[[GFX90A_BC:.+]]"
+// MULTI-D-ONLY-BC-NEXT: "amdgcn-amd-amdhsa" - "AMDGCN::Linker", inputs: ["[[GFX908_BC]]", "[[GFX90A_BC]]"], output: "a.out"
+
 //
 // Check to ensure that we can use '-fsyntax-only' for HIP output with the new
 // driver.

clang/test/Driver/hipspv-link-static-library.hip

@@ -0,0 +1,28 @@
+// Test HIPSPV static device library linking
+// REQUIRES: system-linux
+// UNSUPPORTED: system-windows
+
+// Create a dummy archive to test SDL linking
+// RUN: rm -rf %t && mkdir %t
+// RUN: touch %t/dummy.bc  
+// RUN: llvm-ar cr %t/libSDL.a %t/dummy.bc
+
+// Test that -l options are passed to llvm-link for --offload=spirv64
+// RUN: %clang -### --target=x86_64-linux-gnu --offload=spirv64 \
+// RUN:   --hip-path=%S/Inputs/hipspv -nohipwrapperinc %s \
+// RUN:   -L%t -lSDL \
+// RUN: 2>&1 | FileCheck -check-prefix=SDL-LINK %s
+
+// Test that .a files are properly unbundled and passed to llvm-link  
+// RUN: %clang -### --target=x86_64-linux-gnu --offload=spirv64 \
+// RUN:   --hip-path=%S/Inputs/hipspv -nohipwrapperinc %s \
+// RUN:   %t/libSDL.a \
+// RUN: 2>&1 | FileCheck -check-prefix=SDL-ARCHIVE %s
+
+// Verify that the input files are added before the SDL files in llvm-link command
+// This tests the ordering fix to match HIPAMD behavior
+// SDL-LINK: "{{.*}}clang-offload-bundler" "-unbundle" "-type=a" "-input={{.*}}libSDL.a" "-targets=hip-spirv64-unknown-unknown-unknown-generic" "-output=[[SDL_A:.*\.a]]" "-allow-missing-bundles"
+// SDL-LINK: "{{.*}}llvm-link" "{{.*}}.bc" "[[SDL_A]]" "-o"
+
+// SDL-ARCHIVE: "{{.*}}clang-offload-bundler" "-unbundle" "-type=a" "-input={{.*}}libSDL.a" "-targets=hip-spirv64-unknown-unknown-unknown-generic" "-output=[[SDL_A:.*\.a]]" "-allow-missing-bundles"  
+// SDL-ARCHIVE: "{{.*}}llvm-link" "{{.*}}.bc" "[[SDL_A]]" "-o"