Author: yaxunl Date: Fri Jun 14 08:54:47 2019 New Revision: 363414 URL: http://llvm.org/viewvc/llvm-project?rev=363414&view=rev Log: [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 Added: cfe/trunk/test/CodeGenCUDA/amdgpu-hip-implicit-kernarg.cu Modified: cfe/trunk/lib/CodeGen/TargetInfo.cpp Modified: cfe/trunk/lib/CodeGen/TargetInfo.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/TargetInfo.cpp?rev=363414&r1=363413&r2=363414&view=diff ============================================================================== --- cfe/trunk/lib/CodeGen/TargetInfo.cpp (original) +++ cfe/trunk/lib/CodeGen/TargetInfo.cpp Fri Jun 14 08:54:47 2019 @@ -7868,7 +7868,8 @@ void AMDGPUTargetCodeGenInfo::setTargetA 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"); Added: cfe/trunk/test/CodeGenCUDA/amdgpu-hip-implicit-kernarg.cu URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/CodeGenCUDA/amdgpu-hip-implicit-kernarg.cu?rev=363414&view=auto ============================================================================== --- cfe/trunk/test/CodeGenCUDA/amdgpu-hip-implicit-kernarg.cu (added) +++ cfe/trunk/test/CodeGenCUDA/amdgpu-hip-implicit-kernarg.cu Fri Jun 14 08:54:47 2019 @@ -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" _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits