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() : 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"); 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"