[SYCL] Add a module flag for device compilations (#13880)

Add a module flag "-sycl-device" for SYCL device compilations. 
The flag is not added for host compilations.  This helps optimization 
reports identify target device compilations.

Author: premanandrao

clang/lib/CodeGen/CodeGenModule.cpp

@@ -1308,8 +1308,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)

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 {{.*}}}
 //.

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"