From cabce71845f442fdf5fde2459b94efae8aef241f Mon Sep 17 00:00:00 2001 From: Yaxun Liu Date: Fri, 14 Jun 2019 15:54:47 +0000 Subject: [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 llvm-svn: 363414 --- clang/lib/CodeGen/TargetInfo.cpp | 3 ++- clang/test/CodeGenCUDA/amdgpu-hip-implicit-kernarg.cu | 8 ++++++++ 2 files changed, 10 insertions(+), 1 deletion(-) create mode 100644 clang/test/CodeGenCUDA/amdgpu-hip-implicit-kernarg.cu diff --git a/clang/lib/CodeGen/TargetInfo.cpp b/clang/lib/CodeGen/TargetInfo.cpp index f39764d1a46..4ef57786f44 100644 --- a/clang/lib/CodeGen/TargetInfo.cpp +++ b/clang/lib/CodeGen/TargetInfo.cpp @@ -7868,7 +7868,8 @@ void AMDGPUTargetCodeGenInfo::setTargetAttributes( const auto *ReqdWGS = M.getLangOpts().OpenCL ? FD->getAttr() : nullptr; - if (M.getLangOpts().OpenCL && FD->hasAttr() && + if (((M.getLangOpts().OpenCL && FD->hasAttr()) || + (M.getLangOpts().HIP && FD->hasAttr())) && (M.getTriple().getOS() == llvm::Triple::AMDHSA)) F->addFnAttr("amdgpu-implicitarg-num-bytes", "48"); diff --git a/clang/test/CodeGenCUDA/amdgpu-hip-implicit-kernarg.cu b/clang/test/CodeGenCUDA/amdgpu-hip-implicit-kernarg.cu new file mode 100644 index 00000000000..8f730ac14c5 --- /dev/null +++ b/clang/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" -- cgit v1.2.3