Remove unnecessary changes

Signed-off-by: Mikhail Lychkov <[email protected]>

Author: mlychkov

clang/test/CodeGenSYCL/Inputs/sycl.hpp

@@ -1,18 +1,7 @@
 #pragma once
 
-typedef __UINT8_TYPE__ uint8_t;
-typedef __SIZE_TYPE__ size_t;
-
 #define ATTR_SYCL_KERNEL __attribute__((sycl_kernel))
 
-#ifndef __SYCL_ALWAYS_INLINE
-#if __has_attribute(always_inline)
-#define __SYCL_ALWAYS_INLINE __attribute__((always_inline))
-#else
-#define __SYCL_ALWAYS_INLINE
-#endif
-#endif // __SYCL_ALWAYS_INLINE
-
 // Dummy runtime classes to model SYCL API.
 namespace cl {
 namespace sycl {
@@ -505,19 +494,5 @@ class image {
   }
 };
 
-extern "C" SYCL_EXTERNAL __attribute__((opencl_local)) uint8_t *
-__sycl_allocateLocalMemory(size_t Size, size_t Alignment);
-
-template <typename T>
-__attribute__((opencl_local)) T *
-    __SYCL_ALWAYS_INLINE
-    group_local_memory() {
-#ifdef __SYCL_DEVICE_ONLY__
-  __attribute__((opencl_local)) uint8_t *AllocatedMem =
-      __sycl_allocateLocalMemory(sizeof(T), alignof(T));
-  return (__attribute__((opencl_local)) T *)AllocatedMem;
-#endif
-}
-
 } // namespace sycl
 } // namespace cl

clang/test/CodeGenSYCL/group-local-memory.cpp

@@ -1,44 +1,8 @@
-// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -S -emit-llvm %s -o - | FileCheck --check-prefixes=CHECK,CHECK-UNOPT %s
-// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -S -emit-llvm %s -o - | FileCheck %s
-
-// CHECK: [[WGLOCALMEM_1:@WGLocalMem.*]] = internal addrspace(3) global [8 x i8] undef, align 8
-// CHECK: [[WGLOCALMEM_2:@WGLocalMem.*]] = internal addrspace(3) global [4 x i8] undef, align 4
-// CHECK: [[WGLOCALMEM_3:@WGLocalMem.*]] = internal addrspace(3) global [128 x i8] undef, align 4
-
-// CHECK-NOT: __sycl_allocateLocalMemory
-
-#include "Inputs/sycl.hpp"
-
-constexpr size_t WgSize = 32;
-constexpr size_t WgCount = 4;
-constexpr size_t Size = WgSize * WgCount;
-
-class KernelA;
-class KernelB;
-
-using namespace cl::sycl;
-
-int main() {
-  queue Q;
-  {
-    Q.submit([&](handler &cgh) {
-      cgh.parallel_for<KernelA>(
-          range<1>(Size), [=](item<1> Item) {
-            auto *Ptr1 = group_local_memory<long>();
-            // CHECK-UNOPT: i8 addrspace(3)* getelementptr inbounds ([8 x i8], [8 x i8] addrspace(3)* [[WGLOCALMEM_1]], i32 0, i32 0)
-            auto *Ptr2 = group_local_memory<float>();
-            // CHECK-UNOPT: i8 addrspace(3)* getelementptr inbounds ([4 x i8], [4 x i8] addrspace(3)* [[WGLOCALMEM_2]], i32 0, i32 0)
-          });
-    });
-  }
-
-  {
-    Q.submit([&](handler &cgh) {
-      cgh.parallel_for<KernelB>(
-          range<1>(Size), [=](item<1> Item) {
-            auto *Ptr3 = group_local_memory<int[WgSize]>();
-            // CHECK-UNOPT: i8 addrspace(3)* getelementptr inbounds ([128 x i8], [128 x i8] addrspace(3)* [[WGLOCALMEM_3]], i32 0, i32 0)
-          });
-    });
-  }
-}
+// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown-sycldevice \
+// RUN:     -S -emit-llvm -mllvm -debug-pass=Structure -disable-llvm-passes \
+// RUN:     -o - %s 2>&1 | FileCheck %s
+// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown-sycldevice \
+// RUN:     -S -emit-llvm -mllvm -debug-pass=Structure -o - %s 2>&1 \
+// RUN:     | FileCheck %s
+
+// CHECK: Replace __sycl_allocateLocalMemory with allocation of memory in local address space

clang/test/CodeGenSYCL/kernel-by-reference.cpp

@@ -15,15 +15,15 @@ int simple_add(int i) {
 int main() {
   queue q;
 #if defined(SYCL2020)
-  // expected-warning@Inputs/sycl.hpp:{{[0-9]+}} {{Passing kernel functions by value is deprecated in SYCL 2020}}
+  // expected-warning@Inputs/sycl.hpp:301 {{Passing kernel functions by value is deprecated in SYCL 2020}}
   // expected-note@+3 {{in instantiation of function template specialization}}
 #endif
   q.submit([&](handler &h) {
     h.single_task_2017<class sycl2017>([]() { simple_add(10); });
   });
 
 #if defined(SYCL2017)
-  // expected-warning@Inputs/sycl.hpp:{{[0-9]+}} {{Passing of kernel functions by reference is a SYCL 2020 extension}}
+  // expected-warning@Inputs/sycl.hpp:296 {{Passing of kernel functions by reference is a SYCL 2020 extension}}
   // expected-note@+3 {{in instantiation of function template specialization}}
 #endif
   q.submit([&](handler &h) {