Merge remote-tracking branch 'origin/sycl' into benie/cl-binary-type-…

…intermediate

.github/workflows/sycl_containers.yaml → .github/workflows/sycl-containers.yaml

@@ -12,14 +12,14 @@ on:
  - 'devops/dependencies.json'
  - 'devops/scripts/install_drivers.sh'
  - 'devops/scripts/install_build_tools.sh'
- - '.github/workflows/sycl_containers.yaml'
+ - '.github/workflows/sycl-containers.yaml'
  pull_request:
  paths:
  - 'devops/containers/**'
  - 'devops/dependencies.json'
  - 'devops/scripts/install_drivers.sh'
  - 'devops/scripts/install_build_tools.sh'
- - '.github/workflows/sycl_containers.yaml'
+ - '.github/workflows/sycl-containers.yaml'
 
 jobs:
  base_image_ubuntu2204:

.github/workflows/gh_pages.yml → .github/workflows/sycl-docs.yml

@@ -7,7 +7,7 @@ on:
  branches:
  - sycl
  paths:
- - '.github/workflows/gh_pages.yml'
+ - '.github/workflows/sycl-docs.yml'
  - 'clang/docs/**'
  - 'sycl/doc/**'
 

.github/workflows/sycl_linux_matrix_e2e_on_nightly.yml → .github/workflows/sycl-linux-matrix-e2e-on-nightly.yml

@@ -48,7 +48,7 @@ jobs:
  image: ghcr.io/intel/llvm/sycl_ubuntu2204_nightly:build
  image_options: -u 1001 --gpus all --cap-add SYS_ADMIN
  target_devices: ext_oneapi_cuda:gpu
- uses: ./.github/workflows/sycl_linux_run_tests.yml
+ uses: ./.github/workflows/sycl-linux-run-tests.yml
  with:
  name: ${{ matrix.name }}
  runner: ${{ matrix. runner }}
@@ -62,7 +62,7 @@ jobs:
 
  aws_start:
  name: AWS Start
- uses: ./.github/workflows/sycl_aws.yml
+ uses: ./.github/workflows/sycl-aws.yml
  secrets: inherit
  with:
  mode: start
@@ -71,7 +71,7 @@ jobs:
  linux_e2e_on_nightly_aws:
  name: '[AWS][CUDA] E2E on Nightly'
  needs: [aws_start]
- uses: ./.github/workflows/sycl_linux_run_tests.yml
+ uses: ./.github/workflows/sycl-linux-run-tests.yml
  with:
  name: CUDA
  runner: '["aws-cuda_${{ github.run_id }}-${{ github.run_attempt }}"]'
@@ -86,7 +86,7 @@ jobs:
  name: AWS Stop
  needs: [aws_start, linux_e2e_on_nightly_aws]
  if: always()
- uses: ./.github/workflows/sycl_aws.yml
+ uses: ./.github/workflows/sycl-aws.yml
  secrets: inherit
  with:
  mode: stop

.github/workflows/sycl_linux_precommit_aws.yml → .github/workflows/sycl-linux-precommit-aws.yml

@@ -54,7 +54,7 @@ jobs:
 
  e2e-cuda:
  needs: [aws-start]
- uses: ./.github/workflows/sycl_linux_run_tests.yml
+ uses: ./.github/workflows/sycl-linux-run-tests.yml
  with:
  name: CUDA E2E
  runner: '["aws_cuda-${{ github.event.workflow_run.id }}-${{ github.event.workflow_run.run_attempt }}"]'

.github/workflows/sycl_linux_precommit.yml → .github/workflows/sycl-linux-precommit.yml

@@ -18,8 +18,8 @@ on:
  - 'clang/docs/**'
  - '**.md'
  - '**.rst'
- - '.github/workflows/sycl_windows_*.yml'
- - '.github/workflows/sycl_macos_*.yml'
+ - '.github/workflows/sycl-windows-*.yml'
+ - '.github/workflows/sycl-macos-*.yml'
  - 'devops/containers/**'
  - 'devops/actions/build_container/**'
 
@@ -30,12 +30,12 @@ concurrency:
 
 jobs:
  detect_changes:
- uses: ./.github/workflows/sycl_detect_changes.yml
+ uses: ./.github/workflows/sycl-detect-changes.yml
 
  build:
  needs: [detect_changes]
  if: always() && success()
- uses: ./.github/workflows/sycl_linux_build.yml
+ uses: ./.github/workflows/sycl-linux-build.yml
  with:
  build_ref: ${{ github.sha }}
  merge_ref: ''
@@ -75,7 +75,7 @@ jobs:
  install_drivers: ${{ contains(needs.detect_changes.outputs.filters, 'drivers') }}
  extra_lit_opts: --param matrix-xmx8=True --param gpu-intel-dg2=True
  env: '{"LIT_FILTER":"Matrix/"}'
- uses: ./.github/workflows/sycl_linux_run_tests.yml
+ uses: ./.github/workflows/sycl-linux-run-tests.yml
  with:
  name: ${{ matrix.name }}
  runner: ${{ matrix. runner }}
@@ -111,7 +111,7 @@ jobs:
  runner: '["Linux", "gen12"]'
  - name: Perf tests on Intel Arc A-Series Graphics system
  runner: '["Linux", "arc"]'
- uses: ./.github/workflows/sycl_linux_run_tests.yml
+ uses: ./.github/workflows/sycl-linux-run-tests.yml
  with:
  name: ${{ matrix.name }}
  runner: ${{ matrix. runner }}

.github/workflows/sycl_nightly.yml → .github/workflows/sycl-nightly.yml

@@ -8,7 +8,7 @@ on:
 jobs:
  ubuntu2204_build:
  if: github.repository == 'intel/llvm'
- uses: ./.github/workflows/sycl_linux_build.yml
+ uses: ./.github/workflows/sycl-linux-build.yml
  secrets: inherit
  with:
  build_cache_root: "/__w/"
@@ -59,7 +59,7 @@ jobs:
  image: ghcr.io/intel/llvm/ubuntu2204_build:latest
  image_options: -u 1001 --gpus all --cap-add SYS_ADMIN
  target_devices: ext_oneapi_cuda:gpu
- uses: ./.github/workflows/sycl_linux_run_tests.yml
+ uses: ./.github/workflows/sycl-linux-run-tests.yml
  with:
  name: ${{ matrix.name }}
  runner: ${{ matrix.runner }}
@@ -74,7 +74,7 @@ jobs:
  sycl_toolchain_decompress_command: ${{ needs.ubuntu2204_build.outputs.artifact_decompress_command }}
 
  build-win:
- uses: ./.github/workflows/sycl_windows_build.yml
+ uses: ./.github/workflows/sycl-windows-build.yml
  if: github.repository == 'intel/llvm'
  with:
  retention-days: 90
@@ -89,7 +89,7 @@ jobs:
  always()
  && !cancelled()
  && needs.build.outputs.build_conclusion == 'success'
- uses: ./.github/workflows/sycl_windows_run_tests.yml
+ uses: ./.github/workflows/sycl-windows-run-tests.yml
  with:
  name: Intel GEN12 Graphics with Level Zero
  runner: '["Windows","gen12"]'

.github/workflows/sycl_post_commit.yml → .github/workflows/sycl-post-commit.yml

@@ -12,18 +12,18 @@ on:
  - sycl
  - sycl-devops-pr/**
  paths:
- - .github/workflows/sycl_post_commit.yml
- - .github/workflows/sycl_linux_build.yml
- - .github/workflows/sycl_linux_run_tests.yml
- - .github/workflows/sycl_macos_build_and_test.yml
+ - .github/workflows/sycl-post-commit.yml
+ - .github/workflows/sycl-linux-build.yml
+ - .github/workflows/sycl-linux-run-tests.yml
+ - .github/workflows/sycl-macos-build-and-test.yml
  - ./devops/actions/cleanup
  - ./devops/actions/cached_checkout
 
 jobs:
  build-lin:
  name: Linux (Self build + shared libraries + no-assertions)
  if: github.repository == 'intel/llvm'
- uses: ./.github/workflows/sycl_linux_build.yml
+ uses: ./.github/workflows/sycl-linux-build.yml
  with:
  build_cache_root: "/__w/llvm"
  build_cache_suffix: sprod_shared
@@ -63,7 +63,7 @@ jobs:
  env: '{"LIT_FILTER":"PerformanceTests/"}'
  extra_lit_opts: -a -j 1 --param enable-perf-tests=True
  target_devices: all
- uses: ./.github/workflows/sycl_linux_run_tests.yml
+ uses: ./.github/workflows/sycl-linux-run-tests.yml
  with:
  name: ${{ matrix.name }}
  runner: ${{ matrix. runner }}
@@ -87,7 +87,7 @@ jobs:
  always()
  && success()
  && github.repository == 'intel/llvm'
- uses: ./.github/workflows/sycl_windows_build.yml
+ uses: ./.github/workflows/sycl-windows-build.yml
 
  e2e-win:
  needs: build-win
@@ -96,7 +96,7 @@ jobs:
  always()
  && !cancelled()
  && needs.build-win.outputs.build_conclusion == 'success'
- uses: ./.github/workflows/sycl_windows_run_tests.yml
+ uses: ./.github/workflows/sycl-windows-run-tests.yml
  with:
  name: Intel GEN12 Graphics with Level Zero
  runner: '["Windows","gen12"]'
@@ -105,4 +105,4 @@ jobs:
  macos_default:
  name: macOS
  if: github.repository == 'intel/llvm'
- uses: ./.github/workflows/sycl_macos_build_and_test.yml
+ uses: ./.github/workflows/sycl-macos-build-and-test.yml

.github/workflows/sycl_windows_precommit.yml → .github/workflows/sycl-windows-precommit.yml

@@ -16,9 +16,9 @@ on:
  - 'clang/docs/**'
  - '**.md'
  - '**.rst'
- - '.github/workflows/sycl_linux_*.yml'
- - '.github/workflows/sycl_precommit_aws.yml'
- - '.github/workflows/sycl_macos_*.yml'
+ - '.github/workflows/sycl-linux-*.yml'
+ - '.github/workflows/sycl-precommit-aws.yml'
+ - '.github/workflows/sycl-macos-*.yml'
  - 'devops/containers/**'
  - 'devops/actions/build_container/**'
 
@@ -32,14 +32,14 @@ concurrency:
 
 jobs:
  detect_changes:
- uses: ./.github/workflows/sycl_detect_changes.yml
+ uses: ./.github/workflows/sycl-detect-changes.yml
 
  build:
  needs: [detect_changes]
  if: |
  always() && success()
  && github.repository == 'intel/llvm'
- uses: ./.github/workflows/sycl_windows_build.yml
+ uses: ./.github/workflows/sycl-windows-build.yml
  with:
  changes: ${{ needs.detect_changes.outputs.filters }}
 
@@ -50,7 +50,7 @@ jobs:
  always()
  && !cancelled()
  && needs.build.outputs.build_conclusion == 'success'
- uses: ./.github/workflows/sycl_windows_run_tests.yml
+ uses: ./.github/workflows/sycl-windows-run-tests.yml
  with:
  name: Intel GEN12 Graphics with Level Zero
  runner: '["Windows","gen12"]'

README.md

@@ -10,8 +10,8 @@ For general contribution process see [CONTRIBUTING.md](./CONTRIBUTING.md)
 
 [![oneAPI logo](https://spec.oneapi.io/oneapi-logo-white-scaled.jpg)](https://www.oneapi.io/)
 
-[![SYCL Post Commit](https://github.com/intel/llvm/actions/workflows/sycl_post_commit.yml/badge.svg?branch=sycl)](https://github.com/intel/llvm/actions/workflows/sycl_post_commit.yml)
-[![Generate Doxygen documentation](https://github.com/intel/llvm/actions/workflows/gh_pages.yml/badge.svg?branch=sycl)](https://github.com/intel/llvm/actions/workflows/gh_pages.yml)
+[![SYCL Post Commit](https://github.com/intel/llvm/actions/workflows/sycl-post-commit.yml/badge.svg?branch=sycl)](https://github.com/intel/llvm/actions/workflows/sycl-post-commit.yml)
+[![Generate Doxygen documentation](https://github.com/intel/llvm/actions/workflows/sycl-docs.yml/badge.svg?branch=sycl)](https://github.com/intel/llvm/actions/workflows/sycl-docs.yml)
 
 The DPC++ is a LLVM-based compiler project that implements compiler and runtime
 support for the SYCL\* language. The project is hosted in the

clang/include/clang/Driver/Options.td

@@ -8404,6 +8404,9 @@ def : CLFlag<"O1">, Alias<_SLASH_O>, AliasArgs<["1"]>,
  HelpText<"Optimize for size (like /Og /Os /Oy /Ob2 /GF /Gy)">;
 def : CLFlag<"O2">, Alias<_SLASH_O>, AliasArgs<["2"]>,
  HelpText<"Optimize for speed (like /Og /Oi /Ot /Oy /Ob2 /GF /Gy)">;
+def : CLFlag<"O3">, Alias<_SLASH_O>, AliasArgs<["3"]>,
+ HelpText<"Optimize for maximum speed and enable more aggressive optimizations"
+ " that may not improve performance on some programs">;
 def : CLFlag<"Ob0">, Alias<_SLASH_O>, AliasArgs<["b0"]>,
  HelpText<"Disable function inlining">;
 def : CLFlag<"Ob1">, Alias<_SLASH_O>, AliasArgs<["b1"]>,

clang/lib/CodeGen/CGBuiltin.cpp

@@ -23271,21 +23271,11 @@ static bool hasFuncNameRequestedFPAccuracy(StringRef Name,
  return (FuncMapIt != LangOpts.FPAccuracyFuncMap.end());
 }
 
-llvm::CallInst *CodeGenFunction::EmitFPBuiltinIndirectCall(
+llvm::CallInst *CodeGenFunction::MaybeEmitFPBuiltinofFD(
  llvm::FunctionType *IRFuncTy, const SmallVectorImpl<llvm::Value *> &IRArgs,
- llvm::Value *FnPtr, const FunctionDecl *FD) {
- llvm::Function *Func;
+ llvm::Value *FnPtr, StringRef Name, unsigned FDBuiltinID) {
  unsigned FPAccuracyIntrinsicID = 0;
- StringRef Name;
- if (CurrentBuiltinID == 0) {
- // Even if the current function doesn't have a clang builtin, create
- // an 'fpbuiltin-max-error' attribute for it; unless it's marked with
- // an NoBuiltin attribute.
- if (FD->hasAttr<NoBuiltinAttr>() ||
- !FD->getNameInfo().getName().isIdentifier())
- return nullptr;
-
- Name = FD->getName();
+ if (FDBuiltinID == 0) {
  FPAccuracyIntrinsicID =
  llvm::StringSwitch<unsigned>(Name)
  .Case("fadd", llvm::Intrinsic::fpbuiltin_fadd)
@@ -23300,9 +23290,7 @@ llvm::CallInst *CodeGenFunction::EmitFPBuiltinIndirectCall(
  } else {
  // The function has a clang builtin. Create an attribute for it
  // only if it has an fpbuiltin intrinsic.
- unsigned BuiltinID = getCurrentBuiltinID();
- Name = CGM.getContext().BuiltinInfo.getName(BuiltinID);
- switch (BuiltinID) {
+ switch (FDBuiltinID) {
  default:
  // If the function has a clang builtin but doesn't have an
  // fpbuiltin, it will be generated with no 'fpbuiltin-max-error'
@@ -23384,7 +23372,8 @@ llvm::CallInst *CodeGenFunction::EmitFPBuiltinIndirectCall(
  const LangOptions &LangOpts = getLangOpts();
  if (hasFuncNameRequestedFPAccuracy(Name, LangOpts) ||
  !LangOpts.FPAccuracyVal.empty()) {
- Func = CGM.getIntrinsic(FPAccuracyIntrinsicID, IRArgs[0]->getType());
+ llvm::Function *Func =
+ CGM.getIntrinsic(FPAccuracyIntrinsicID, IRArgs[0]->getType());
  return CreateBuiltinCallWithAttr(*this, Name, Func, ArrayRef(IRArgs),
  FPAccuracyIntrinsicID);
  }

clang/lib/CodeGen/CGCall.cpp

@@ -5707,8 +5707,9 @@ RValue CodeGenFunction::EmitCall(const CGFunctionInfo &CallInfo,
  if (!getLangOpts().FPAccuracyFuncMap.empty() ||
  !getLangOpts().FPAccuracyVal.empty()) {
  const auto *FD = dyn_cast_if_present<FunctionDecl>(TargetDecl);
- if (FD) {
- CI = EmitFPBuiltinIndirectCall(IRFuncTy, IRCallArgs, CalleePtr, FD);
+ if (FD && FD->getNameInfo().getName().isIdentifier()) {
+ CI = MaybeEmitFPBuiltinofFD(IRFuncTy, IRCallArgs, CalleePtr,
+ FD->getName(), FD->getBuiltinID());
  if (CI)
  return RValue::get(CI);
  }

clang/lib/CodeGen/CodeGenFunction.h

@@ -4434,9 +4434,10 @@ class CodeGenFunction : public CodeGenTypeCache {
  RValue EmitIntelSYCLPtrAnnotationBuiltin(const CallExpr *E);
 
  llvm::CallInst *
- EmitFPBuiltinIndirectCall(llvm::FunctionType *IRFuncTy,
- const SmallVectorImpl<llvm::Value *> &IRArgs,
- llvm::Value *FnPtr, const FunctionDecl *FD);
+ MaybeEmitFPBuiltinofFD(llvm::FunctionType *IRFuncTy,
+ const SmallVectorImpl<llvm::Value *> &IRArgs,
+ llvm::Value *FnPtr, StringRef Name,
+ unsigned FDBuiltinID);
 
  enum class MSVCIntrin;
  llvm::Value *EmitMSVCBuiltinExpr(MSVCIntrin BuiltinID, const CallExpr *E);

clang/lib/CodeGen/CodeGenTypes.cpp

@@ -342,8 +342,8 @@ llvm::Type *CodeGenTypes::ConvertSYCLJointMatrixINTELType(RecordDecl *RD) {
  if (CompTy->isStructTy()) {
  StringRef LlvmTyName = CompTy->getStructName();
  // Emit half/int16/float for sycl[::*]::{half,bfloat16,tf32}
- if (LlvmTyName.startswith("class.sycl::") ||
- LlvmTyName.startswith("class.__sycl_internal::"))
+ if (LlvmTyName.starts_with("class.sycl::") ||
+ LlvmTyName.starts_with("class.__sycl_internal::"))
  LlvmTyName = LlvmTyName.rsplit("::").second;
  if (LlvmTyName == "half") {
  CompTy = llvm::Type::getHalfTy(getLLVMContext());

clang/lib/Driver/Driver.cpp

@@ -3518,7 +3518,7 @@ getLinkerArgs(Compilation &C, DerivedArgList &Args, bool IncludeObj = false) {
  // manner than the OpenMP processing. We should try and refactor this
  // to use the OpenMP flow (adding -l<name> to the llvm-link step)
  auto resolveStaticLib = [&](StringRef LibName, bool IsStatic) -> bool {
- if (!LibName.startswith("-l"))
+ if (!LibName.starts_with("-l"))
  return false;
  for (auto &LPath : LibPaths) {
  if (!IsStatic) {
@@ -3663,7 +3663,7 @@ static bool IsSYCLDeviceLibObj(std::string ObjFilePath, bool isMSVCEnv) {
  StringRef ObjFileName = llvm::sys::path::filename(ObjFilePath);
  StringRef ObjSuffix = isMSVCEnv ? ".obj" : ".o";
  bool Ret =
- (ObjFileName.startswith("libsycl-") && ObjFileName.endswith(ObjSuffix))
+ (ObjFileName.starts_with("libsycl-") && ObjFileName.ends_with(ObjSuffix))
  ? true
  : false;
  return Ret;

clang/lib/Driver/ToolChains/Linux.cpp

@@ -344,7 +344,7 @@ Linux::Linux(const Driver &D, const llvm::Triple &Triple, const ArgList &Args)
  // The deprecated -DLLVM_ENABLE_PROJECTS=libcxx configuration installs
  // libc++.so in D.Dir+"/../lib/". Detect this path.
  // TODO Remove once LLVM_ENABLE_PROJECTS=libcxx is unsupported.
- if (StringRef(D.Dir).startswith(SysRoot) &&
+ if (StringRef(D.Dir).starts_with(SysRoot) &&
  (Args.hasArg(options::OPT_fsycl) ||
  D.getVFS().exists(D.Dir + "/../lib/libsycl.so")))
  addPathIfExists(D, D.Dir + "/../lib", Paths);