Refactor Module and Function APIs.

- hipFunction_t is now returned by value.  This eliminates dynamic
      allocation / memory management complexity in the module.  Removed
the kernel
      name so the structure is just 16 bytes now.

    - Moved the hsa_executable_load_module and hsa_executable_freeze
      calls to the hipModuleLoad and hipModuleLoadData calls.

    - Apply sharedMemBytes in hipModuleLaunchKernel to group segment
      size (not private).


[ROCm/hip commit: 4d29885be3]

projects/hip/include/hip/hcc_detail/hip_runtime_api.h

@@ -33,7 +33,7 @@ THE SOFTWARE.
 
 #include <hip/hcc_detail/host_defines.h>
 #include <hip/hip_runtime_api.h>
-//#include "hip/hip_hcc.h"
+#include <hsa/hsa.h>
 
 #if defined (__HCC__) &&  (__hcc_workweek__ < 16155)
 #error("This version of HIP requires a newer version of HCC.");
@@ -72,7 +72,12 @@ typedef struct ihipIpcEventHandle_t *hipIpcEventHandle_t;
 
 typedef struct ihipModule_t *hipModule_t;
 
-typedef struct ihipFunction_t *hipFunction_t;
+struct ihipModuleSymbol_t{
+    hsa_executable_symbol_t 	_hsaSymbol;
+    uint64_t 			_object;             // The kernel object.
+};
+
+typedef struct ihipModuleSymbol_t hipFunction_t;
 
 typedef void* hipDeviceptr_t;
 

projects/hip/src/hip_hcc.h

@@ -383,26 +383,6 @@ public:
     size_t psize;
 };
 
-class ihipFunction_t{
-public:
-    ihipFunction_t(const char *name) {
-        size_t nameSz = strlen(name);
-        char *kernelName = (char*)malloc(nameSz);
-        strncpy(kernelName, name, nameSz);
-        _kernelName = kernelName;
-    };
-
-    ~ihipFunction_t() {
-        if (_kernelName) {
-            free((void*)_kernelName);
-            _kernelName = NULL;
-        };
-    };
-public:
-    const char             *_kernelName;
-    hsa_executable_symbol_t _kernelSymbol;
-    uint64_t _kernel;
-};
 
 class ihipModule_t {
 public:
@@ -412,20 +392,7 @@ public:
   void *ptr;
   size_t size;
 
-  ihipModule_t() : executable(), object(), fileName(), ptr(nullptr), size(0), hipFunctionTable() {}
-  ~ihipModule_t() {
-    for (int i = 0; i < hipFunctionTable.size(); ++i) {
-      ihipFunction_t *func = hipFunctionTable[i];
-      delete func;
-    }
-    hipFunctionTable.clear();
-  }
-
-  void registerFunction(ihipFunction_t* func) {
-    hipFunctionTable.push_back(func);
-  }
-private:
-  std::vector<ihipFunction_t*> hipFunctionTable;
+  ihipModule_t() : executable(), object(), fileName(), ptr(nullptr), size(0) {}
 };
 
 template <typename MUTEX_TYPE>

projects/hip/src/hip_module.cpp

@@ -35,6 +35,11 @@ THE SOFTWARE.
 
 //TODO Use Pool APIs from HCC to get memory regions.
 
+#define CHECK_HSA(hsaStatus, hipStatus) \
+if (hsaStatus != HSA_STATUS_SUCCESS) {\
+    return ihipLogStatus(hipStatus);\
+}
+
 namespace hipdrv {
 
     hsa_status_t findSystemRegions(hsa_region_t region, void *data){
@@ -154,9 +159,13 @@ hipError_t hipModuleLoad(hipModule_t *module, const char *fname){
             }
 
             status = hsa_executable_create(HSA_PROFILE_FULL, HSA_EXECUTABLE_STATE_UNFROZEN, NULL, &(*module)->executable);
-            if(status != HSA_STATUS_SUCCESS){
-                return ihipLogStatus(hipErrorNotInitialized);
-            }
+            CHECK_HSA(status, hipErrorNotInitialized);
+
+            status = hsa_executable_load_code_object((*module)->executable, agent, (*module)->object, NULL);
+            CHECK_HSA(status, hipErrorNotInitialized);
+
+            status = hsa_executable_freeze((*module)->executable, NULL);
+            CHECK_HSA(status, hipErrorNotInitialized);
         }
     }
 
@@ -189,7 +198,7 @@ hipError_t hipModuleUnload(hipModule_t hmod)
     return ihipLogStatus(ret);
 }
 
-hipError_t ihipModuleGetFunction(hipFunction_t *func, hipModule_t hmod, const char *name){
+hipError_t ihipModuleGetSymbol(hipFunction_t *func, hipModule_t hmod, const char *name){
     auto ctx = ihipGetTlsDefaultCtx();
     hipError_t ret = hipSuccess;
 
@@ -201,27 +210,21 @@ hipError_t ihipModuleGetFunction(hipFunction_t *func, hipModule_t hmod, const ch
         ret = hipErrorInvalidContext;
 
     }else{
-        *func = new ihipFunction_t(name);
-        hmod->registerFunction(*func);
         int deviceId = ctx->getDevice()->_deviceId;
         ihipDevice_t *currentDevice = ihipGetDevice(deviceId);
         hsa_agent_t gpuAgent = (hsa_agent_t)currentDevice->_hsaAgent;
 
-        hsa_status_t status;
-        status = hsa_executable_load_code_object(hmod->executable, gpuAgent, hmod->object, NULL);
-        if(status != HSA_STATUS_SUCCESS){
-            return ihipLogStatus(hipErrorNotInitialized);
-        }
 
-        status = hsa_executable_freeze(hmod->executable, NULL);
-        status = hsa_executable_get_symbol(hmod->executable, NULL, name, gpuAgent, 0, &(*func)->_kernelSymbol);
+
+        hsa_status_t status;
+        status = hsa_executable_get_symbol(hmod->executable, NULL, name, gpuAgent, 0, &(func->_hsaSymbol));
         if(status != HSA_STATUS_SUCCESS){
             return ihipLogStatus(hipErrorNotFound);
         }
 
-        status = hsa_executable_symbol_get_info((*func)->_kernelSymbol,
+        status = hsa_executable_symbol_get_info(func->_hsaSymbol,
                                    HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT,
-                                   &(*func)->_kernel);
+                                   &func->_object);
 
         if(status != HSA_STATUS_SUCCESS){
             return ihipLogStatus(hipErrorNotFound);
@@ -234,7 +237,7 @@ hipError_t ihipModuleGetFunction(hipFunction_t *func, hipModule_t hmod, const ch
 hipError_t hipModuleGetFunction(hipFunction_t *hfunc, hipModule_t hmod,
                                 const char *name){
     HIP_INIT_API(hfunc, hmod, name);
-    return ihipModuleGetFunction(hfunc, hmod, name);
+    return ihipModuleGetSymbol(hfunc, hmod, name);
 }
 
 
@@ -275,7 +278,7 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f,
         }
 
         uint32_t groupSegmentSize;
-        hsa_status_t status = hsa_executable_symbol_get_info(f->_kernelSymbol,
+        hsa_status_t status = hsa_executable_symbol_get_info(f._hsaSymbol,
                                                              HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZE,
                                                              &groupSegmentSize);
         if(status != HSA_STATUS_SUCCESS){
@@ -283,20 +286,19 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f,
         }
 
         uint32_t privateSegmentSize;
-        status = hsa_executable_symbol_get_info(f->_kernelSymbol,
+        status = hsa_executable_symbol_get_info(f._hsaSymbol,
                                                 HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE,
                                                 &privateSegmentSize);
         if(status != HSA_STATUS_SUCCESS){
             return ihipLogStatus(hipErrorNotFound);
         }
-        privateSegmentSize += sharedMemBytes;
 
 
         /*
           Kernel argument preparation.
         */
         grid_launch_parm lp;
-        hStream = ihipPreLaunchKernel(hStream, 0, 0, &lp, f->_kernelName);
+        hStream = ihipPreLaunchKernel(hStream, 0, 0, &lp, "ModuleKernel");
 
 #if USE_DISPATCH_HSA_KERNEL
 
@@ -315,7 +317,7 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f,
         aql.grid_size_z = blockDimZ * gridDimZ;
         aql.group_segment_size = groupSegmentSize;
         aql.private_segment_size = privateSegmentSize;
-        aql.kernel_object = f->_kernel;
+        aql.kernel_object = f._object;
         aql.setup = 3 << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS;
         aql.header =   (HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE) |
                                     (1 << HSA_PACKET_HEADER_BARRIER) |
@@ -356,7 +358,7 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f,
           Launch AQL packet
         */
         hStream->launchModuleKernel(*lp.av, signal, blockDimX, blockDimY, blockDimZ,
-                  gridDimX, gridDimY, gridDimZ, groupSegmentSize, privateSegmentSize, kern, kernArgSize, f->_kernel);
+                  gridDimX, gridDimY, gridDimZ, groupSegmentSize, privateSegmentSize, kern, kernArgSize, f->_object);
 
 
         /*
@@ -384,7 +386,7 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f,
 #endif // USE_DISPATCH_HSA_KERNEL
 
 
-        ihipPostLaunchKernel(f->_kernelName, hStream, lp);
+        ihipPostLaunchKernel("ModuleKernel", hStream, lp);
 
     }
 
@@ -405,9 +407,9 @@ hipError_t hipModuleGetGlobal(hipDeviceptr_t *dptr, size_t *bytes,
     }
     else{
         hipFunction_t func;
-        ihipModuleGetFunction(&func, hmod, name);
+        ihipModuleGetSymbol(&func, hmod, name);
         *bytes = PrintSymbolSizes(hmod->ptr, name) + sizeof(amd_kernel_code_t);
-        *dptr = reinterpret_cast<void*>(func->_kernel);
+        *dptr = reinterpret_cast<void*>(func._object);
         return ihipLogStatus(ret);
     }
 }
@@ -419,7 +421,7 @@ hipError_t hipModuleLoadData(hipModule_t *module, const void *image)
     hipError_t ret = hipSuccess;
     if(image == NULL || module == NULL){
         return ihipLogStatus(hipErrorNotInitialized);
-    }else{
+    } else {
         auto ctx = ihipGetTlsDefaultCtx();
         *module = new ihipModule_t;
         int deviceId = ctx->getDevice()->_deviceId;
@@ -452,9 +454,13 @@ hipError_t hipModuleLoadData(hipModule_t *module, const void *image)
         }
 
         status = hsa_executable_create(HSA_PROFILE_FULL, HSA_EXECUTABLE_STATE_UNFROZEN, NULL, &(*module)->executable);
-        if(status != HSA_STATUS_SUCCESS){
-            return ihipLogStatus(hipErrorNotInitialized);
-        }
+        CHECK_HSA(status, hipErrorNotInitialized);
+
+        status = hsa_executable_load_code_object((*module)->executable, agent, (*module)->object, NULL);
+        CHECK_HSA(status, hipErrorNotInitialized);
+
+        status = hsa_executable_freeze((*module)->executable, NULL);
+        CHECK_HSA(status, hipErrorNotInitialized);
     }
     return ihipLogStatus(ret);
 }

projects/hip/src/trace_helper.h

@@ -72,6 +72,15 @@ inline std::string ToString(hipEvent_t v)
     return ss.str();
 };
 
+//  hipEvent_t specialization. TODO - maybe add an event ID for debug?
+template <>
+inline std::string ToString(hipFunction_t v) 
+{
+    std::ostringstream ss;
+    ss << "0x" << std::hex << v._object;
+    return ss.str();
+};
+
 
 
 //  hipStream_t