[AMDGPU] Enable the implicit arguments for HIP (CLANG)

[cdevadas]

Summary

Enable 48-bytes of implicit arguments for HIP as well. Earlier it was enabled for OpenCL. This code is specific to AMDGPU target.

[yaxunl]

Currently HIP and CUDA share the same test directories, so better put the test in CodeGenCUDA.

[cdevadas]

Moved the test to CodeGenCUDA directory.

[hliao]

test/CodeGenCUDA/amdgpu-hip-implicit-kernarg.cu

8

For a single check, you don't need CHECK-DAG.

[cdevadas]

simplified the check in the test case.

lib/CodeGen/TargetInfo.cpp

void AMDGPUTargetCodeGenInfo::setTargetAttributes(
   if (!FD)
     return;
 
   llvm::Function *F = cast<llvm::Function>(GV);
 
   const auto *ReqdWGS = M.getLangOpts().OpenCL ?
     FD->getAttr<ReqdWorkGroupSizeAttr>() : nullptr;
 
-  if (M.getLangOpts().OpenCL && FD->hasAttr<OpenCLKernelAttr>() &&
+  if (((M.getLangOpts().OpenCL && FD->hasAttr<OpenCLKernelAttr>()) ||
+      (M.getLangOpts().HIP && FD->hasAttr<CUDAGlobalAttr>())) &&
       (M.getTriple().getOS() == llvm::Triple::AMDHSA))
     F->addFnAttr("amdgpu-implicitarg-num-bytes", "48");
 
   const auto *FlatWGS = FD->getAttr<AMDGPUFlatWorkGroupSizeAttr>();
   if (ReqdWGS || FlatWGS) {
     unsigned Min = 0;
     unsigned Max = 0;
     if (FlatWGS) {

test/CodeGenCUDA/amdgpu-hip-implicit-kernarg.cu

+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -emit-llvm -x hip -o - %s | FileCheck %s
+#include "Inputs/cuda.h"
+
+__global__ void hip_kernel_temp() {
+}
+
+// CHECK: attributes #0 = { noinline nounwind optnone "amdgpu-implicitarg-num-bytes"="48"