From: Aaron Enye Shi <enye.shi@gmail.com>
Date: Tue, 2 Apr 2019 20:10:18 +0000 (+0000)
Subject: [HIP-Clang] Fat binary should not be produced for non GPU code
X-Git-Tag: llvmorg-10-init~8646
X-Git-Url: http://review.tizen.org/git/?a=commitdiff_plain;h=13d8e929409fb1a1f8a2318250923c90f3abac7c;p=platform%2Fupstream%2Fllvm.git

[HIP-Clang] Fat binary should not be produced for non GPU code

Skip producing the fat binary functions for HIP when no device code is present.

Reviewers: yaxunl

Differential Review: https://reviews.llvm.org/D60141

llvm-svn: 357520
---

diff --git a/clang/lib/CodeGen/CGCUDANV.cpp b/clang/lib/CodeGen/CGCUDANV.cpp
index 0602601..cd969dc8 100644
--- a/clang/lib/CodeGen/CGCUDANV.cpp
+++ b/clang/lib/CodeGen/CGCUDANV.cpp
@@ -472,6 +472,8 @@ llvm::Function *CGNVCUDARuntime::makeModuleCtorFunction() {
   StringRef CudaGpuBinaryFileName = CGM.getCodeGenOpts().CudaGpuBinaryFileName;
   if (CudaGpuBinaryFileName.empty() && !IsHIP)
     return nullptr;
+  if (IsHIP && EmittedKernels.empty() && DeviceVars.empty())
+    return nullptr;
 
   // void __{cuda|hip}_register_globals(void* handle);
   llvm::Function *RegisterGlobalsFunc = makeRegisterGlobalsFn();
diff --git a/clang/test/CodeGenCUDA/device-stub.cu b/clang/test/CodeGenCUDA/device-stub.cu
index 387a787..00d7fa4 100644
--- a/clang/test/CodeGenCUDA/device-stub.cu
+++ b/clang/test/CodeGenCUDA/device-stub.cu
@@ -228,13 +228,19 @@ void hostfunc(void) { kernelfunc<<<1, 1>>>(1, 1, 1); }
 // device-side globals, but we still need to register GPU binary.
 // Skip GPU binary string first.
 // CUDANOGLOBALS: @{{.*}} = private constant{{.*}}
-// HIPNOGLOBALS: @{{.*}} = internal constant{{.*}}
+// HIPNOGLOBALS-NOT: @{{.*}} = internal constant{{.*}}
 // NOGLOBALS-NOT: define internal void @__{{.*}}_register_globals
-// NOGLOBALS: define internal void @__[[PREFIX:cuda|hip]]_module_ctor
-// NOGLOBALS: call{{.*}}[[PREFIX]]RegisterFatBinary{{.*}}__[[PREFIX]]_fatbin_wrapper
+// CUDANOGLOBALS: define internal void @__[[PREFIX:cuda|hip]]_module_ctor
+// CUDANOGLOBALS: call{{.*}}[[PREFIX]]RegisterFatBinary{{.*}}__[[PREFIX]]_fatbin_wrapper
 // NOGLOBALS-NOT: call void @__[[PREFIX]]_register_globals
-// NOGLOBALS: define internal void @__[[PREFIX]]_module_dtor
-// NOGLOBALS: call void @__[[PREFIX]]UnregisterFatBinary
+// CUDANOGLOBALS: define internal void @__[[PREFIX]]_module_dtor
+// CUDANOGLOBALS: call void @__[[PREFIX]]UnregisterFatBinary
+
+// There should be no fat binary functions when no device-code is found for HIP.
+// HIPNOGLOBALS-NOT: define internal void @__[[PREFIX:cuda|hip]]_module_ctor
+// HIPNOGLOBALS-NOT: call{{.*}}[[PREFIX]]RegisterFatBinary{{.*}}__[[PREFIX]]_fatbin_wrapper
+// HIPNOGLOBALS-NOT: define internal void @__[[PREFIX]]_module_dtor
+// HIPNOGLOBALS-NOT: call void @__[[PREFIX]]UnregisterFatBinary
 
 // There should be no constructors/destructors if we have no GPU binary.
 // NOGPUBIN-NOT: define internal void @__[[PREFIX]]_register_globals