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..b354702ca47de --- /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 -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 +// 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; +}