Page MenuHomePhabricator

D62244.diff
No OneTemporary

File Metadata

Created
Tue, Jul 16, 9:18 PM

D62244.diff

Index: cfe/trunk/lib/CodeGen/TargetInfo.cpp
===================================================================
--- cfe/trunk/lib/CodeGen/TargetInfo.cpp
+++ cfe/trunk/lib/CodeGen/TargetInfo.cpp
@@ -7868,7 +7868,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: cfe/trunk/test/CodeGenCUDA/amdgpu-hip-implicit-kernarg.cu
===================================================================
--- cfe/trunk/test/CodeGenCUDA/amdgpu-hip-implicit-kernarg.cu
+++ cfe/trunk/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"

Event Timeline