[HIP][CodeGen] Making local variable use correct address space for CGDecl

[hdelan]

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

cc @npmiller

[Fznamznon]

Thanks for the fix. The change LGTM, but could you please add a test?
Also, please note that HIP AMD GPU Test Suite is failing (with some of tests xpassing though).

[hdelan]

Here is PR for removing XFAILs in test suite intel/llvm-test-suite#763

[hdelan]

Test added in clang

[Fznamznon]

Thanks!

[smanna12]

LGTM

[bader]

@hdelan, I think we need to disable HierPar/hier_par_wgscope.cpp on hip back-end in addition to intel/llvm-test-suite#763. It hangs with this patch. Please, add // UNSUPPORTED: hip comment to disable it.

[hdelan]

@hdelan, I think we need to disable HierPar/hier_par_wgscope.cpp on hip back-end in addition to intel/llvm-test-suite#763. It hangs with this patch. Please, add // UNSUPPORTED: hip comment to disable it.

Thanks, change made in intel/llvm-test-suite#763