intel/llvm
1
0
Fork 0
mirror of https://github.com/intel/llvm.git synced 2026-01-29 12:53:33 +08:00
Code Issues Packages Projects Releases Wiki Activity

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

Browse Source 
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
This commit is contained in:
Yaxun Liu
2019-06-14 15:54:47 +00:00
parent 573ffd88a0
commit cabce71845
2 changed files with 10 additions and 1 deletions
Download Patch File Download Diff File Expand all files Collapse all files

3
clang/lib/CodeGen/TargetInfo.cpp
View File

@@ -7868,7 +7868,8 @@ void AMDGPUTargetCodeGenInfo::setTargetAttributes(
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");

8
clang/test/CodeGenCUDA/amdgpu-hip-implicit-kernarg.cu Normal file
View File

@@ -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"