merge main into amd-staging (llvm#726)

Author: ronlieb

clang-tools-extra/clangd/test/CMakeLists.txt

@@ -5,6 +5,10 @@ set(CLANGD_TEST_DEPS
   split-file
   )
 
+if (LLVM_INCLUDE_BENCHMARKS)
+  list(APPEND CLANGD_TEST_DEPS IndexBenchmark)
+endif()
+
 if(CLANGD_BUILD_XPC)
   list(APPEND CLANGD_TEST_DEPS clangd-xpc-test-client)
   list(APPEND CLANGD_TEST_DEPS ClangdXpcUnitTests)

clang-tools-extra/clangd/test/include-cleaner-batch-fix.test

@@ -7,7 +7,9 @@
 # RUN: cp -r %S/Inputs/include-cleaner %t/include
 # RUN: echo '-I%t/include' > %t/compile_flags.txt
 # Create a config file enabling include-cleaner features.
-# RUN: echo $'Diagnostics:\n  UnusedIncludes: Strict\n  MissingIncludes: Strict' >> %t/clangd/config.yaml
+# RUN: echo 'Diagnostics:' > %t/clangd/config.yaml
+# RUN: echo '  UnusedIncludes: Strict' >> %t/clangd/config.yaml
+# RUN: echo '  MissingIncludes: Strict' >> %t/clangd/config.yaml
 
 # RUN: env XDG_CONFIG_HOME=%t clangd -lit-test -enable-config --compile-commands-dir=%t < %s | FileCheck -strict-whitespace %s
 {"jsonrpc":"2.0","id":0,"method":"initialize","params":{"processId":123,"rootPath":"clangd","capabilities":{"workspace":{"workspaceEdit":{"documentChanges":true, "changeAnnotationSupport":{"groupsOnLabel":true}}}},"trace":"off"}}

clang-tools-extra/clangd/test/index-tools.test

@@ -1,6 +1,7 @@
+# Paths are not constructed correctly for the test to run on Windows.
+# UNSUPPORTED: system-windows
+# REQUIRES: have-benchmarks
 # RUN: clangd-indexer %p/Inputs/BenchmarkSource.cpp -- -I%p/Inputs > %t.index
-# FIXME: By default, benchmarks are excluded from the list of default targets hence not built. Find a way to depend on benchmarks to run the next command.
-# REQUIRES: shell
-# RUN: if [ -f %clangd-benchmark-dir/IndexBenchmark ]; then %clangd-benchmark-dir/IndexBenchmark %t.index %p/Inputs/requests.json --benchmark_min_time=0.01 ; fi
+# RUN: %clangd-benchmark-dir/IndexBenchmark %t.index %p/Inputs/requests.json --benchmark_min_time=0.01
 # Pass invalid JSON file and check that IndexBenchmark fails to parse it.
-# RUN: if [ -f %clangd-benchmark-dir/IndexBenchmark ]; then not %clangd-benchmark-dir/IndexBenchmark %t.index %t --benchmark_min_time=0.01 ; fi
+# RUN: not %clangd-benchmark-dir/IndexBenchmark %t.index %t --benchmark_min_time=0.01

clang-tools-extra/clangd/test/lit.cfg.py

@@ -1,4 +1,5 @@
 import lit.llvm
+import lit.util
 
 lit.llvm.initialize(lit_config, config)
 lit.llvm.llvm_config.clang_setup()
@@ -37,6 +38,9 @@ def calculate_arch_features(arch_string):
 if config.have_zlib:
     config.available_features.add("zlib")
 
+if lit.util.pythonize_bool(config.have_benchmarks):
+    config.available_features.add("have-benchmarks")
+
 # It is not realistically possible to account for all options that could
 # possibly be present in system and user configuration files, so disable
 # default configs for the test runs.

clang-tools-extra/clangd/test/lit.site.cfg.py.in

@@ -19,6 +19,7 @@ config.clangd_build_dexp = @CLANGD_BUILD_DEXP@
 config.clangd_enable_remote = @CLANGD_ENABLE_REMOTE@
 config.clangd_tidy_checks = @CLANGD_TIDY_CHECKS@
 config.have_zlib = @LLVM_ENABLE_ZLIB@
+config.have_benchmarks = "@LLVM_INCLUDE_BENCHMARKS@"
 
 # Delegate logic to lit.cfg.py.
 lit_config.load_config(config, "@CMAKE_CURRENT_SOURCE_DIR@/lit.cfg.py")

clang-tools-extra/clangd/test/system-include-extractor.test

@@ -5,7 +5,8 @@
 
 # Create a bin directory to store the mock-driver and add it to the path
 # RUN: mkdir -p %t.dir/bin
-# RUN: export PATH=%t.dir/bin:$PATH
+# RUN: %python -c "print(__import__('os').environ['PATH'])" > %t.path
+# RUN: export PATH=%t.dir/bin:%{readfile:%t.path}
 # Generate a mock-driver that will print %temp_dir%/my/dir and
 # %temp_dir%/my/dir2 as include search paths.
 # RUN: echo '#!/bin/sh' >> %t.dir/bin/my_driver.sh

clang/include/clang/AST/ASTContext.h

@@ -488,6 +488,10 @@ class ASTContext : public RefCountedBase<ASTContext> {
 
   /// Declaration for the CUDA cudaConfigureCall function.
   FunctionDecl *cudaConfigureCallDecl = nullptr;
+  /// Declaration for the CUDA cudaGetParameterBuffer function.
+  FunctionDecl *cudaGetParameterBufferDecl = nullptr;
+  /// Declaration for the CUDA cudaLaunchDevice function.
+  FunctionDecl *cudaLaunchDeviceDecl = nullptr;
 
   /// Keeps track of all declaration attributes.
   ///
@@ -1641,6 +1645,18 @@ class ASTContext : public RefCountedBase<ASTContext> {
     return cudaConfigureCallDecl;
   }
 
+  void setcudaGetParameterBufferDecl(FunctionDecl *FD) {
+    cudaGetParameterBufferDecl = FD;
+  }
+
+  FunctionDecl *getcudaGetParameterBufferDecl() {
+    return cudaGetParameterBufferDecl;
+  }
+
+  void setcudaLaunchDeviceDecl(FunctionDecl *FD) { cudaLaunchDeviceDecl = FD; }
+
+  FunctionDecl *getcudaLaunchDeviceDecl() { return cudaLaunchDeviceDecl; }
+
   /// Returns true iff we need copy/dispose helpers for the given type.
   bool BlockRequiresCopying(QualType Ty, const VarDecl *D);
 

clang/include/clang/Basic/BuiltinsX86.td

@@ -3160,15 +3160,15 @@ let Features = "avx512bw", Attributes = [NoThrow, Const] in {
   def kshiftridi : X86Builtin<"unsigned long long int(unsigned long long int, _Constant unsigned int)">;
 }
 
-let Features = "avx512dq", Attributes = [NoThrow, Const] in {
+let Features = "avx512dq", Attributes = [NoThrow, Const, Constexpr] in {
   def kmovb : X86Builtin<"unsigned char(unsigned char)">;
 }
 
-let Features = "avx512f", Attributes = [NoThrow, Const] in {
+let Features = "avx512f", Attributes = [NoThrow, Const, Constexpr] in {
   def kmovw : X86Builtin<"unsigned short(unsigned short)">;
 }
 
-let Features = "avx512bw", Attributes = [NoThrow, Const] in {
+let Features = "avx512bw", Attributes = [NoThrow, Const, Constexpr] in {
   def kmovd : X86Builtin<"unsigned int(unsigned int)">;
   def kmovq : X86Builtin<"unsigned long long int(unsigned long long int)">;
 }

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -9528,6 +9528,8 @@ def err_kern_is_nonstatic_method : Error<
   "kernel function %0 must be a free function or static member function">;
 def err_config_scalar_return : Error<
   "CUDA special function '%0' must have scalar return type">;
+def err_config_pointer_return
+    : Error<"CUDA special function '%0' must have pointer return type">;
 def err_kern_call_not_global_function : Error<
   "kernel call to non-global function %0">;
 def err_global_call_not_config : Error<
@@ -13770,4 +13772,9 @@ def warn_comparison_in_enum_initializer : Warning<
 def note_enum_compare_typo_suggest : Note<
   "use '%0' to perform a bitwise shift">;
 
+def err_cuda_device_kernel_launch_not_supported
+    : Error<"device-side kernel call/launch is not supported">;
+def err_cuda_device_kernel_launch_require_rdc
+    : Error<"kernel launch from __device__ or __global__ function requires "
+            "relocatable device code (i.e. requires -fgpu-rdc)">;
 } // end of sema component.

clang/include/clang/Sema/SemaCUDA.h

@@ -273,6 +273,11 @@ class SemaCUDA : public SemaBase {
   /// of the function that will be called to configure kernel call, with the
   /// parameters specified via <<<>>>.
   std::string getConfigureFuncName() const;
+  /// Return the name of the parameter buffer allocation function for the
+  /// device kernel launch.
+  std::string getGetParameterBufferFuncName() const;
+  /// Return the name of the device kernel launch function.
+  std::string getLaunchDeviceFuncName() const;
 
   /// Record variables that are potentially ODR-used in CUDA/HIP.
   void recordPotentialODRUsedVariable(MultiExprArg Args,