From 77fe6e7faf64f1fc1b8c7ef17985783e6b5bc45c Mon Sep 17 00:00:00 2001
From: Steve Merritt <steve.merritt@intel.com>
Date: Wed, 26 Jun 2019 10:08:12 -0700
Subject: [PATCH] [SYCL] remove incorrect source correlation from kernel
 instructions

Clear incorrect source locations from several instructions in the IR for the
kernel.  This includes the call to __init(), the call to operator() and the
return statement.

Signed-off-by: Steve Merritt <steve.merritt@intel.com>
---
 clang/lib/Sema/SemaSYCL.cpp                   | 17 ++++++--
 .../CodeGenSYCL/debug-info-srcpos-kernel.cpp  | 39 +++++++++++--------
 2 files changed, 37 insertions(+), 19 deletions(-)

diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 516d37dc2b05d..1af8ea665f191 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -385,12 +385,23 @@ class KernelBodyTransform : public TreeTransform<KernelBodyTransform> {
       auto NewDecl = MappingPair.second;
       return DeclRefExpr::Create(
           SemaRef.getASTContext(), DRE->getQualifierLoc(),
-          DRE->getTemplateKeywordLoc(), NewDecl, false, DRE->getNameInfo(),
+          DRE->getTemplateKeywordLoc(), NewDecl, false,
+          DeclarationNameInfo(DRE->getNameInfo().getName(), SourceLocation(),
+                              DRE->getNameInfo().getInfo()),
           NewDecl->getType(), DRE->getValueKind());
     }
     return DRE;
   }
 
+  StmtResult RebuildCompoundStmt(SourceLocation LBraceLoc,
+                                 MultiStmtArg Statements,
+                                 SourceLocation RBraceLoc,
+                                 bool IsStmtExpr) {
+    // Build a new compound statement but clear the source locations.
+    return getSema().ActOnCompoundStmt(SourceLocation(), SourceLocation(),
+                                       Statements, IsStmtExpr);
+  }
+
 private:
   std::pair<DeclaratorDecl *, DeclaratorDecl *> MappingPair;
   Sema &SemaRef;
@@ -520,8 +531,8 @@ static CompoundStmt *CreateOpenCLKernelBody(Sema &S,
         auto ME = MemberExpr::Create(
             S.Context, SpecialObjME, false, SourceLocation(),
             NestedNameSpecifierLoc(), SourceLocation(), InitMethod, MethodDAP,
-            InitMethod->getNameInfo(), nullptr, InitMethod->getType(),
-            VK_LValue, OK_Ordinary);
+            DeclarationNameInfo(InitMethod->getDeclName(), SourceLocation()),
+            nullptr, InitMethod->getType(), VK_LValue, OK_Ordinary);
 
         // Not referenced -> not emitted
         S.MarkFunctionReferenced(SourceLocation(), InitMethod, true);
diff --git a/clang/test/CodeGenSYCL/debug-info-srcpos-kernel.cpp b/clang/test/CodeGenSYCL/debug-info-srcpos-kernel.cpp
index 51390aa8c7a15..cbc7fd1fd896c 100644
--- a/clang/test/CodeGenSYCL/debug-info-srcpos-kernel.cpp
+++ b/clang/test/CodeGenSYCL/debug-info-srcpos-kernel.cpp
@@ -1,31 +1,38 @@
-// RUN: DISABLE_INFER_AS=1 %clang --sycl %s -S -emit-llvm -g -o - | FileCheck %s --check-prefixes CHECK,CHECK-OLD
-// RUN: %clang --sycl %s -S -emit-llvm -g -o - | FileCheck %s --check-prefixes CHECK,CHECK-NEW
+// RUN: %clang --sycl %s -S -I %S/Inputs -emit-llvm -g -o - | FileCheck %s
 //
-// Verify the SYCL kernel routine is marked artificial.
+// Verify the SYCL kernel routine is marked artificial and has no source
+// correlation.
 //
-// Since it has no source correlation of its own, the SYCL kernel needs to be
-// marked artificial or it will inherit source correlation from the surrounding
-// code.
+// The SYCL kernel should have no source correlation of its own, so it needs
+// to be marked artificial or it will inherit source correlation from the
+// surrounding code.
 //
 
-template <typename Name, typename Func>
-__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
+#include <sycl.hpp>
+
+template <typename name, typename Func>
+__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
   kernelFunc();
 }
 
 int main() {
-  int value = 0;
-  int* captured = &value;
-  kernel_single_task<class kernel_function>([=]() {
-      *captured = 1;
-      });
+  cl::sycl::sampler Sampler;
+  kernel<class use_kernel_for_test>([=]() {
+    Sampler.use();
+  });
   return 0;
 }
 
-// CHECK-OLD: define{{.*}} spir_kernel {{.*}}void @_ZTSZ4mainE15kernel_function(i32*{{.*}}){{.*}} !dbg [[KERNEL:![0-9]+]] {{.*}}{
-// CHECK-NEW: define{{.*}} spir_kernel {{.*}}void @_ZTSZ4mainE15kernel_function(i32 addrspace(4)*{{.*}}){{.*}} !dbg [[KERNEL:![0-9]+]] {{.*}}{
+// CHECK: define{{.*}} spir_kernel {{.*}}19use_kernel_for_test({{.*}}){{.*}} !dbg [[KERNEL:![0-9]+]] {{.*}}{
+// CHECK: getelementptr inbounds %"class.{{.*}}.anon"{{.*}} !dbg [[LINE_A0:![0-9]+]]
+// CHECK: call spir_func void {{.*}}6__init{{.*}} !dbg [[LINE_A0]]
+// CHECK: call spir_func void @"_ZZ4mainENK3$_0clEv"{{.*}} !dbg [[LINE_B0:![0-9]+]]
+// CHECK: ret void, !dbg [[LINE_A0]]
 // CHECK: [[FILE:![0-9]+]] = !DIFile(filename: "{{.*}}debug-info-srcpos-kernel.cpp"{{.*}})
-// CHECK: [[KERNEL]] = {{.*}}!DISubprogram(name: "_ZTSZ4mainE15kernel_function"
+// CHECK: [[KERNEL]] = {{.*}}!DISubprogram(name: "{{.*}}19use_kernel_for_test"
 // CHECK-SAME: scope: [[FILE]]
 // CHECK-SAME: file: [[FILE]]
 // CHECK-SAME: flags: DIFlagArtificial | DIFlagPrototyped
+// CHECK: [[LINE_A0]] = !DILocation(line: 0
+// CHECK: [[LINE_B0]] = !DILocation(line: 0
+