[HIP][CodeGen] Make local variable use correct address space for CGDecl (#5374)

The code:
```
cgh.parallel_for_work_group<class WkGrp>(
            sycl::range<1>{N / 2}, sycl::range<1>{2}, [=](sycl::group<1> myGroup) {
            auto floo = 2.0;
            myGroup.parallel_for_work_item(
                [&](sycl::h_item<1> it) { acc[it.get_global_id()] = floo; });
            });
```
Was failing for the HIP backend. 

The variable `floo` is in local memory. Since the default address space for variables is `private`, the address space needs to be changed to `local`. The line `LangAS AS = GetGlobalVarAddressSpace(&D); ` correctly gets the appropriate address space.  However, when checking for the value of the address space, the new address space is not used, rather the old one 

```
if (Ty.getAddressSpace() == LangAS::opencl_local ||
      Ty.getAddressSpace() == LangAS::sycl_local ||
```
This results in `floo` being initialized as 

```
Init = EmitNullConstant(Ty);
```
Which is incorrect. Instead of 
```
Init = llvm::UndefValue::get(LTy);
```

On AMD, `floo` not being an `UndefValue` throws an assert later on in the compilation chain in `llvm/lib/Target/AMDGPU/AMDGPUAsmPrinter.cpp`

```
void AMDGPUAsmPrinter::emitGlobalVariable(const GlobalVariable *GV) {
  if (GV->getAddressSpace() == AMDGPUAS::LOCAL_ADDRESS) {
    if (GV->hasInitializer() && !isa<UndefValue>(GV->getInitializer())) {
      OutContext.reportError({},
                             Twine(GV->getName()) +
                                 ": unsupported initializer for address space");
      return;
    }
```

Therefore it is necessary to use the address space returned by `GetGlobalVarAddressSpace(&D)`, instead of the default address space when seeing whether the variable should be initialized as an UndefValue or a NullConstant. This results in the proper initialization of this kind of local variable as an `UndefValue`.

Author: hdelan

clang/lib/CodeGen/CGDecl.cpp

@@ -266,8 +266,7 @@ llvm::Constant *CodeGenModule::getOrCreateStaticVarDecl(
   // OpenCL/SYCL variables in local address space and CUDA shared
   // variables cannot have an initializer.
   llvm::Constant *Init = nullptr;
-  if (Ty.getAddressSpace() == LangAS::opencl_local ||
-      Ty.getAddressSpace() == LangAS::sycl_local ||
+  if (AS == LangAS::opencl_local || AS == LangAS::sycl_local ||
       D.hasAttr<CUDASharedAttr>() || D.hasAttr<LoaderUninitializedAttr>())
     Init = llvm::UndefValue::get(LTy);
   else

clang/test/CodeGenSYCL/wg_init.cpp

@@ -0,0 +1,22 @@
+// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s
+
+// This test checks that a local variable initialized within a
+// parallel_for_work_group scope is initialized as an UndefValue in addrspace(3)
+// in LLVM IR.
+
+#include "Inputs/sycl.hpp"
+
+using namespace sycl;
+
+int main() {
+  queue q;
+  q.submit([&](handler &h) {
+    h.parallel_for_work_group<class kernel>(
+        range<1>{1}, range<1>{1}, [=](group<1> G) {
+          int WG_VAR = 10;
+        });
+  });
+  return 0;
+}
+
+// CHECK: @{{.*}}WG_VAR = internal addrspace(3) global {{.*}} undef, {{.*}}