[Clang][OpenMP] Process debug info for reduction - fix SWDEV-535727 (llvm#2748)

Author: ronlieb

clang/lib/CodeGen/CGOpenMPRuntime.cpp

@@ -9720,6 +9720,17 @@ static void emitTargetCallKernelLaunch(
         ++RedVarCount;
       }
     }
+    // Process debug info.
+    if (CGF.CGM.getCodeGenOpts().getDebugInfo() !=
+        llvm::codegenoptions::NoDebugInfo) {
+      auto FillInfoMap = [&](MappableExprsHandler::MappingExprInfo &MapExpr) {
+        return emitMappingInformation(CGF, OMPBuilder, MapExpr);
+      };
+
+      CombinedInfo.Names.resize(CombinedInfo.Exprs.size());
+      llvm::transform(CombinedInfo.Exprs, CombinedInfo.Names.begin(),
+                      FillInfoMap);
+    }
   }
 
   CGOpenMPRuntime::TargetDataInfo Info;

clang/test/OpenMP/xteam_red_debug_info.c

@@ -0,0 +1,16 @@
+// RUN: %clang -g %s -fopenmp --offload-arch=gfx90a -S --offload-host-only -emit-llvm -o - | FileCheck %s
+
+void test_xteam_red_debug_info() {
+  int N = 100000;
+  double c[N];
+  double sum = 0.0;
+  #pragma omp target teams distribute parallel for reduction(+: sum)
+  for (int i=0; i<N; i++){
+    sum += c[i];
+  }
+  sum = sum/(double)N;
+}
+
+// CHECK:       @.offload_sizes = private unnamed_addr constant [6 x i64]
+// CHECK-NEXT:  @.offload_maptypes = private unnamed_addr constant [6 x i64]
+// CHECK-NEXT:  @.offload_mapnames = private constant [6 x ptr]