From 783fe2e01b08de28ff88daab8ef2b4b299679bd7 Mon Sep 17 00:00:00 2001
From: Tony Tye <Tony.Tye@amd.com>
Date: Sun, 10 Jan 2021 02:32:45 +0000
Subject: [PATCH] Delete roc::Program::hsaDevice

Delete roc::Program::hsaDevice and access directly from device
associated with program. This allows to be clear if the device is a
NullDevice which has no meaningful HSA agent backend device.

Change-Id: I81f96aff47bf9b8166d0ff6a5efc7c01f0fb6de3
---
 rocclr/device/rocm/rockernel.cpp  | 11 +++++------
 rocclr/device/rocm/rocprogram.hpp |  3 ---
 2 files changed, 5 insertions(+), 9 deletions(-)

diff --git a/rocclr/device/rocm/rockernel.cpp b/rocclr/device/rocm/rockernel.cpp
index 25498a4e12..273dee2e31 100644
--- a/rocclr/device/rocm/rockernel.cpp
+++ b/rocclr/device/rocm/rockernel.cpp
@@ -71,7 +71,7 @@ bool LightningKernel::init() {
   // Get the kernel code handle
   hsa_status_t hsaStatus;
   hsa_executable_symbol_t symbol;
-  hsa_agent_t agent = program()->hsaDevice();
+  hsa_agent_t agent = program()->rocDevice().getBackendDevice();
   hsaStatus = hsa_executable_get_symbol_by_name(program()->hsaExecutable(),
                                                 symbolName().c_str(),
                                                 &agent, &symbol);
@@ -142,8 +142,8 @@ bool LightningKernel::init() {
   }
 
   uint32_t wavefront_size = 0;
-  if (hsa_agent_get_info(program()->hsaDevice(), HSA_AGENT_INFO_WAVEFRONT_SIZE, &wavefront_size) !=
-      HSA_STATUS_SUCCESS) {
+  if (hsa_agent_get_info(program()->rocDevice().getBackendDevice(), HSA_AGENT_INFO_WAVEFRONT_SIZE,
+                         &wavefront_size) != HSA_STATUS_SUCCESS) {
     DevLogPrintfError("[ROC][Kernel] Cannot get Wavefront Size, failed with hsa_status: %d \n",
                       hsaStatus);
     return false;
@@ -177,8 +177,6 @@ bool LightningKernel::init() {
 #if defined(WITH_COMPILER_LIB)
 bool HSAILKernel::init() {
   acl_error errorCode;
-  // compile kernel down to ISA
-  hsa_agent_t hsaDevice = program()->hsaDevice();
   // Pull out metadata from the ELF
   size_t sizeOfArgList;
   aclCompiler* compileHandle = program()->rocDevice().compiler();
@@ -222,7 +220,8 @@ bool HSAILKernel::init() {
   }
 
   uint32_t wavefront_size = 0;
-  hsa_status_t hsaStatus = hsa_agent_get_info(program()->hsaDevice(), HSA_AGENT_INFO_WAVEFRONT_SIZE, &wavefront_size);
+  hsa_status_t hsaStatus = hsa_agent_get_info(program()->rocDevice().getBackendDevice(),
+                                              HSA_AGENT_INFO_WAVEFRONT_SIZE, &wavefront_size);
   if (HSA_STATUS_SUCCESS != hsaStatus) {
     DevLogPrintfError("Could not get Wave Info Size: %d, failed with hsa_status: %d \n",
                       errorCode, hsaStatus);
diff --git a/rocclr/device/rocm/rocprogram.hpp b/rocclr/device/rocm/rocprogram.hpp
index cc74aa6320..2cc4bb6ab4 100644
--- a/rocclr/device/rocm/rocprogram.hpp
+++ b/rocclr/device/rocm/rocprogram.hpp
@@ -57,9 +57,6 @@ class Program : public device::Program {
     return static_cast<const Device&>(device());
   }
 
-  //! Returns the hsaBinary associated with the program
-  hsa_agent_t hsaDevice() const { return rocNullDevice().getBackendDevice(); }
-
   hsa_executable_t hsaExecutable() const {
     assert(!isNull());
     return hsaExecutable_;