From 60473ae542f6b388e96a5c5407b2bbb64cc60325 Mon Sep 17 00:00:00 2001
From: Rahul Garg <rahul.garg@amd.com>
Date: Fri, 17 Jul 2020 00:04:36 +0000
Subject: [PATCH] Handle size 0 symbols

Change-Id: I1629c5027ec5f9af48c8e9e0696829b844423096
---
 rocclr/device/rocm/rocprogram.cpp | 51 ++++++++++++++++---------------
 1 file changed, 27 insertions(+), 24 deletions(-)

diff --git a/rocclr/device/rocm/rocprogram.cpp b/rocclr/device/rocm/rocprogram.cpp
index ef9156d798..977ce3442f 100644
--- a/rocclr/device/rocm/rocprogram.cpp
+++ b/rocclr/device/rocm/rocprogram.cpp
@@ -197,32 +197,35 @@ bool Program::createGlobalVarObj(amd::Memory** amd_mem_obj, void** device_pptr,
     return false;
   }
 
-  /* Find HSA Symbol Address */
-  status = hsa_executable_symbol_get_info(global_symbol,
+  // Handle size 0 symbols
+  if (*bytes != 0) {
+    // Find HSA Symbol Address
+    status = hsa_executable_symbol_get_info(global_symbol,
                                           HSA_EXECUTABLE_SYMBOL_INFO_VARIABLE_ADDRESS, device_pptr);
-  if (status != HSA_STATUS_SUCCESS) {
-    buildLog_ += "Error: Failed to find the Symbol Address : ";
-    buildLog_ += hsa_strerror(status);
-    buildLog_ += "\n";
-    return false;
+    if (status != HSA_STATUS_SUCCESS) {
+      buildLog_ += "Error: Failed to find the Symbol Address : ";
+      buildLog_ += hsa_strerror(status);
+      buildLog_ += "\n";
+      return false;
+    }
+
+    roc_device = static_cast<const roc::Device*>(&dev());
+    *amd_mem_obj = new(roc_device->context()) amd::Buffer(roc_device->context(), 0, *bytes,
+                                                          *device_pptr);
+
+    if (*amd_mem_obj == nullptr) {
+      buildLog_ += "[OCL] Failed to create a mem object!";
+      buildLog_ += "\n";
+      return false;
+    }
+
+    if (!((*amd_mem_obj)->create(nullptr))) {
+      buildLog_ += "[OCL] failed to create a svm hidden buffer!";
+      buildLog_ += "\n";
+      (*amd_mem_obj)->release();
+      return false;
+    }
   }
-
-  roc_device = static_cast<const roc::Device*>(&dev());
-  *amd_mem_obj = new(roc_device->context()) amd::Buffer(roc_device->context(), 0, *bytes, *device_pptr);
-
-  if (*amd_mem_obj == nullptr) {
-    buildLog_ += "[OCL] Failed to create a mem object!";
-    buildLog_ += "\n";
-    return false;
-  }
-
-  if (!((*amd_mem_obj)->create(nullptr))) {
-    buildLog_ += "[OCL] failed to create a svm hidden buffer!";
-    buildLog_ += "\n";
-    (*amd_mem_obj)->release();
-    return false;
-  }
-
   return true;
 }