From 88fe42534c5e0c91f2e05fb1651b4d187d5a5118 Mon Sep 17 00:00:00 2001
From: Mariya Podchishchaeva <mariya.podchishchaeva@intel.com>
Date: Tue, 11 Aug 2020 15:28:43 +0300
Subject: [PATCH] [SYCL] Add a test for wrapped pointer kernel args

Before structs decomposition USM pointers that are wrapped by structs
were passsed with generic address space. It caused performance issues.
After structs decomposition all pointers are passed with global address
space, this test is added to save this behaviour.
---
 .../test/CodeGenSYCL/pointers-in-structs.cpp  | 36 +++++++++++++++++++
 1 file changed, 36 insertions(+)
 create mode 100644 clang/test/CodeGenSYCL/pointers-in-structs.cpp

diff --git a/clang/test/CodeGenSYCL/pointers-in-structs.cpp b/clang/test/CodeGenSYCL/pointers-in-structs.cpp
new file mode 100644
index 0000000000000..ca6251801eb69
--- /dev/null
+++ b/clang/test/CodeGenSYCL/pointers-in-structs.cpp
@@ -0,0 +1,36 @@
+// RUN: %clang_cc1 -fsycl -fsycl-is-device -I %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
+
+// This test checks that compiler generates correct address spaces for pointer
+// kernel arguments that are wrapped by struct.
+
+#include <sycl.hpp>
+
+struct A {
+  float *F;
+};
+
+struct B {
+  int *F1;
+  float *F2;
+  A F3;
+};
+
+int main() {
+  B Obj;
+  cl::sycl::kernel_single_task<class structs>(
+      [=]() {
+        (void)Obj;
+      });
+  float A = 1;
+  float *Ptr = &A;
+  auto Lambda = [=]() {
+    *Ptr += 1;
+  };
+  cl::sycl::kernel_single_task<class lambdas>([=]() {
+    Lambda();
+  });
+  return 0;
+}
+
+// CHECK: define spir_kernel void @{{.*}}structs{{.*}}(i32 addrspace(1)* %_arg_F1, float addrspace(1)* %_arg_F2, float addrspace(1)* %_arg_F)
+// CHECK: define spir_kernel void @{{.*}}lambdas{{.*}}(float addrspace(1)* %_arg_)