[SYCL][NATIVECPU] Materialize floating point atomic builtins

Author: PietroGhg

clang/lib/Frontend/InitPreprocessor.cpp

@@ -1513,7 +1513,8 @@ static void InitializePredefinedMacros(const TargetInfo &TI,
     const llvm::Triple::SubArchType DeviceSubArch = DeviceTriple.getSubArch();
     if (DeviceTriple.isNVPTX() || DeviceTriple.isAMDGPU() ||
         (DeviceTriple.isSPIR() &&
-         DeviceSubArch != llvm::Triple::SPIRSubArch_fpga))
+         DeviceSubArch != llvm::Triple::SPIRSubArch_fpga) ||
+        LangOpts.SYCLIsNativeCPU)
       Builder.defineMacro("SYCL_USE_NATIVE_FP_ATOMICS");
     // Enable generation of USM address spaces for FPGA.
     if (DeviceSubArch == llvm::Triple::SPIRSubArch_fpga) {

llvm/include/llvm/SYCLLowerIR/FAtomicsNativeCPU.h

@@ -0,0 +1,29 @@
+//===------- FAtomicsNativeCPU.h - Materializes FP Atomics ----------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+//
+// A transformation pass that materializes floating points atomics by emitting
+// corresponding atomicrmw instruction.
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include "llvm/IR/Module.h"
+#include "llvm/IR/PassManager.h"
+
+namespace llvm {
+
+class ModulePass;
+
+class FAtomicsNativeCPU
+    : public PassInfoMixin<FAtomicsNativeCPU> {
+public:
+  PreservedAnalyses run(Module &M, ModuleAnalysisManager &MAM);
+};
+
+} // namespace llvm

llvm/lib/SYCLNativeCPUUtils/CMakeLists.txt

@@ -4,6 +4,7 @@ add_llvm_component_library(LLVMSYCLNativeCPUUtils
   RenameKernelSYCLNativeCPU.cpp
   ConvertToMuxBuiltinsSYCLNativeCPU.cpp
   FixABIMuxBuiltinsSYCLNativeCPU.cpp
+  FAtomicsNativeCPU.cpp
 
   ADDITIONAL_HEADER_DIRS
   ${LLVM_MAIN_INCLUDE_DIR}/llvm/SYCLLowerIR

llvm/lib/SYCLNativeCPUUtils/FAtomicsNativeCPU.cpp

@@ -0,0 +1,47 @@
+//===------- FAtomicsNativeCPU.cpp - Materializes FP Atomics --------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+//
+// A transformation pass that materializes floating points atomics by emitting
+// corresponding atomicrmw instruction.
+//
+//===----------------------------------------------------------------------===//
+
+#include "llvm/SYCLLowerIR/FAtomicsNativeCPU.h"
+#include "llvm/IR/IRBuilder.h"
+#include "llvm/IR/Instructions.h"
+#include "llvm/IR/LLVMContext.h"
+#include "llvm/Support/Alignment.h"
+#include "llvm/Support/AtomicOrdering.h"
+
+using namespace llvm;
+
+PreservedAnalyses FAtomicsNativeCPU::run(Module &M,
+                                         ModuleAnalysisManager &MAM) {
+  bool ModuleChanged = false;
+  auto &Ctx = M.getContext();
+  for (auto &F : M) {
+    AtomicRMWInst::BinOp OpCode;
+    if (F.getName().starts_with("_Z21__spirv_AtomicFAddEXT")) {
+      OpCode = AtomicRMWInst::BinOp::FAdd;
+    } else if (F.getName().starts_with("_Z21__spirv_AtomicFMinEXT")) {
+      OpCode = AtomicRMWInst::BinOp::FMin;
+    } else if (F.getName().starts_with("_Z21__spirv_AtomicFMaxEXT")) {
+      OpCode = AtomicRMWInst::BinOp::FMax;
+    } else {
+      continue;
+    }
+
+    BasicBlock *BB = BasicBlock::Create(Ctx, "entry", &F);
+    IRBuilder<> Builder(BB);
+    auto A =
+        Builder.CreateAtomicRMW(OpCode, F.getArg(0), F.getArg(3), MaybeAlign(),
+                                AtomicOrdering::Monotonic, SyncScope::System);
+    Builder.CreateRet(A);
+  }
+  return ModuleChanged ? PreservedAnalyses::none() : PreservedAnalyses::all();
+}

llvm/lib/SYCLNativeCPUUtils/PipelineSYCLNativeCPU.cpp

@@ -11,6 +11,7 @@
 // When NATIVECPU_USE_OCK is set, adds passes from the oneAPI Construction Kit.
 //
 //===----------------------------------------------------------------------===//
+#include "llvm/SYCLLowerIR/FAtomicsNativeCPU.h"
 #include "llvm/SYCLLowerIR/ConvertToMuxBuiltinsSYCLNativeCPU.h"
 #include "llvm/SYCLLowerIR/FixABIMuxBuiltinsSYCLNativeCPU.h"
 #include "llvm/SYCLLowerIR/PrepareSYCLNativeCPU.h"
@@ -70,6 +71,7 @@ void llvm::sycl::utils::addSYCLNativeCPUBackendPasses(
     OptimizationLevel OptLevel) {
   MPM.addPass(SpecConstantsPass(SpecConstantsPass::HandlingMode::emulation));
   MPM.addPass(ConvertToMuxBuiltinsSYCLNativeCPUPass());
+  MPM.addPass(FAtomicsNativeCPU());
 #ifdef NATIVECPU_USE_OCK
   MPM.addPass(compiler::utils::PrepareBarriersPass());
   MPM.addPass(compiler::utils::TransferKernelMetadataPass());