From: Yaxun Liu <Yaxun.Liu@amd.com>
Date: Fri, 14 Jun 2019 15:54:47 +0000 (+0000)
Subject: [AMDGPU] Enable the implicit arguments for HIP (CLANG)
X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=6279011307b083e8bc984ee021fba12dee435010;p=clang

[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


git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@363414 91177308-0d34-0410-b5e6-96231b3b80d8
---

diff --git a/lib/CodeGen/TargetInfo.cpp b/lib/CodeGen/TargetInfo.cpp
index f39764d1a4..4ef57786f4 100644
--- a/lib/CodeGen/TargetInfo.cpp
+++ b/lib/CodeGen/TargetInfo.cpp
@@ -7868,7 +7868,8 @@ void AMDGPUTargetCodeGenInfo::setTargetAttributes(
   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");
 
diff --git a/test/CodeGenCUDA/amdgpu-hip-implicit-kernarg.cu b/test/CodeGenCUDA/amdgpu-hip-implicit-kernarg.cu
new file mode 100644
index 0000000000..8f730ac14c
--- /dev/null
+++ b/test/CodeGenCUDA/amdgpu-hip-implicit-kernarg.cu
@@ -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"