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

[Christudasan Devadasan via Phabricator via cfe-commits]

cdevadas updated this revision to Diff 203975.
cdevadas added a comment.

simplified the check in the test case.


Repository:
  rC Clang

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D62244/new/

https://reviews.llvm.org/D62244

Files:
  lib/CodeGen/TargetInfo.cpp
  test/CodeGenCUDA/amdgpu-hip-implicit-kernarg.cu


Index: test/CodeGenCUDA/amdgpu-hip-implicit-kernarg.cu
===================================================================
--- /dev/null
+++ test/CodeGenCUDA/amdgpu-hip-implicit-kernarg.cu
@@ -0,0 +1,7 @@
+// 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"
Index: lib/CodeGen/TargetInfo.cpp
===================================================================
--- lib/CodeGen/TargetInfo.cpp
+++ lib/CodeGen/TargetInfo.cpp
@@ -7853,7 +7853,8 @@
   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");
 



Index: test/CodeGenCUDA/amdgpu-hip-implicit-kernarg.cu
===================================================================
--- /dev/null
+++ test/CodeGenCUDA/amdgpu-hip-implicit-kernarg.cu
@@ -0,0 +1,7 @@
+// 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"
Index: lib/CodeGen/TargetInfo.cpp
===================================================================
--- lib/CodeGen/TargetInfo.cpp
+++ lib/CodeGen/TargetInfo.cpp
@@ -7853,7 +7853,8 @@
   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");
 

_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits