================
@@ -0,0 +1,48 @@
+/// Check that the x86 intrinsic headers parse and compile for a SPIR-V device
+/// when the auxiliary host target is x86-64 MSVC, which is what the sse/sse2
+/// device features derived from that host are for.
+
+// RUN: %clang_cc1 -triple spirv64-unknown-unknown -aux-triple
x86_64-pc-windows-msvc \
+// RUN: -fsycl-is-device -ffreestanding -emit-llvm -o - %s | FileCheck %s
+
+/// With sse/sse2 disabled, the always_inline intrinsics cannot be inlined into
+/// device code; this is the failure the derived features prevent.
+/// Codegen stops at the first failing function, so only add() is diagnosed.
+// RUN: %clang_cc1 -triple spirv64-unknown-unknown -aux-triple
x86_64-pc-windows-msvc \
+// RUN: -fsycl-is-device -ffreestanding -target-feature -sse -target-feature
-sse2 \
+// RUN: -emit-llvm -verify=no-sse -o - %s
+
+/// Intrinsics that lower to x86 builtins are rejected for the device.
+// RUN: %clang_cc1 -triple spirv64-unknown-unknown -aux-triple
x86_64-pc-windows-msvc \
+// RUN: -fsycl-is-device -ffreestanding -DX86_BUILTIN -emit-llvm
-verify=x86-builtin \
+// RUN: -o - %s
+
+#include <immintrin.h>
+
+// CHECK-LABEL: define {{.*}}spir_func noundef float @_Z3addff
+[[clang::sycl_external]] float add(float x, float y) {
+ // no-sse-error@+1 {{always_inline function '_mm_set1_ps' requires target
feature 'sse', but would be inlined into function 'add' that is compiled
without support for 'sse'}}
+ __m128 a = _mm_set1_ps(x);
+ // no-sse-error@+1 {{always_inline function '_mm_set1_ps' requires target
feature 'sse', but would be inlined into function 'add' that is compiled
without support for 'sse'}}
+ __m128 b = _mm_set1_ps(y);
+ // CHECK: fadd <4 x float>
----------------
schittir wrote:
> I think DiagIfDeviceCode() should suppress the error if you don't actually
> use the relevant math functions? At least, that's how I remember it's
> supposed to work.
I think it should. Maybe something like this. Needs work. Will follow-up.
Sema::DiagnoseUseOfDecl (SemaExpr.cpp), which runs on every call:
```
if (getLangOpts().SYCLIsDevice)
if (auto *FD = dyn_cast<FunctionDecl>(D))
if (const auto *TA = FD->getAttr<TargetAttr>())
if (/* TA requests features the device target doesn't support */)
SYCL().DiagIfDeviceCode(Loc, diag::err_sycl_target_attr_callee) << FD;
```
https://github.com/llvm/llvm-project/pull/227665
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
