Merge branch 'main' into cuf_tma_load_llvmotr

Author: clementval

.ci/premerge_advisor_explain.py

@@ -129,7 +129,8 @@ def main(
         # If the job succeeds and there is not an existing comment, we
         # should not write one to reduce noise.
         comments = []
-    with open("comments", "w") as comment_file_handle:
+    comments_file_name = f"comments-{platform.system()}-{platform.machine()}"
+    with open(comments_file_name, "w") as comment_file_handle:
         json.dump(comments, comment_file_handle)
 
 

.github/workflows/premerge.yaml

@@ -124,9 +124,9 @@ jobs:
         if: ${{ always() && !startsWith(matrix.runs-on, 'depot-ubuntu-24.04-arm') }}
         continue-on-error: true
         with:
-          name: workflow-args
+          name: workflow-args-x86-linux
           path: |
-            comments
+            comments-Linux-x86_64
 
   premerge-checks-windows:
     name: Build and Test Windows
@@ -185,6 +185,14 @@ jobs:
           path: artifacts/
           retention-days: 5
           include-hidden-files: 'true'
+      - name: Upload Comment
+        uses: actions/upload-artifact@330a01c490aca151604b8cf639adc76d48f6c5d4 # v5.0.0
+        if: always()
+        continue-on-error: true
+        with:
+          name: workflow-args-windows
+          path: |
+            comments-Windows-x86_64
 
   premerge-check-macos:
     name: MacOS Premerge Checks

clang/lib/Tooling/DependencyScanning/DependencyScannerImpl.cpp

@@ -638,7 +638,7 @@ dependencies::initializeScanInstanceDependencyCollector(
 }
 
 bool DependencyScanningAction::runInvocation(
-    std::unique_ptr<CompilerInvocation> Invocation,
+    std::string Executable, std::unique_ptr<CompilerInvocation> Invocation,
     IntrusiveRefCntPtr<llvm::vfs::FileSystem> FS,
     std::shared_ptr<PCHContainerOperations> PCHContainerOps,
     DiagnosticConsumer *DiagConsumer) {
@@ -654,9 +654,12 @@ bool DependencyScanningAction::runInvocation(
   if (Scanned) {
     // Scanning runs once for the first -cc1 invocation in a chain of driver
     // jobs. For any dependent jobs, reuse the scanning result and just
-    // update the LastCC1Arguments to correspond to the new invocation.
+    // update the new invocation.
     // FIXME: to support multi-arch builds, each arch requires a separate scan
-    setLastCC1Arguments(std::move(OriginalInvocation));
+    if (MDC)
+      MDC->applyDiscoveredDependencies(OriginalInvocation);
+    Consumer.handleBuildCommand(
+        {Executable, OriginalInvocation.getCC1CommandLine()});
     return true;
   }
 
@@ -701,8 +704,12 @@ bool DependencyScanningAction::runInvocation(
   // ExecuteAction is responsible for calling finish.
   DiagConsumerFinished = true;
 
-  if (Result)
-    setLastCC1Arguments(std::move(OriginalInvocation));
+  if (Result) {
+    if (MDC)
+      MDC->applyDiscoveredDependencies(OriginalInvocation);
+    Consumer.handleBuildCommand(
+        {Executable, OriginalInvocation.getCC1CommandLine()});
+  }
 
   return Result;
 }

clang/lib/Tooling/DependencyScanning/DependencyScannerImpl.h

@@ -38,38 +38,23 @@ class DependencyScanningAction {
       std::optional<StringRef> ModuleName = std::nullopt)
       : Service(Service), WorkingDirectory(WorkingDirectory),
         Consumer(Consumer), Controller(Controller), DepFS(std::move(DepFS)) {}
-  bool runInvocation(std::unique_ptr<CompilerInvocation> Invocation,
+  bool runInvocation(std::string Executable,
+                     std::unique_ptr<CompilerInvocation> Invocation,
                      IntrusiveRefCntPtr<llvm::vfs::FileSystem> FS,
                      std::shared_ptr<PCHContainerOperations> PCHContainerOps,
                      DiagnosticConsumer *DiagConsumer);
 
   bool hasScanned() const { return Scanned; }
   bool hasDiagConsumerFinished() const { return DiagConsumerFinished; }
 
-  /// Take the cc1 arguments corresponding to the most recent invocation used
-  /// with this action. Any modifications implied by the discovered dependencies
-  /// will have already been applied.
-  std::vector<std::string> takeLastCC1Arguments() {
-    std::vector<std::string> Result;
-    std::swap(Result, LastCC1Arguments); // Reset LastCC1Arguments to empty.
-    return Result;
-  }
-
 private:
-  void setLastCC1Arguments(CompilerInvocation &&CI) {
-    if (MDC)
-      MDC->applyDiscoveredDependencies(CI);
-    LastCC1Arguments = CI.getCC1CommandLine();
-  }
-
   DependencyScanningService &Service;
   StringRef WorkingDirectory;
   DependencyConsumer &Consumer;
   DependencyActionController &Controller;
   IntrusiveRefCntPtr<DependencyScanningWorkerFilesystem> DepFS;
   std::optional<CompilerInstance> ScanInstanceStorage;
   std::shared_ptr<ModuleDepCollector> MDC;
-  std::vector<std::string> LastCC1Arguments;
   bool Scanned = false;
   bool DiagConsumerFinished = false;
 };

clang/lib/Tooling/DependencyScanning/DependencyScanningWorker.cpp

@@ -76,18 +76,14 @@ static bool createAndRunToolInvocation(
     DependencyScanningAction &Action,
     IntrusiveRefCntPtr<llvm::vfs::FileSystem> FS,
     std::shared_ptr<clang::PCHContainerOperations> &PCHContainerOps,
-    DiagnosticsEngine &Diags, DependencyConsumer &Consumer) {
+    DiagnosticsEngine &Diags) {
   auto Invocation = createCompilerInvocation(CommandLine, Diags);
   if (!Invocation)
     return false;
 
-  if (!Action.runInvocation(std::move(Invocation), std::move(FS),
-                            PCHContainerOps, Diags.getClient()))
-    return false;
-
-  std::vector<std::string> Args = Action.takeLastCC1Arguments();
-  Consumer.handleBuildCommand({CommandLine[0], std::move(Args)});
-  return true;
+  return Action.runInvocation(CommandLine[0], std::move(Invocation),
+                              std::move(FS), PCHContainerOps,
+                              Diags.getClient());
 }
 
 bool DependencyScanningWorker::scanDependencies(
@@ -112,9 +108,9 @@ bool DependencyScanningWorker::scanDependencies(
 
   bool Success = false;
   if (CommandLine[1] == "-cc1") {
-    Success = createAndRunToolInvocation(
-        CommandLine, Action, FS, PCHContainerOps,
-        *DiagEngineWithCmdAndOpts.DiagEngine, Consumer);
+    Success =
+        createAndRunToolInvocation(CommandLine, Action, FS, PCHContainerOps,
+                                   *DiagEngineWithCmdAndOpts.DiagEngine);
   } else {
     Success = forEachDriverJob(
         CommandLine, *DiagEngineWithCmdAndOpts.DiagEngine, FS,
@@ -128,7 +124,7 @@ bool DependencyScanningWorker::scanDependencies(
             return true;
           }
 
-          // Insert -cc1 comand line options into Argv
+          // Insert -cc1 command line options into Argv
           std::vector<std::string> Argv;
           Argv.push_back(Cmd.getExecutable());
           llvm::append_range(Argv, Cmd.getArguments());
@@ -139,7 +135,7 @@ bool DependencyScanningWorker::scanDependencies(
           // dependency scanning filesystem.
           return createAndRunToolInvocation(
               std::move(Argv), Action, FS, PCHContainerOps,
-              *DiagEngineWithCmdAndOpts.DiagEngine, Consumer);
+              *DiagEngineWithCmdAndOpts.DiagEngine);
         });
   }
 

flang/lib/Optimizer/Builder/CUDAIntrinsicCall.cpp

@@ -53,6 +53,8 @@ static const char __ldlu_r2x2[] = "__ldlu_r2x2_";
 static const char __ldlu_r4x4[] = "__ldlu_r4x4_";
 static const char __ldlu_r8x2[] = "__ldlu_r8x2_";
 
+static constexpr unsigned kTMAAlignment = 16;
+
 // CUDA specific intrinsic handlers.
 static constexpr IntrinsicHandler cudaHandlers[]{
     {"__ldca_i4x4",
@@ -1505,7 +1507,7 @@ static void genTMABulkLoad(fir::FirOpBuilder &builder, mlir::Location loc,
   mlir::Value size = mlir::arith::MulIOp::create(builder, loc, nelem, eleSize);
   auto llvmPtrTy = mlir::LLVM::LLVMPointerType::get(builder.getContext());
   barrier = builder.createConvert(loc, llvmPtrTy, barrier);
-  setAlignment(dst, 16);
+  setAlignment(dst, kTMAAlignment);
   dst = builder.createConvert(loc, llvmPtrTy, dst);
   src = builder.createConvert(loc, llvmPtrTy, src);
   mlir::NVVM::InlinePtxOp::create(
@@ -1609,6 +1611,7 @@ static void genTMABulkStore(fir::FirOpBuilder &builder, mlir::Location loc,
                             mlir::Value src, mlir::Value dst, mlir::Value count,
                             mlir::Value eleSize) {
   mlir::Value size = mlir::arith::MulIOp::create(builder, loc, eleSize, count);
+  setAlignment(src, kTMAAlignment);
   src = convertPtrToNVVMSpace(builder, loc, src,
                               mlir::NVVM::NVVMMemorySpace::Shared);
   dst = convertPtrToNVVMSpace(builder, loc, dst,

flang/test/Lower/CUDA/cuda-device-proc.cuf

@@ -675,6 +675,7 @@ attributes(global) subroutine test_tma_bulk_store_c4(c, n)
 end subroutine
 
 ! CHECK-LABEL: func.func @_QPtest_tma_bulk_store_c4
+! CHECK: cuf.shared_memory !fir.array<1024xcomplex<f32>> align 16 {bindc_name = "tmpa", uniq_name = "_QFtest_tma_bulk_store_c4Etmpa"} -> !fir.ref<!fir.array<1024xcomplex<f32>>>
 ! CHECK: nvvm.cp.async.bulk.global.shared.cta %{{.*}}, %{{.*}}, %{{.*}} : <1>, <3>
 ! CHECK: nvvm.inline_ptx "cp.async.bulk.commit_group;"
 ! CHECK: nvvm.cp.async.bulk.wait_group 0
@@ -688,6 +689,7 @@ attributes(global) subroutine test_tma_bulk_store_c8(c, n)
 end subroutine
 
 ! CHECK-LABEL: func.func @_QPtest_tma_bulk_store_c8
+! CHECK: cuf.shared_memory !fir.array<1024xcomplex<f64>> align 16 {bindc_name = "tmpa", uniq_name = "_QFtest_tma_bulk_store_c8Etmpa"} -> !fir.ref<!fir.array<1024xcomplex<f64>>>
 ! CHECK: nvvm.cp.async.bulk.global.shared.cta %{{.*}}, %{{.*}}, %{{.*}} : <1>, <3>
 ! CHECK: nvvm.inline_ptx "cp.async.bulk.commit_group;"
 ! CHECK: nvvm.cp.async.bulk.wait_group 0
@@ -701,6 +703,7 @@ attributes(global) subroutine test_tma_bulk_store_i4(c, n)
 end subroutine
 
 ! CHECK-LABEL: func.func @_QPtest_tma_bulk_store_i4
+! CHECK: cuf.shared_memory !fir.array<1024xi32> align 16 {bindc_name = "tmpa", uniq_name = "_QFtest_tma_bulk_store_i4Etmpa"} -> !fir.ref<!fir.array<1024xi32>>
 ! CHECK: nvvm.cp.async.bulk.global.shared.cta %{{.*}}, %{{.*}}, %{{.*}} : <1>, <3>
 ! CHECK: nvvm.inline_ptx "cp.async.bulk.commit_group;"
 ! CHECK: nvvm.cp.async.bulk.wait_group 0
@@ -714,6 +717,7 @@ attributes(global) subroutine test_tma_bulk_store_i8(c, n)
 end subroutine
 
 ! CHECK-LABEL: func.func @_QPtest_tma_bulk_store_i8
+! CHECK: cuf.shared_memory !fir.array<1024xi64> align 16 {bindc_name = "tmpa", uniq_name = "_QFtest_tma_bulk_store_i8Etmpa"} -> !fir.ref<!fir.array<1024xi64>>
 ! CHECK: nvvm.cp.async.bulk.global.shared.cta %{{.*}}, %{{.*}}, %{{.*}} : <1>, <3>
 ! CHECK: nvvm.inline_ptx "cp.async.bulk.commit_group;"
 ! CHECK: nvvm.cp.async.bulk.wait_group 0
@@ -728,6 +732,7 @@ attributes(global) subroutine test_tma_bulk_store_r2(c, n)
 end subroutine
 
 ! CHECK-LABEL: func.func @_QPtest_tma_bulk_store_r2
+! CHECK: cuf.shared_memory !fir.array<1024xf16> align 16 {bindc_name = "tmpa", uniq_name = "_QFtest_tma_bulk_store_r2Etmpa"} -> !fir.ref<!fir.array<1024xf16>>
 ! CHECK: nvvm.cp.async.bulk.global.shared.cta %{{.*}}, %{{.*}}, %{{.*}} : <1>, <3>
 ! CHECK: nvvm.inline_ptx "cp.async.bulk.commit_group;"
 ! CHECK: nvvm.cp.async.bulk.wait_group 0
@@ -741,6 +746,7 @@ attributes(global) subroutine test_tma_bulk_store_r4(c, n)
 end subroutine
 
 ! CHECK-LABEL: func.func @_QPtest_tma_bulk_store_r4
+! CHECK: cuf.shared_memory !fir.array<1024xf32> align 16 {bindc_name = "tmpa", uniq_name = "_QFtest_tma_bulk_store_r4Etmpa"} -> !fir.ref<!fir.array<1024xf32>>
 ! CHECK: nvvm.cp.async.bulk.global.shared.cta %{{.*}}, %{{.*}}, %{{.*}} : <1>, <3>
 ! CHECK: nvvm.inline_ptx "cp.async.bulk.commit_group;"
 ! CHECK: nvvm.cp.async.bulk.wait_group 0
@@ -754,6 +760,7 @@ attributes(global) subroutine test_tma_bulk_store_r8(c, n)
 end subroutine
 
 ! CHECK-LABEL: func.func @_QPtest_tma_bulk_store_r8
+! CHECK: cuf.shared_memory !fir.array<1024xf64> align 16 {bindc_name = "tmpa", uniq_name = "_QFtest_tma_bulk_store_r8Etmpa"} -> !fir.ref<!fir.array<1024xf64>>
 ! CHECK: nvvm.cp.async.bulk.global.shared.cta %{{.*}}, %{{.*}}, %{{.*}} : <1>, <3>
 ! CHECK: nvvm.inline_ptx "cp.async.bulk.commit_group;"
 ! CHECK: nvvm.cp.async.bulk.wait_group 0

flang/test/Lower/OpenMP/Todo/reduction-task.f90

@@ -8,5 +8,5 @@ subroutine reduction_task()
 
   !$omp parallel reduction(task, +:i)
   i = i + 1
-  !$omp end parallel 
+  !$omp end parallel
 end subroutine reduction_task

flang/test/Lower/OpenMP/allocatable-array-bounds.f90

@@ -47,7 +47,7 @@
 subroutine read_write_section()
     integer, allocatable :: sp_read(:)
     integer, allocatable :: sp_write(:)
-    allocate(sp_read(10)) 
+    allocate(sp_read(10))
     allocate(sp_write(10))
     sp_write = (/0,0,0,0,0,0,0,0,0,0/)
     sp_read = (/1,2,3,4,5,6,7,8,9,10/)

flang/test/Lower/OpenMP/allocatable-map.f90

@@ -6,7 +6,7 @@
 !HLFIRDIALECT: %[[POINTER_MAP:.*]] = omp.map.info var_ptr(%[[POINTER]]#1 : !fir.ref<!fir.box<!fir.ptr<i32>>>, !fir.box<!fir.ptr<i32>>) map_clauses(always, to) capture(ByRef) members(%[[POINTER_MAP_MEMBER]] : [0] : !fir.llvm_ptr<!fir.ref<i32>>) -> !fir.ref<!fir.box<!fir.ptr<i32>>> {name = "point"}
 !HLFIRDIALECT: omp.target map_entries(%[[POINTER_MAP]] -> {{.*}}, %[[POINTER_MAP_MEMBER]] -> {{.*}} : !fir.ref<!fir.box<!fir.ptr<i32>>>, !fir.llvm_ptr<!fir.ref<i32>>) {
 subroutine pointer_routine()
-    integer, pointer :: point 
+    integer, pointer :: point
 !$omp target map(tofrom:point)
     point = 1
 !$omp end target