From 92f7701aeef71942dce22a132b2d0dec54e8c447 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Thu, 16 Mar 2023 10:03:53 -0700 Subject: [PATCH 1/2] [SYCL] Reenable fp64 aspect runtime check This commit reenables the fp64 aspect check when ensuring compatability of a given device image, following the relaxation of aspect propagation for the aspect. Signed-off-by: Larsen, Steffen --- .../program_manager/program_manager.cpp | 4 ++-- .../relaxed_fp64_propagation.cpp | 24 +++++++++++++++++++ 2 files changed, 26 insertions(+), 2 deletions(-) create mode 100644 sycl/test/optional_kernel_features/relaxed_fp64_propagation.cpp diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index 86800d49711f8..75065098daa94 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -567,7 +567,7 @@ RT::PiProgram ProgramManager::getBuiltPIProgram( while (AIt != AEnd) { auto Aspect = static_cast(*AIt); // Strict check for fp64 is disabled temporarily to avoid confusion. - if (Aspect != aspect::fp64 && !Dev->has(Aspect)) + if (!Dev->has(Aspect)) throw sycl::exception(errc::kernel_not_supported, "Required aspect " + getAspectNameStr(Aspect) + " is not supported on the device"); @@ -2309,7 +2309,7 @@ bool doesDevSupportDeviceRequirements(const device &Dev, while (!Aspects.empty()) { aspect Aspect = Aspects.consume(); // Strict check for fp64 is disabled temporarily to avoid confusion. - if (Aspect != aspect::fp64 && !Dev.has(Aspect)) + if (!Dev.has(Aspect)) return false; } } diff --git a/sycl/test/optional_kernel_features/relaxed_fp64_propagation.cpp b/sycl/test/optional_kernel_features/relaxed_fp64_propagation.cpp new file mode 100644 index 0000000000000..a5cb7f5ecde85 --- /dev/null +++ b/sycl/test/optional_kernel_features/relaxed_fp64_propagation.cpp @@ -0,0 +1,24 @@ +// RUN: %clangxx %s -S -o %t_opt.ll -fsycl-device-only -Xclang -verify -Xclang -verify-ignore-unexpected=note +// RUN: FileCheck %s --input-file %t_opt.ll --check-prefix=CHECK-OPT +// RUN: %clangxx %s -S -O0 -o %t_noopt.ll -fsycl-device-only -Xclang -verify -Xclang -verify-ignore-unexpected=note +// RUN: FileCheck %s --input-file %t_noopt.ll --check-prefix=CHECK-NOOPT + +// Tests that an optimization that removes the use of double still produces a +// warning. + +// CHECK-OPT-NOT: double +// CHECK-NOOPT: double + +#include + +int main() { + sycl::queue Q; + // expected-warning-re@+1 {{function '{{.*}}' uses aspect 'fp64' not listed in its 'sycl::device_has' attribute}} + Q.single_task([=]() [[sycl::device_has()]] { + // Double will be optimized out as LoweredFloat can be set directly to a + // lowered value. + double Double = 3.14; + volatile float LoweredFloat = Double; + }); + return 0; +} From 8733950734ead0d3beeb07d102b2eb0286b9e099 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Wed, 22 Mar 2023 12:25:39 -0700 Subject: [PATCH 2/2] Switch to no-sycl-early-optimizations to avoid unrecognized -O0 Signed-off-by: Larsen, Steffen --- sycl/test/optional_kernel_features/relaxed_fp64_propagation.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test/optional_kernel_features/relaxed_fp64_propagation.cpp b/sycl/test/optional_kernel_features/relaxed_fp64_propagation.cpp index a5cb7f5ecde85..b354702ca47de 100644 --- a/sycl/test/optional_kernel_features/relaxed_fp64_propagation.cpp +++ b/sycl/test/optional_kernel_features/relaxed_fp64_propagation.cpp @@ -1,6 +1,6 @@ // RUN: %clangxx %s -S -o %t_opt.ll -fsycl-device-only -Xclang -verify -Xclang -verify-ignore-unexpected=note // RUN: FileCheck %s --input-file %t_opt.ll --check-prefix=CHECK-OPT -// RUN: %clangxx %s -S -O0 -o %t_noopt.ll -fsycl-device-only -Xclang -verify -Xclang -verify-ignore-unexpected=note +// RUN: %clangxx %s -S -fno-sycl-early-optimizations -o %t_noopt.ll -fsycl-device-only -Xclang -verify -Xclang -verify-ignore-unexpected=note // RUN: FileCheck %s --input-file %t_noopt.ll --check-prefix=CHECK-NOOPT // Tests that an optimization that removes the use of double still produces a