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

[Yaxun Liu via cfe-commits]

Author: yaxunl
Date: Fri Jun 14 08:54:47 2019
New Revision: 363414

URL: http://llvm.org/viewvc/llvm-project?rev=363414&view=rev
Log:
[AMDGPU] Enable the implicit arguments for HIP (CLANG)

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

Differential Revision: https://reviews.llvm.org/D62244

Added:
    cfe/trunk/test/CodeGenCUDA/amdgpu-hip-implicit-kernarg.cu
Modified:
    cfe/trunk/lib/CodeGen/TargetInfo.cpp

Modified: cfe/trunk/lib/CodeGen/TargetInfo.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/TargetInfo.cpp?rev=363414&r1=363413&r2=363414&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/TargetInfo.cpp (original)
+++ cfe/trunk/lib/CodeGen/TargetInfo.cpp Fri Jun 14 08:54:47 2019
@@ -7868,7 +7868,8 @@ void AMDGPUTargetCodeGenInfo::setTargetA
   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");
 

Added: cfe/trunk/test/CodeGenCUDA/amdgpu-hip-implicit-kernarg.cu
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenCUDA/amdgpu-hip-implicit-kernarg.cu?rev=363414&view=auto
==============================================================================
--- cfe/trunk/test/CodeGenCUDA/amdgpu-hip-implicit-kernarg.cu (added)
+++ cfe/trunk/test/CodeGenCUDA/amdgpu-hip-implicit-kernarg.cu Fri Jun 14 08:54:47 2019
@@ -0,0 +1,8 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -emit-llvm -x hip -o - %s | FileCheck %s
+#include "Inputs/cuda.h"
+
+__global__ void hip_kernel_temp() {
+}
+
+// CHECK: attributes {{.*}} = {{.*}} "amdgpu-implicitarg-num-bytes"="48"