Skip to content

Commit 0106136

Browse files
MrSidimsMrSidims
authored
[SYCL] Pass foffload-fp32-prec-[div/sqrt] options to device's BE (#16107)
Signed-off-by: Sidorov, Dmitry <[email protected]>
1 parent c0c91aa commit 0106136

File tree

7 files changed

+118
-1
lines changed

7 files changed

+118
-1
lines changed

‎clang/lib/Driver/ToolChains/SYCL.cpp

Lines changed: 10 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -2061,9 +2061,18 @@ void SYCLToolChain::AddImpliedTargetArgs(const llvm::Triple &Triple,
20612061
if (Args.hasFlag(options::OPT_ftarget_export_symbols,
20622062
options::OPT_fno_target_export_symbols, false))
20632063
BeArgs.push_back("-library-compilation");
2064-
} else if (IsJIT)
2064+
// -foffload-fp32-prec-[sqrt/div]
2065+
if (Args.hasArg(options::OPT_foffload_fp32_prec_div) ||
2066+
Args.hasArg(options::OPT_foffload_fp32_prec_sqrt))
2067+
BeArgs.push_back("-ze-fp32-correctly-rounded-divide-sqrt");
2068+
} else if (IsJIT) {
20652069
// -ftarget-compile-fast JIT
20662070
Args.AddLastArg(BeArgs, options::OPT_ftarget_compile_fast);
2071+
// -foffload-fp32-prec-div JIT
2072+
Args.AddLastArg(BeArgs, options::OPT_foffload_fp32_prec_div);
2073+
// -foffload-fp32-prec-sqrt JIT
2074+
Args.AddLastArg(BeArgs, options::OPT_foffload_fp32_prec_sqrt);
2075+
}
20672076
if (IsGen) {
20682077
for (auto [DeviceName, BackendArgStr] : PerDeviceArgs) {
20692078
CmdArgs.push_back("-device_options");

‎clang/test/Driver/sycl-foffload-fp32-prec-div-old-model.cpp

Lines changed: 13 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,13 @@
1+
// Test SYCL -foffload-fp32-prec-div
2+
3+
// RUN: %clang -### -fsycl --no-offload-new-driver \
4+
// RUN: -fsycl-targets=spir64_gen -foffload-fp32-prec-div %s 2>&1 \
5+
// RUN: | FileCheck -check-prefix=AOT %s
6+
7+
// RUN: %clang -### -fsycl --no-offload-new-driver \
8+
// RUN: -foffload-fp32-prec-div %s 2>&1 \
9+
// RUN: | FileCheck -check-prefix=JIT %s
10+
11+
// AOT: "-ze-fp32-correctly-rounded-divide-sqrt"
12+
13+
// JIT: clang-offload-wrapper{{.*}} "-compile-opts={{.*}}-foffload-fp32-prec-div"

‎clang/test/Driver/sycl-foffload-fp32-prec-div.cpp

Lines changed: 13 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,13 @@
1+
// Test SYCL -foffload-fp32-prec-div
2+
3+
// RUN: %clang -### -fsycl --offload-new-driver \
4+
// RUN: -fsycl-targets=spir64_gen -foffload-fp32-prec-div %s 2>&1 \
5+
// RUN: | FileCheck -check-prefix=AOT %s
6+
7+
// RUN: %clang -### -fsycl --offload-new-driver \
8+
// RUN: -foffload-fp32-prec-div %s 2>&1 \
9+
// RUN: | FileCheck -check-prefix=JIT %s
10+
11+
// AOT: clang-offload-packager{{.*}} "--image=file={{.*}}.bc,triple=spir64_gen-unknown-unknown,arch={{.*}},kind=sycl,compile-opts=-options -ze-fp32-correctly-rounded-divide-sqrt{{.*}}"
12+
13+
// JIT: clang-offload-packager{{.*}} "--image=file={{.*}}.bc,triple=spir64-unknown-unknown,arch={{.*}}compile-opts={{.*}}-foffload-fp32-prec-div"

‎clang/test/Driver/sycl-foffload-fp32-prec-sqrt-old-model.cpp

Lines changed: 13 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,13 @@
1+
// Test SYCL -foffload-fp32-prec-sqrt
2+
3+
// RUN: %clang -### -fsycl --no-offload-new-driver \
4+
// RUN: -fsycl-targets=spir64_gen -foffload-fp32-prec-sqrt %s 2>&1 \
5+
// RUN: | FileCheck -check-prefix=AOT %s
6+
7+
// RUN: %clang -### -fsycl --no-offload-new-driver \
8+
// RUN: -foffload-fp32-prec-sqrt %s 2>&1 \
9+
// RUN: | FileCheck -check-prefix=JIT %s
10+
11+
// AOT: "-ze-fp32-correctly-rounded-divide-sqrt"
12+
13+
// JIT: clang-offload-wrapper{{.*}} "-compile-opts={{.*}}-foffload-fp32-prec-sqrt"

‎clang/test/Driver/sycl-foffload-fp32-prec-sqrt.cpp

Lines changed: 13 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,13 @@
1+
// Test SYCL -foffload-fp32-prec-sqrt
2+
3+
// RUN: %clang -### -fsycl --offload-new-driver \
4+
// RUN: -fsycl-targets=spir64_gen -foffload-fp32-prec-sqrt %s 2>&1 \
5+
// RUN: | FileCheck -check-prefix=AOT %s
6+
7+
// RUN: %clang -### -fsycl --offload-new-driver \
8+
// RUN: -foffload-fp32-prec-sqrt %s 2>&1 \
9+
// RUN: | FileCheck -check-prefix=JIT %s
10+
11+
// AOT: clang-offload-packager{{.*}} "--image=file={{.*}}.bc,triple=spir64_gen-unknown-unknown,arch={{.*}},kind=sycl,compile-opts=-options -ze-fp32-correctly-rounded-divide-sqrt{{.*}}"
12+
13+
// JIT: clang-offload-packager{{.*}} "--image=file={{.*}}.bc,triple=spir64-unknown-unknown,arch={{.*}}compile-opts={{.*}}-foffload-fp32-prec-sqrt"

‎sycl/source/detail/program_manager/program_manager.cpp

Lines changed: 10 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -446,6 +446,16 @@ static void appendCompileOptionsFromImage(std::string &CompileOpts,
446446
CompileOpts = NewCompileOpts;
447447
OptPos = CompileOpts.find(TargetRegisterAllocMode);
448448
}
449+
constexpr std::string_view ReplaceOpts[] = {"-foffload-fp32-prec-div",
450+
"-foffload-fp32-prec-sqrt"};
451+
for (const std::string_view Opt : ReplaceOpts) {
452+
if (auto Pos = CompileOpts.find(Opt); Pos != std::string::npos) {
453+
const char *BackendOption = nullptr;
454+
PlatformImpl->getBackendOption(std::string(Opt).c_str(),
455+
&BackendOption);
456+
CompileOpts.replace(Pos, Opt.length(), BackendOption);
457+
}
458+
}
449459
}
450460
}
451461

‎sycl/test-e2e/KernelAndProgram/fp32-precise-fdiv.cpp

Lines changed: 46 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,46 @@
1+
// RUN: %{build} -Wno-error=unused-command-line-argument -foffload-fp32-prec-div -foffload-fp32-prec-sqrt -o %t_with.out
2+
// RUN: %{build} -Wno-error=unused-command-line-argument -foffload-fp32-prec-div -o %t_with_div.out
3+
// RUN: %{build} -Wno-error=unused-command-line-argument -foffload-fp32-prec-sqrt -o %t_with_sqrt.out
4+
// RUN: %{build} -o %t_without.out
5+
6+
// RUN: env SYCL_UR_TRACE=2 %{run} %t_with.out 2>&1 | FileCheck %if hip || cuda %{ --check-prefix=CHECK-WITHOUT %} %else %{ --check-prefix=CHECK-WITH %} %s
7+
// RUN: env SYCL_UR_TRACE=2 %{run} %t_with_div.out 2>&1 | FileCheck %if hip || cuda %{ --check-prefix=CHECK-WITHOUT %} %else %{ --check-prefix=CHECK-WITH %} %s
8+
// RUN: env SYCL_UR_TRACE=2 %{run} %t_with_sqrt.out 2>&1 | FileCheck %if hip || cuda %{ --check-prefix=CHECK-WITHOUT %} %else %{ --check-prefix=CHECK-WITH %} %s
9+
// RUN: env SYCL_UR_TRACE=2 %{run} %t_without.out 2>&1 | FileCheck --implicit-check-not=fp32-correctly-rounded-divide-sqrt %s
10+
11+
// CHECK-WITH: <--- urProgramBuild
12+
// CHECK-WITH-SAME: fp32-correctly-rounded-divide-sqrt
13+
14+
// CHECK-WITHOUT-NOT: <--- urProgramBuild{{.*}}fp32-correctly-rounded-divide-sqrt{{.*}} -> UR_RESULT_SUCCESS
15+
// CHECK-WITHOUT: <--- urProgramBuild{{.*}} -> UR_RESULT_SUCCESS
16+
17+
#include <sycl/detail/core.hpp>
18+
19+
int main() {
20+
sycl::buffer<size_t, 1> Buffer(4);
21+
22+
sycl::queue Queue;
23+
24+
sycl::range<1> NumOfWorkItems{Buffer.size()};
25+
26+
Queue.submit([&](sycl::handler &cgh) {
27+
sycl::accessor Accessor{Buffer, cgh, sycl::write_only};
28+
cgh.parallel_for<class FillBuffer>(NumOfWorkItems, [=](sycl::id<1> WIid) {
29+
Accessor[WIid] = WIid.get(0);
30+
});
31+
});
32+
33+
sycl::host_accessor HostAccessor{Buffer, sycl::read_only};
34+
35+
bool MismatchFound = false;
36+
for (size_t I = 0; I < Buffer.size(); ++I) {
37+
if (HostAccessor[I] != I) {
38+
std::cout << "The result is incorrect for element: " << I
39+
<< " , expected: " << I << " , got: " << HostAccessor[I]
40+
<< std::endl;
41+
MismatchFound = true;
42+
}
43+
}
44+
45+
return MismatchFound;
46+
}

0 commit comments

Comments
 (0)