Merge branch 'main' into lldb-mcp-multi

Author: ashgti

.ci/all_requirements.txt

@@ -0,0 +1,80 @@
+# 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
+"""Caches .lit_test_times.txt files between premerge invocations.
+
+.lit_test_times.txt files are used by lit to order tests to best take advantage
+of parallelism. Having them around and up to date can result in a ~15%
+improvement in test times. This script downloading cached test time files and
+uploading new versions to the GCS buckets used for caching.
+"""
+
+import sys
+import os
+import logging
+import multiprocessing.pool
+import pathlib
+import glob
+
+from google.cloud import storage
+
+GCS_PARALLELISM = 100
+
+
+def _maybe_upload_timing_file(bucket, timing_file_path):
+    if os.path.exists(timing_file_path):
+        timing_file_blob = bucket.blob("lit_timing/" + timing_file_path)
+        timing_file_blob.upload_from_filename(timing_file_path)
+
+
+def upload_timing_files(storage_client, bucket_name: str):
+    bucket = storage_client.bucket(bucket_name)
+    with multiprocessing.pool.ThreadPool(GCS_PARALLELISM) as thread_pool:
+        futures = []
+        for timing_file_path in glob.glob("**/.lit_test_times.txt", recursive=True):
+            futures.append(
+                thread_pool.apply_async(
+                    _maybe_upload_timing_file, (bucket, timing_file_path)
+                )
+            )
+        for future in futures:
+            future.get()
+    print("Done uploading")
+
+
+def _maybe_download_timing_file(blob):
+    file_name = blob.name.removeprefix("lit_timing/")
+    pathlib.Path(os.path.dirname(file_name)).mkdir(parents=True, exist_ok=True)
+    blob.download_to_filename(file_name)
+
+
+def download_timing_files(storage_client, bucket_name: str):
+    bucket = storage_client.bucket(bucket_name)
+    blobs = bucket.list_blobs(prefix="lit_timing")
+    with multiprocessing.pool.ThreadPool(GCS_PARALLELISM) as thread_pool:
+        futures = []
+        for timing_file_blob in blobs:
+            futures.append(
+                thread_pool.apply_async(
+                    _maybe_download_timing_file, (timing_file_blob,)
+                )
+            )
+        for future in futures:
+            future.get()
+    print("Done downloading")
+
+
+if __name__ == "__main__":
+    if len(sys.argv) != 2:
+        logging.fatal("Expected usage is cache_lit_timing_files.py <upload/download>")
+        sys.exit(1)
+    action = sys.argv[1]
+    storage_client = storage.Client()
+    bucket_name = os.environ["CACHE_GCS_BUCKET"]
+    if action == "download":
+        download_timing_files(storage_client, bucket_name)
+    elif action == "upload":
+        upload_timing_files(storage_client, bucket_name)
+    else:
+        logging.fatal("Expected usage is cache_lit_timing_files.py <upload/download>")
+        sys.exit(1)

.ci/cache_lit_timing_files.py

@@ -1 +1,2 @@
 junitparser==3.2.0
+google-cloud-storage==3.3.0

.ci/requirements.txt

@@ -2,7 +2,7 @@ FROM docker.io/library/ubuntu:24.04 as base
 ENV LLVM_SYSROOT=/opt/llvm
 
 FROM base as stage1-toolchain
-ENV LLVM_VERSION=20.1.8
+ENV LLVM_VERSION=21.1.0
 
 RUN apt-get update && \
     apt-get install -y \

.github/workflows/containers/github-action-ci/Dockerfile

@@ -87,10 +87,6 @@ class NVPTXTargetCodeGenInfo : public TargetCodeGenInfo {
   static void addNVVMMetadata(llvm::GlobalValue *GV, StringRef Name,
                               int Operand);
 
-  static void
-  addGridConstantNVVMMetadata(llvm::GlobalValue *GV,
-                              const SmallVectorImpl<int> &GridConstantArgs);
-
 private:
   static void emitBuiltinSurfTexDeviceCopy(CodeGenFunction &CGF, LValue Dst,
                                            LValue Src) {
@@ -265,27 +261,24 @@ void NVPTXTargetCodeGenInfo::setTargetAttributes(
     // By default, all functions are device functions
     if (FD->hasAttr<DeviceKernelAttr>() || FD->hasAttr<CUDAGlobalAttr>()) {
       // OpenCL/CUDA kernel functions get kernel metadata
-      // Create !{<func-ref>, metadata !"kernel", i32 1} node
       // And kernel functions are not subject to inlining
       F->addFnAttr(llvm::Attribute::NoInline);
       if (FD->hasAttr<CUDAGlobalAttr>()) {
-        SmallVector<int, 10> GCI;
+        F->setCallingConv(llvm::CallingConv::PTX_Kernel);
+
         for (auto IV : llvm::enumerate(FD->parameters()))
           if (IV.value()->hasAttr<CUDAGridConstantAttr>())
-            // For some reason arg indices are 1-based in NVVM
-            GCI.push_back(IV.index() + 1);
-        // Create !{<func-ref>, metadata !"kernel", i32 1} node
-        F->setCallingConv(llvm::CallingConv::PTX_Kernel);
-        addGridConstantNVVMMetadata(F, GCI);
+            F->addParamAttr(
+                IV.index(),
+                llvm::Attribute::get(F->getContext(), "nvvm.grid_constant"));
       }
       if (CUDALaunchBoundsAttr *Attr = FD->getAttr<CUDALaunchBoundsAttr>())
         M.handleCUDALaunchBoundsAttr(F, Attr);
     }
   }
   // Attach kernel metadata directly if compiling for NVPTX.
-  if (FD->hasAttr<DeviceKernelAttr>()) {
+  if (FD->hasAttr<DeviceKernelAttr>())
     F->setCallingConv(llvm::CallingConv::PTX_Kernel);
-  }
 }
 
 void NVPTXTargetCodeGenInfo::addNVVMMetadata(llvm::GlobalValue *GV,
@@ -305,29 +298,6 @@ void NVPTXTargetCodeGenInfo::addNVVMMetadata(llvm::GlobalValue *GV,
   MD->addOperand(llvm::MDNode::get(Ctx, MDVals));
 }
 
-void NVPTXTargetCodeGenInfo::addGridConstantNVVMMetadata(
-    llvm::GlobalValue *GV, const SmallVectorImpl<int> &GridConstantArgs) {
-
-  llvm::Module *M = GV->getParent();
-  llvm::LLVMContext &Ctx = M->getContext();
-
-  // Get "nvvm.annotations" metadata node
-  llvm::NamedMDNode *MD = M->getOrInsertNamedMetadata("nvvm.annotations");
-
-  SmallVector<llvm::Metadata *, 5> MDVals = {llvm::ConstantAsMetadata::get(GV)};
-  if (!GridConstantArgs.empty()) {
-    SmallVector<llvm::Metadata *, 10> GCM;
-    for (int I : GridConstantArgs)
-      GCM.push_back(llvm::ConstantAsMetadata::get(
-          llvm::ConstantInt::get(llvm::Type::getInt32Ty(Ctx), I)));
-    MDVals.append({llvm::MDString::get(Ctx, "grid_constant"),
-                   llvm::MDNode::get(Ctx, GCM)});
-  }
-
-  // Append metadata to nvvm.annotations
-  MD->addOperand(llvm::MDNode::get(Ctx, MDVals));
-}
-
 bool NVPTXTargetCodeGenInfo::shouldEmitStaticExternCAliases() const {
   return false;
 }

clang/lib/CodeGen/Targets/NVPTX.cpp

@@ -19,13 +19,9 @@ void foo() {
   tkernel_const<S><<<1,1>>>({});
   tkernel<const S><<<1,1>>>(1, {});
 }
-//.
-//.
-// CHECK: [[META0:![0-9]+]] = !{ptr @_Z6kernel1Sii, !"grid_constant", [[META1:![0-9]+]]}
-// CHECK: [[META1]] = !{i32 1, i32 3}
-// CHECK: [[META2:![0-9]+]] = !{ptr @_Z13tkernel_constIK1SEvT_, !"grid_constant", [[META3:![0-9]+]]}
-// CHECK: [[META3]] = !{i32 1}
-// CHECK: [[META4:![0-9]+]] = !{ptr @_Z13tkernel_constI1SEvT_, !"grid_constant", [[META3]]}
-// CHECK: [[META5:![0-9]+]] = !{ptr @_Z7tkernelIK1SEviT_, !"grid_constant", [[META6:![0-9]+]]}
-// CHECK: [[META6]] = !{i32 2}
-//.
+
+// CHECK: define dso_local ptx_kernel void @_Z6kernel1Sii(ptr noundef byval(%struct.S) align 1 "nvvm.grid_constant" %gc_arg1, i32 noundef %arg2, i32 noundef "nvvm.grid_constant" %gc_arg3)
+// CHECK: define ptx_kernel void @_Z13tkernel_constIK1SEvT_(ptr noundef byval(%struct.S) align 1 "nvvm.grid_constant" %arg)
+// CHECK: define ptx_kernel void @_Z13tkernel_constI1SEvT_(ptr noundef byval(%struct.S) align 1 "nvvm.grid_constant" %arg)
+// CHECK: define ptx_kernel void @_Z7tkernelIK1SEviT_(i32 noundef %dummy, ptr noundef byval(%struct.S) align 1 "nvvm.grid_constant" %arg)
+

clang/test/CodeGenCUDA/grid-constant.cu

@@ -29,7 +29,6 @@
 #include "llvm/Support/Debug.h"
 #include <list>
 #include <map>
-#include <sstream>
 
 template <typename T>
 static Fortran::semantics::Scope *GetScope(
@@ -61,6 +60,13 @@ template <typename T> class DirectiveAttributeVisitor {
         parser::OmpDefaultmapClause::ImplicitBehavior>
         defaultMap;
 
+    std::optional<Symbol::Flag> FindSymbolWithDSA(const Symbol &symbol) {
+      if (auto it{objectWithDSA.find(&symbol)}; it != objectWithDSA.end()) {
+        return it->second;
+      }
+      return std::nullopt;
+    }
+
     bool withinConstruct{false};
     std::int64_t associatedLoopLevel{0};
   };
@@ -75,10 +81,19 @@ template <typename T> class DirectiveAttributeVisitor {
         : std::make_optional<DirContext>(dirContext_.back());
   }
   void PushContext(const parser::CharBlock &source, T dir, Scope &scope) {
-    dirContext_.emplace_back(source, dir, scope);
+    if constexpr (std::is_same_v<T, llvm::acc::Directive>) {
+      dirContext_.emplace_back(source, dir, scope);
+      if (std::size_t size{dirContext_.size()}; size > 1) {
+        std::size_t lastIndex{size - 1};
+        dirContext_[lastIndex].defaultDSA =
+            dirContext_[lastIndex - 1].defaultDSA;
+      }
+    } else {
+      dirContext_.emplace_back(source, dir, scope);
+    }
   }
   void PushContext(const parser::CharBlock &source, T dir) {
-    dirContext_.emplace_back(source, dir, context_.FindScope(source));
+    PushContext(source, dir, context_.FindScope(source));
   }
   void PopContext() { dirContext_.pop_back(); }
   void SetContextDirectiveSource(parser::CharBlock &dir) {
@@ -100,9 +115,21 @@ template <typename T> class DirectiveAttributeVisitor {
     AddToContextObjectWithDSA(symbol, flag, GetContext());
   }
   bool IsObjectWithDSA(const Symbol &symbol) {
-    auto it{GetContext().objectWithDSA.find(&symbol)};
-    return it != GetContext().objectWithDSA.end();
+    return GetContext().FindSymbolWithDSA(symbol).has_value();
+  }
+  bool IsObjectWithVisibleDSA(const Symbol &symbol) {
+    for (std::size_t i{dirContext_.size()}; i != 0; i--) {
+      if (dirContext_[i - 1].FindSymbolWithDSA(symbol).has_value()) {
+        return true;
+      }
+    }
+    return false;
   }
+
+  bool WithinConstruct() {
+    return !dirContext_.empty() && GetContext().withinConstruct;
+  }
+
   void SetContextAssociatedLoopLevel(std::int64_t level) {
     GetContext().associatedLoopLevel = level;
   }
@@ -1573,10 +1600,10 @@ void AccAttributeVisitor::Post(const parser::AccDefaultClause &x) {
 // and adjust the symbol for each Name if necessary
 void AccAttributeVisitor::Post(const parser::Name &name) {
   auto *symbol{name.symbol};
-  if (symbol && !dirContext_.empty() && GetContext().withinConstruct) {
+  if (symbol && WithinConstruct()) {
     symbol = &symbol->GetUltimate();
     if (!symbol->owner().IsDerivedType() && !symbol->has<ProcEntityDetails>() &&
-        !symbol->has<SubprogramDetails>() && !IsObjectWithDSA(*symbol)) {
+        !symbol->has<SubprogramDetails>() && !IsObjectWithVisibleDSA(*symbol)) {
       if (Symbol * found{currScope().FindSymbol(name.source)}) {
         if (symbol != found) {
           name.symbol = found; // adjust the symbol within region
@@ -1959,7 +1986,7 @@ void OmpAttributeVisitor::ResolveSeqLoopIndexInParallelOrTaskConstruct(
 // till OpenMP-5.0 standard.
 // In above both cases we skip the privatization of iteration variables.
 bool OmpAttributeVisitor::Pre(const parser::DoConstruct &x) {
-  if (!dirContext_.empty() && GetContext().withinConstruct) {
+  if (WithinConstruct()) {
     llvm::SmallVector<const parser::Name *> ivs;
     if (x.IsDoNormal()) {
       const parser::Name *iv{GetLoopIndex(x)};
@@ -2685,7 +2712,7 @@ void OmpAttributeVisitor::CreateImplicitSymbols(const Symbol *symbol) {
 void OmpAttributeVisitor::Post(const parser::Name &name) {
   auto *symbol{name.symbol};
 
-  if (symbol && !dirContext_.empty() && GetContext().withinConstruct) {
+  if (symbol && WithinConstruct()) {
     if (IsPrivatizable(symbol) && !IsObjectWithDSA(*symbol)) {
       // TODO: create a separate function to go through the rules for
       //       predetermined, explicitly determined, and implicitly

flang/lib/Semantics/resolve-directives.cpp

@@ -200,3 +200,25 @@ program openacc_parallel_validity
   !$acc end parallel
 
 end program openacc_parallel_validity
+
+subroutine acc_parallel_default_none
+  integer :: i, l
+  real :: a(10,10)
+  l = 10  
+  !$acc parallel default(none)
+  !$acc loop
+  !ERROR: The DEFAULT(NONE) clause requires that 'l' must be listed in a data-mapping clause
+  do i = 1, l
+    !ERROR: The DEFAULT(NONE) clause requires that 'a' must be listed in a data-mapping clause
+    a(1,i) = 1
+  end do
+  !$acc end parallel
+
+  !$acc data copy(a)
+  !$acc parallel loop firstprivate(l) default(none)
+  do i = 1, l
+    a(1,i) = 1
+  end do
+  !$acc end parallel
+  !$acc end data
+end subroutine acc_parallel_default_none

flang/test/Semantics/OpenACC/acc-parallel.f90

@@ -10,6 +10,7 @@ add_subdirectory(lldb-fuzzer EXCLUDE_FROM_ALL)
 
 add_lldb_tool_subdirectory(lldb-instr)
 add_lldb_tool_subdirectory(lldb-dap)
+add_lldb_tool_subdirectory(lldb-mcp)
 if (LLDB_BUILD_LLDBRPC)
   add_lldb_tool_subdirectory(lldb-rpc-gen)
 endif()

lldb/tools/CMakeLists.txt

@@ -0,0 +1,33 @@
+add_lldb_tool(lldb-mcp
+  lldb-mcp.cpp
+
+  LINK_COMPONENTS
+    Option
+    Support
+  LINK_LIBS
+    liblldb
+    lldbHost
+    lldbProtocolMCP
+  )
+
+if(APPLE)
+  configure_file(
+    ${CMAKE_CURRENT_SOURCE_DIR}/lldb-mcp-Info.plist.in
+    ${CMAKE_CURRENT_BINARY_DIR}/lldb-mcp-Info.plist
+    )
+  target_link_options(lldb-mcp
+    PRIVATE LINKER:-sectcreate,__TEXT,__info_plist,${CMAKE_CURRENT_BINARY_DIR}/lldb-mcp-Info.plist)
+endif()
+
+if(LLDB_BUILD_FRAMEWORK)
+  # In the build-tree, we know the exact path to the framework directory.
+  # The installed framework can be in different locations.
+  lldb_setup_rpaths(lldb-mcp
+    BUILD_RPATH
+      "${LLDB_FRAMEWORK_ABSOLUTE_BUILD_DIR}"
+    INSTALL_RPATH
+      "@loader_path/../../../SharedFrameworks"
+      "@loader_path/../../System/Library/PrivateFrameworks"
+      "@loader_path/../../Library/PrivateFrameworks"
+  )
+endif()