From 44094f8780e31c658a2f6d3e7d7f0da4d6a4580e Mon Sep 17 00:00:00 2001 From: Vyacheslav N Klochkov Date: Thu, 11 May 2023 13:00:47 -0700 Subject: [PATCH 1/6] [ESIMD] Enable -fsycl-esimd-force-stateless-mem by default Signed-off-by: Vyacheslav N Klochkov --- clang/include/clang/Driver/Options.td | 4 ++-- clang/lib/Driver/ToolChains/Clang.cpp | 17 +++++++++++------ clang/lib/Frontend/InitPreprocessor.cpp | 4 +++- .../Driver/sycl-esimd-force-stateless-mem.cpp | 19 +++++++++++-------- .../sycl-esimd-force-stateless-mem.cpp | 8 +++++--- llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp | 2 +- llvm/test/SYCLLowerIR/ESIMD/acc_ptr.ll | 17 ++++++++++++----- sycl/test/esimd/lsc.cpp | 4 ++-- 8 files changed, 47 insertions(+), 28 deletions(-) diff --git a/clang/include/clang/Driver/Options.td b/clang/include/clang/Driver/Options.td index f1fe71eec2c19..0f6f86296c125 100644 --- a/clang/include/clang/Driver/Options.td +++ b/clang/include/clang/Driver/Options.td @@ -3170,10 +3170,10 @@ def fsycl_device_only : Flag<["-"], "fsycl-device-only">, Flags<[CoreOption]>, def fsycl_embed_ir : Flag<["-"], "fsycl-embed-ir">, Flags<[CoreOption]>, HelpText<"Embed LLVM IR for runtime kernel fusion">; defm sycl_esimd_force_stateless_mem : BoolFOption<"sycl-esimd-force-stateless-mem", - LangOpts<"SYCLESIMDForceStatelessMem">, DefaultFalse, + LangOpts<"SYCLESIMDForceStatelessMem">, DefaultTrue, PosFlag, + "Enabled by default. (experimental)">, NegFlag, BothFlags<[CC1Option, CoreOption], "">>; diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 96860154ef762..ff5209550eed2 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -5294,6 +5294,10 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA, options::OPT_fno_sycl_unnamed_lambda, true)) CmdArgs.push_back("-fno-sycl-unnamed-lambda"); + if (!Args.hasFlag(options::OPT_fsycl_esimd_force_stateless_mem, + options::OPT_fno_sycl_esimd_force_stateless_mem, true)) + CmdArgs.push_back("-fno-sycl-esimd-force-stateless-mem"); + // Add the Unique ID prefix StringRef UniqueID = D.getSYCLUniqueID(Input.getBaseInput()); if (!UniqueID.empty()) @@ -5390,9 +5394,10 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA, for (auto &Macro : D.getSYCLTargetMacroArgs()) CmdArgs.push_back(Args.MakeArgString(Macro)); } - if (Args.hasFlag(options::OPT_fsycl_esimd_force_stateless_mem, - options::OPT_fno_sycl_esimd_force_stateless_mem, false)) - CmdArgs.push_back("-fsycl-esimd-force-stateless-mem"); + if (IsSYCL && + !Args.hasFlag(options::OPT_fsycl_esimd_force_stateless_mem, + options::OPT_fno_sycl_esimd_force_stateless_mem, true)) + CmdArgs.push_back("-fno-sycl-esimd-force-stateless-mem"); const auto DeviceTraitsMacrosArgs = D.getDeviceTraitsMacrosArgs(); for (const auto &Arg : DeviceTraitsMacrosArgs) { @@ -10065,9 +10070,9 @@ void SYCLPostLink::ConstructJob(Compilation &C, const JobAction &JA, addArgs(CmdArgs, TCArgs, {"-device-globals"}); // Make ESIMD accessors use stateless memory accesses. - if (TCArgs.hasFlag(options::OPT_fsycl_esimd_force_stateless_mem, - options::OPT_fno_sycl_esimd_force_stateless_mem, false)) - addArgs(CmdArgs, TCArgs, {"-lower-esimd-force-stateless-mem"}); + if (TCArgs.hasFlag(options::OPT_fno_sycl_esimd_force_stateless_mem, + options::OPT_fsycl_esimd_force_stateless_mem, false)) + addArgs(CmdArgs, TCArgs, {"-lower-esimd-force-stateless-mem=false"}); // Add output file table file option assert(Output.isFilename() && "output must be a filename"); diff --git a/clang/lib/Frontend/InitPreprocessor.cpp b/clang/lib/Frontend/InitPreprocessor.cpp index 9c56cbee5b890..7348edf4623c3 100644 --- a/clang/lib/Frontend/InitPreprocessor.cpp +++ b/clang/lib/Frontend/InitPreprocessor.cpp @@ -1344,7 +1344,9 @@ static void InitializePredefinedMacros(const TargetInfo &TI, if (LangOpts.SYCLUnnamedLambda) Builder.defineMacro("__SYCL_UNNAMED_LAMBDA__"); - if (LangOpts.SYCLESIMDForceStatelessMem) + // Stateless memory may be enforced only for SYCL device or host. + if ((LangOpts.SYCLIsDevice || LangOpts.SYCLIsHost) && + LangOpts.SYCLESIMDForceStatelessMem) Builder.defineMacro("__ESIMD_FORCE_STATELESS_MEM"); // OpenCL definitions. diff --git a/clang/test/Driver/sycl-esimd-force-stateless-mem.cpp b/clang/test/Driver/sycl-esimd-force-stateless-mem.cpp index 90e506a86f53c..df34824038551 100644 --- a/clang/test/Driver/sycl-esimd-force-stateless-mem.cpp +++ b/clang/test/Driver/sycl-esimd-force-stateless-mem.cpp @@ -1,14 +1,17 @@ /// Verify that the driver option is translated to corresponding options /// to host/device compilation and sycl-post-link. -// RUN: %clang -### -fsycl -fsycl-esimd-force-stateless-mem \ -// RUN: %s 2>&1 | FileCheck -check-prefix=CHECK-PASS-TO-COMPS %s -// CHECK-PASS-TO-COMPS: clang{{.*}} "-fsycl-esimd-force-stateless-mem" -// CHECK-PASS-TO-COMPS: sycl-post-link{{.*}} "-lower-esimd-force-stateless-mem" -// CHECK-PASS-TO-COMPS: clang{{.*}} "-fsycl-is-host" {{.*}}"-fsycl-esimd-force-stateless-mem" -" -/// Verify that stateless memory accesses mapping is not enforced by default +// Case1: Check that the enforcing is turned on by default. +// Actually, only sycl-post-link gets an additional flag in this case. // RUN: %clang -### -fsycl %s 2>&1 | FileCheck -check-prefix=CHECK-DEFAULT %s -// CHECK-DEFAULT-NOT: clang{{.*}} "-fsycl-esimd-force-stateless-mem" +// CHECK-DEFAULT-NOT: clang{{.*}} "sycl-esimd-force-stateless-mem" // CHECK-DEFAULT-NOT: sycl-post-link{{.*}} "-lower-esimd-force-stateless-mem" +// CHECK-DEFAULT-NOT: clang{{.*}} "-fsycl-is-host" {{.*}}"sycl-esimd-force-stateless-mem" + +// Case2: Check that -fno-sycl-esimd-force-stateless-mem is handled correctly - +// i.e. sycl-post-link gets nothing and clang gets corresponding -fno... option. +// RUN: %clang -### -fsycl -fno-sycl-esimd-force-stateless-mem %s 2>&1 | FileCheck -check-prefix=CHECK-NEG %s +// CHECK-NEG: clang{{.*}} "-fno-sycl-esimd-force-stateless-mem" +// CHECK-NEG: sycl-post-link{{.*}} "-lower-esimd-force-stateless-mem=false" +// CHECK-NEG-NOT: clang{{.*}} "-fsycl-is-host" {{.*}}"sycl-esimd-force-stateless-mem" diff --git a/clang/test/Preprocessor/sycl-esimd-force-stateless-mem.cpp b/clang/test/Preprocessor/sycl-esimd-force-stateless-mem.cpp index 61dcfb25f6e7d..007bf29e718d4 100644 --- a/clang/test/Preprocessor/sycl-esimd-force-stateless-mem.cpp +++ b/clang/test/Preprocessor/sycl-esimd-force-stateless-mem.cpp @@ -1,11 +1,13 @@ /// This test checks that the macro __ESIMD_FORCE_STATELESS_MEM is automatically -/// defined only if the option -fsycl-esimd-force-stateless-mem is used. +/// defined by default with -fsycl, -fsycl-is-device or -fsycl-is-host. // RUN: %clang_cc1 %s -fsycl-is-device -fsycl-esimd-force-stateless-mem -E -dM | FileCheck --check-prefix=CHECK-OPT %s +// RUN: %clang_cc1 %s -fsycl-is-device -E -dM | FileCheck --check-prefix=CHECK-OPT %s +// RUN: %clang_cc1 %s -fsycl-is-host -E -dM | FileCheck --check-prefix=CHECK-OPT %s // RUN: %clang_cc1 %s -E -dM | FileCheck --check-prefix=CHECK-NOOPT %s -// RUN: %clang_cc1 %s -fsycl-is-device -E -dM | FileCheck --check-prefix=CHECK-NOOPT %s -// RUN: %clang_cc1 %s -fsycl-is-host -E -dM | FileCheck --check-prefix=CHECK-NOOPT %s +// RUN: %clang_cc1 %s -fsycl-is-device -fno-sycl-esimd-force-stateless-mem -E -dM | FileCheck --check-prefix=CHECK-NOOPT %s +// RUN: %clang_cc1 %s -fsycl-is-host -fno-sycl-esimd-force-stateless-mem -E -dM | FileCheck --check-prefix=CHECK-NOOPT %s // CHECK-OPT:#define __ESIMD_FORCE_STATELESS_MEM 1 // CHECK-NOOPT-NOT:#define __ESIMD_FORCE_STATELESS_MEM 1 diff --git a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp index 4f87d9e70192d..cce04e43f5965 100644 --- a/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp +++ b/llvm/lib/SYCLLowerIR/ESIMD/LowerESIMD.cpp @@ -57,7 +57,7 @@ using namespace llvm::esimd; cl::opt ForceStatelessMem( "lower-esimd-force-stateless-mem", llvm::cl::Optional, llvm::cl::Hidden, llvm::cl::desc("Use stateless API for accessor based API."), - llvm::cl::init(false)); + llvm::cl::init(true)); namespace { SmallPtrSet collectGenXVolatileTypes(Module &); diff --git a/llvm/test/SYCLLowerIR/ESIMD/acc_ptr.ll b/llvm/test/SYCLLowerIR/ESIMD/acc_ptr.ll index c53a19f1af99f..650bd09cb451a 100644 --- a/llvm/test/SYCLLowerIR/ESIMD/acc_ptr.ll +++ b/llvm/test/SYCLLowerIR/ESIMD/acc_ptr.ll @@ -1,4 +1,5 @@ -; RUN: opt -passes=LowerESIMD -S < %s | FileCheck %s +; RUN: opt -passes=LowerESIMD -lower-esimd-force-stateless-mem=false -S < %s | FileCheck --check-prefix=CHECK-NOFORCE --check-prefix=CHECK %s +; RUN: opt -passes=LowerESIMD -lower-esimd-force-stateless-mem=true -S < %s | FileCheck --check-prefix=CHECK-FORCE --check-prefix=CHECK %s ; This test checks that LowerESIMD pass correctly interpretes the ; 'kernel_arg_accessor_ptr' metadata. Particularly, that it generates additional @@ -30,8 +31,14 @@ define weak_odr dso_local spir_kernel void @ESIMDKernel(i32 %_arg_, ptr addrspac ; CHECK: attributes #[[GENX_MAIN]] = { "CMGenxMain" "oclrt"="1" } ; CHECK: !genx.kernels = !{![[GENX_KERNELS:[0-9]+]]} -; CHECK: ![[GENX_KERNELS]] = !{ptr @ESIMDKernel, !"ESIMDKernel", ![[ARG_KINDS:[0-9]+]], i32 0, i32 0, ![[ARG_IO_KINDS:[0-9]+]], ![[ARG_DESCS:[0-9]+]], i32 0, i32 0} -; CHECK: ![[ARG_KINDS]] = !{i32 0, i32 2, i32 2, i32 0, i32 0} -; CHECK: ![[ARG_IO_KINDS]] = !{i32 0, i32 0, i32 0, i32 0, i32 0} -; CHECK: ![[ARG_DESCS]] = !{!"", !"buffer_t", !"buffer_t", !"", !"svmptr_t"} + +; CHECK-NOFORCE: ![[GENX_KERNELS]] = !{ptr @ESIMDKernel, !"ESIMDKernel", ![[ARG_KINDS:[0-9]+]], i32 0, i32 0, ![[ARG_IO_KINDS:[0-9]+]], ![[ARG_DESCS:[0-9]+]], i32 0, i32 0} +; CHECK-FORCE: ![[GENX_KERNELS]] = !{ptr @ESIMDKernel, !"ESIMDKernel", ![[ARG_KINDS:[0-9]+]], i32 0, i32 0, ![[ARG_KINDS]], ![[ARG_DESCS:[0-9]+]], i32 0, i32 0} + +; CHECK-NOFORCE: ![[ARG_KINDS]] = !{i32 0, i32 2, i32 2, i32 0, i32 0} +; CHECK-FORCE: ![[ARG_KINDS]] = !{i32 0, i32 0, i32 0, i32 0, i32 0} +; CHECK-NOFORCE: ![[ARG_IO_KINDS]] = !{i32 0, i32 0, i32 0, i32 0, i32 0} + +; CHECK-NOFORCE: ![[ARG_DESCS]] = !{!"", !"buffer_t", !"buffer_t", !"", !"svmptr_t"} +; CHECK-FORCE: ![[ARG_DESCS]] = !{!"", !"svmptr_t", !"svmptr_t", !"", !"svmptr_t"} diff --git a/sycl/test/esimd/lsc.cpp b/sycl/test/esimd/lsc.cpp index 3115d8b1abded..1d5093abfe83a 100644 --- a/sycl/test/esimd/lsc.cpp +++ b/sycl/test/esimd/lsc.cpp @@ -1,9 +1,9 @@ // RUN: %clangxx -O0 -fsycl -fno-sycl-esimd-force-stateless-mem -fsycl-device-only -Xclang -emit-llvm %s -o %t -// RUN: sycl-post-link -split-esimd -lower-esimd -O0 -S %t -o %t.table +// RUN: sycl-post-link -split-esimd -lower-esimd -lower-esimd-force-stateless-mem=false -O0 -S %t -o %t.table // RUN: FileCheck %s -input-file=%t_esimd_0.ll --check-prefixes=CHECK,CHECK-STATEFUL // RUN: %clangxx -O0 -fsycl -fsycl-esimd-force-stateless-mem -fsycl-device-only -Xclang -emit-llvm %s -o %t -// RUN: sycl-post-link -split-esimd -lower-esimd -lower-esimd-force-stateless-mem -O0 -S %t -o %t.table +// RUN: sycl-post-link -split-esimd -lower-esimd -O0 -S %t -o %t.table // RUN: FileCheck %s -input-file=%t_esimd_0.ll --check-prefixes=CHECK,CHECK-STATELESS // Checks ESIMD intrinsic translation. From 5d41fad8d9b458dffccdf9944be62cb852cce827 Mon Sep 17 00:00:00 2001 From: Vyacheslav N Klochkov Date: Mon, 7 Aug 2023 08:43:39 -0700 Subject: [PATCH 2/6] Support 64-bit offset in scalar_load() and scalar_store() Signed-off-by: Vyacheslav N Klochkov --- sycl/include/sycl/ext/intel/esimd/memory.hpp | 21 ++++++++++++++++---- 1 file changed, 17 insertions(+), 4 deletions(-) diff --git a/sycl/include/sycl/ext/intel/esimd/memory.hpp b/sycl/include/sycl/ext/intel/esimd/memory.hpp index 07b9961b8d68a..fdf70937cf28a 100644 --- a/sycl/include/sycl/ext/intel/esimd/memory.hpp +++ b/sycl/include/sycl/ext/intel/esimd/memory.hpp @@ -620,9 +620,15 @@ scatter(AccessorTy acc, simd offsets, simd vals, /// @return The loaded value. /// template -__ESIMD_API T scalar_load(AccessorTy acc, uint32_t offset) { +__ESIMD_API T scalar_load(AccessorTy acc, +#ifdef __ESIMD_FORCE_STATELESS_MEM + uint64_t offset +#else + uint32_t offset +#endif +) { const simd Res = - gather(acc, simd(offset)); + gather(acc, simd(offset)); return Res[0]; } @@ -634,8 +640,15 @@ __ESIMD_API T scalar_load(AccessorTy acc, uint32_t offset) { /// @param val The stored value. /// template -__ESIMD_API void scalar_store(AccessorTy acc, uint32_t offset, T val) { - scatter(acc, simd(offset), simd(val)); +__ESIMD_API void scalar_store(AccessorTy acc, +#ifdef __ESIMD_FORCE_STATELESS_MEM + uint64_t offset, +#else + uint32_t offset, +#endif + T val) { + scatter(acc, simd(offset), + simd(val)); } /// @anchor usm_gather_rgba From 6559fad2a9a3f67e69a7d048e15e2bf2b1efcdbf Mon Sep 17 00:00:00 2001 From: Vyacheslav N Klochkov Date: Mon, 14 Aug 2023 12:17:55 -0700 Subject: [PATCH 3/6] Fix simd_view_copy_move_assign.cpp testt Signed-off-by: Vyacheslav N Klochkov --- .../ESIMD/api/simd_view_copy_move_assign.cpp | 33 +++++++++---------- 1 file changed, 16 insertions(+), 17 deletions(-) diff --git a/sycl/test-e2e/ESIMD/api/simd_view_copy_move_assign.cpp b/sycl/test-e2e/ESIMD/api/simd_view_copy_move_assign.cpp index c80ec19b397b2..c02cca27430a3 100644 --- a/sycl/test-e2e/ESIMD/api/simd_view_copy_move_assign.cpp +++ b/sycl/test-e2e/ESIMD/api/simd_view_copy_move_assign.cpp @@ -24,34 +24,34 @@ using namespace sycl::ext::intel::esimd; template bool test(queue q, std::string str, F funcUnderTest) { std::cout << "Testing " << str << ", VL = " << VL << " ...\n"; - T A[VL]; - T B[VL]; + size_t Size = 4 * VL; + T A[Size]; + T B[Size]; constexpr unsigned HalfVL = VL > 1 ? (VL / 2) : 1; // The expected result gets the first half of values from B, - int gold[VL]; - for (int i = 0; i < VL; ++i) { + int gold[Size]; + for (int i = 0; i < Size; ++i) { A[i] = -i - 1; B[i] = i + 1; - gold[i] = ((VL > 1) && (i < HalfVL)) ? B[i] : A[i]; + gold[i] = ((VL > 1) && ((i % VL) < HalfVL)) ? B[i] : A[i]; } try { - buffer bufA(A, range<1>(VL)); - buffer bufB(B, range<1>(VL)); - range<1> glob_range{1}; + buffer BufA(A, range<1>(Size)); + buffer BufB(B, range<1>(Size)); q.submit([&](handler &cgh) { - auto PA = bufA.template get_access(cgh); - auto PB = bufB.template get_access(cgh); - cgh.parallel_for(glob_range, [=](id<1> i) SYCL_ESIMD_KERNEL { + auto PA = BufA.template get_access(cgh); + auto PB = BufB.template get_access(cgh); + cgh.parallel_for(range<1>{Size / VL}, [=](id<1> i) SYCL_ESIMD_KERNEL { using namespace sycl::ext::intel::esimd; unsigned int offset = i * VL * sizeof(T); simd va; simd vb; if constexpr (VL == 1) { - va[0] = scalar_load(PA, 0); - vb[0] = scalar_load(PB, 0); + va[0] = scalar_load(PA, offset); + vb[0] = scalar_load(PB, offset); } else { va.copy_from(PA, offset); vb.copy_from(PB, offset); @@ -62,7 +62,7 @@ bool test(queue q, std::string str, F funcUnderTest) { funcUnderTest(va_view, vb_view); if constexpr (VL == 1) { - scalar_store(PB, 0, (T)va[0]); + scalar_store(PB, offset, (T)va[0]); } else { va.copy_to(PA, offset); } @@ -74,7 +74,7 @@ bool test(queue q, std::string str, F funcUnderTest) { } int err_cnt = 0; - for (unsigned i = 0; i < VL; ++i) { + for (unsigned i = 0; i < Size; ++i) { if (A[i] != gold[i]) { err_cnt++; std::cout << "failed at index " << i << ": " << A[i] << " != " << gold[i] @@ -157,8 +157,7 @@ template bool testT(queue &q) { int main(void) { queue q(esimd_test::ESIMDSelector, esimd_test::createExceptionHandler()); auto dev = q.get_device(); - std::cout << "Running on " << dev.get_info() - << "\n"; + esimd_test::printTestLabel(q); bool passed = true; passed &= testT(q); passed &= testT(q); From c67e786a98449b0e72c86b4e2ab5920240de5500 Mon Sep 17 00:00:00 2001 From: Vyacheslav N Klochkov Date: Tue, 15 Aug 2023 10:35:57 -0700 Subject: [PATCH 4/6] Fix the test intrins_trans.cpp Signed-off-by: Vyacheslav N Klochkov --- sycl/test/esimd/intrins_trans.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/test/esimd/intrins_trans.cpp b/sycl/test/esimd/intrins_trans.cpp index b761aec084645..89e26b04db1b2 100644 --- a/sycl/test/esimd/intrins_trans.cpp +++ b/sycl/test/esimd/intrins_trans.cpp @@ -1,9 +1,9 @@ // RUN: %clangxx -O0 -fsycl -fsycl-device-only -fno-sycl-esimd-force-stateless-mem -Xclang -emit-llvm %s -o %t -// RUN: sycl-post-link -split-esimd -lower-esimd -O0 -S %t -o %t.table +// RUN: sycl-post-link -split-esimd -lower-esimd -lower-esimd-force-stateless-mem=false -O0 -S %t -o %t.table // RUN: FileCheck %s -input-file=%t_esimd_0.ll --check-prefixes=CHECK,CHECK-STATEFUL // RUN: %clangxx -O0 -fsycl -fsycl-device-only -fsycl-esimd-force-stateless-mem -Xclang -emit-llvm %s -o %t -// RUN: sycl-post-link -split-esimd -lower-esimd -lower-esimd-force-stateless-mem -O0 -S %t -o %t.table +// RUN: sycl-post-link -split-esimd -lower-esimd -lower-esimd-force-stateless-mem=true -O0 -S %t -o %t.table // RUN: FileCheck %s -input-file=%t_esimd_0.ll --check-prefixes=CHECK,CHECK-STATELESS // Checks ESIMD intrinsic translation with opaque pointers. From 2ae07866fb95217f65e8e328e0519aacb791707b Mon Sep 17 00:00:00 2001 From: "Klochkov, Vyacheslav N" Date: Thu, 4 Jan 2024 11:55:25 -0800 Subject: [PATCH 5/6] Address reviewer's comments: removed duplicated code from Clang.cpp Signed-off-by: Klochkov, Vyacheslav N --- clang/lib/Driver/ToolChains/Clang.cpp | 4 ---- 1 file changed, 4 deletions(-) diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index bdc27d49bbadf..babf605a03ac4 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -5450,10 +5450,6 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA, if (Args.hasArg(options::OPT_fno_sycl_esimd_build_host_code)) CmdArgs.push_back("-fno-sycl-esimd-build-host-code"); } - if (IsSYCL && - !Args.hasFlag(options::OPT_fsycl_esimd_force_stateless_mem, - options::OPT_fno_sycl_esimd_force_stateless_mem, true)) - CmdArgs.push_back("-fno-sycl-esimd-force-stateless-mem"); const auto DeviceTraitsMacrosArgs = D.getDeviceTraitsMacrosArgs(); for (const auto &Arg : DeviceTraitsMacrosArgs) { From 6e20adf285a2fde3b617bceee0259373a7f68289 Mon Sep 17 00:00:00 2001 From: Vyacheslav Klochkov Date: Thu, 15 Feb 2024 14:08:09 -0600 Subject: [PATCH 6/6] Update clang/test/Preprocessor/sycl-esimd-force-stateless-mem.cpp Co-authored-by: premanandrao --- clang/test/Preprocessor/sycl-esimd-force-stateless-mem.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/test/Preprocessor/sycl-esimd-force-stateless-mem.cpp b/clang/test/Preprocessor/sycl-esimd-force-stateless-mem.cpp index 007bf29e718d4..d8db457c50762 100644 --- a/clang/test/Preprocessor/sycl-esimd-force-stateless-mem.cpp +++ b/clang/test/Preprocessor/sycl-esimd-force-stateless-mem.cpp @@ -1,5 +1,5 @@ /// This test checks that the macro __ESIMD_FORCE_STATELESS_MEM is automatically -/// defined by default with -fsycl, -fsycl-is-device or -fsycl-is-host. +/// defined by default with -fsycl-is-device or -fsycl-is-host. // RUN: %clang_cc1 %s -fsycl-is-device -fsycl-esimd-force-stateless-mem -E -dM | FileCheck --check-prefix=CHECK-OPT %s // RUN: %clang_cc1 %s -fsycl-is-device -E -dM | FileCheck --check-prefix=CHECK-OPT %s