diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 7f47a8af721e8..00f906d5ea6b1 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -1296,8 +1296,10 @@ void CodeGenModule::Release() { getTarget().getTargetOpts().NVVMCudaPrecSqrt); } - if (LangOpts.SYCLIsDevice && LangOpts.SYCLIsNativeCPU) { - getModule().addModuleFlag(llvm::Module::Error, "is-native-cpu", 1); + if (LangOpts.SYCLIsDevice) { + getModule().addModuleFlag(llvm::Module::Error, "sycl-device", 1); + if (LangOpts.SYCLIsNativeCPU) + getModule().addModuleFlag(llvm::Module::Error, "is-native-cpu", 1); } if (LangOpts.EHAsynch) diff --git a/clang/test/CodeGenSYCL/function-attrs.cpp b/clang/test/CodeGenSYCL/function-attrs.cpp index 3f2b2467374f2..3a82764e46cae 100644 --- a/clang/test/CodeGenSYCL/function-attrs.cpp +++ b/clang/test/CodeGenSYCL/function-attrs.cpp @@ -5,7 +5,7 @@ int foo(); // CHECK-LABEL: define linkonce_odr spir_func void @_Z3barv( -// CHECK-SAME: ) #[[ATTR2:[0-9]+]] !srcloc !5 { +// CHECK-SAME: ) #[[ATTR2:[0-9]+]] !srcloc ![[MD1:[0-9]+]] { // CHECK-NEXT: entry: // CHECK-NEXT: [[A:%.*]] = alloca i32, align 4 // CHECK-NEXT: [[A_ASCAST:%.*]] = addrspacecast ptr [[A]] to ptr addrspace(4) @@ -18,7 +18,7 @@ void bar() { } // CHECK-LABEL: define linkonce_odr spir_func noundef i32 @_Z3foov( -// CHECK-SAME: ) #[[ATTR2]] !srcloc !6 { +// CHECK-SAME: ) #[[ATTR2]] !srcloc ![[MD2:[0-9]+]] { // CHECK-NEXT: entry: // CHECK-NEXT: [[RETVAL:%.*]] = alloca i32, align 4 // CHECK-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr [[RETVAL]] to ptr addrspace(4) @@ -43,11 +43,11 @@ int main() { // CHECK: attributes #2 = { convergent mustprogress noinline norecurse nounwind optnone "no-trapping-math"="true" "stack-protector-buffer-size"="8" } // CHECK: attributes #3 = { convergent nounwind } //. -// CHECK: !0 = !{i32 1, !"wchar_size", i32 4} -// CHECK: !1 = !{i32 1, i32 2} -// CHECK: !2 = !{i32 4, i32 100000} -// CHECK: !3 = !{i32 {{.*}}} -// CHECK: !4 = !{} -// CHECK: !5 = !{i32 {{.*}}} -// CHECK: !6 = !{i32 {{.*}}} +// CHECK: !{{[0-9]+}} = !{i32 1, !"wchar_size", i32 4} +// CHECK: !{{[0-9]+}} = !{i32 1, i32 2} +// CHECK: !{{[0-9]+}} = !{i32 4, i32 100000} +// CHECK: !{{[0-9]+}} = !{i32 {{.*}}} +// CHECK: !{{[0-9]+}} = !{} +// CHECK: ![[MD1]] = !{i32 {{.*}}} +// CHECK: ![[MD2]] = !{i32 {{.*}}} //. diff --git a/clang/test/CodeGenSYCL/sycl-device-module-flag.cpp b/clang/test/CodeGenSYCL/sycl-device-module-flag.cpp new file mode 100644 index 0000000000000..b123e96266fe6 --- /dev/null +++ b/clang/test/CodeGenSYCL/sycl-device-module-flag.cpp @@ -0,0 +1,16 @@ +// RUN: %clang_cc1 -fsycl-is-device -emit-llvm -internal-isystem %S/Inputs -o - %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-host -emit-llvm -internal-isystem %S/Inputs -o - %s | FileCheck %s --check-prefixes=HOST + +// This test checks if the sycl-device module flag is created for device +// compilations and not for host compilations. +#include "sycl.hpp" + +void foo() { + sycl::handler h; + h.single_task([]() {}); +} + +// Check for the presence of sycl-device module flag in device +// compilations and its absence in host compilations. +// CHECK: !{{[0-9]*}} = !{i32 1, !"sycl-device", i32 1} +// HOST-NOT: !"sycl-device"