Added hipLaunchModuleKernel and new error codes

- hipLaunchModuleKernel maps to cuLaunchKernel
- Whole lot of new error codes added for the use of driver api
 - KernelParams arguments is not yet supported
 - hipLaunchModuleKernel is a synchronous api (will change eventually)
 - All the commands in a stream will wait on host when hipLaunchModuleKernel is called on it

Change-Id: Ib4a4fae1db06fbb3a81d5a5575b026aa821264ed

src/hip_module.cpp

@@ -24,7 +24,10 @@ THE SOFTWARE.
 #include "hcc_detail/trace_helper.h"
 #include <fstream>
 
-hsa_status_t FindRegions(hsa_region_t region, void *data){
+//TODO Use Pool APIs from HCC to get memory regions.
+
+namespace hipdrv{
+hsa_status_t findSystemRegions(hsa_region_t region, void *data){
   hsa_region_segment_t segment_id;
   hsa_region_get_info(region, HSA_REGION_INFO_SEGMENT, &segment_id);
 
@@ -44,28 +47,59 @@ hsa_status_t FindRegions(hsa_region_t region, void *data){
   return HSA_STATUS_SUCCESS;
 }
 
+hsa_status_t findKernArgRegions(hsa_region_t region, void *data){
+  hsa_region_segment_t segment_id;
+  hsa_region_get_info(region, HSA_REGION_INFO_SEGMENT, &segment_id);
+
+  if(segment_id != HSA_REGION_SEGMENT_GLOBAL){
+    return HSA_STATUS_SUCCESS;
+  }
+
+  hsa_region_global_flag_t flags;
+  hsa_region_get_info(region, HSA_REGION_INFO_GLOBAL_FLAGS, &flags);
+
+  hsa_region_t *reg = (hsa_region_t*)data;
+
+  if(flags & HSA_REGION_GLOBAL_FLAG_KERNARG){
+    *reg = region;
+  }
+
+  return HSA_STATUS_SUCCESS;
+}
+
+
+}
+
+
+
 hipError_t hipModuleLoad(hipModule *module, const char *fname){
     HIP_INIT_API(fname);
     hipError_t ret = hipSuccess;
+    if(module == NULL){
+        ret = hipErrorInvalidValue;
+    }
     auto ctx = ihipGetTlsDefaultCtx();
     if(ctx == nullptr){
-        ret = hipErrorInvalidDevice;
+        ret = hipErrorInvalidContext;
     }else{
         int deviceId = ctx->getDevice()->_deviceId;
         ihipDevice_t *currentDevice = ihipGetDevice(deviceId);
         std::ifstream in(fname, std::ios::binary | std::ios::ate);
         if(!in){
-            std::cout<<"Couldn't read file "<<fname<<std::endl;
+            return hipErrorFileNotFound;
         }else{
             size_t size = std::string::size_type(in.tellg());
             void *p = NULL;
             hsa_agent_t agent = currentDevice->_hsaAgent;
             hsa_region_t sysRegion;
-            hsa_status_t status = hsa_agent_iterate_regions(agent, FindRegions, &sysRegion);
-            assert(status == HSA_STATUS_SUCCESS);
+            hsa_status_t status = hsa_agent_iterate_regions(agent, hipdrv::findSystemRegions, &sysRegion);
             status = hsa_memory_allocate(sysRegion, size, (void**)&p);
+            if(status != HSA_STATUS_SUCCESS){
+                return hipErrorOutOfMemory;
+            }
             char *ptr = (char*)p;
             if(!ptr){
+                return hipErrorOutOfMemory;
                 std::cout<<"Error: failed to allocate memory for code object"<<std::endl;
             }
             hsa_code_object_t obj;
@@ -73,8 +107,10 @@ hipError_t hipModuleLoad(hipModule *module, const char *fname){
             std::copy(std::istreambuf_iterator<char>(in),
                       std::istreambuf_iterator<char>(), ptr);
             status = hsa_code_object_deserialize(ptr, size, NULL, &obj);
+            if(status != HSA_STATUS_SUCCESS){
+                return hipErrorSharedObjectInitFailed;
+            }
             *module = obj.handle;
-            assert(status == HSA_STATUS_SUCCESS);
         }
     }
     return ret;
@@ -84,31 +120,135 @@ hipError_t hipModuleGetFunction(hipFunction *func, hipModule hmod, const char *n
     HIP_INIT_API(name);
     auto ctx = ihipGetTlsDefaultCtx();
     hipError_t ret = hipSuccess;
+    if(name == nullptr || hmod == 0){
+        return hipErrorInvalidValue;
+    }
     if(ctx == nullptr){
-      ret = hipErrorInvalidDevice;
+        ret = hipErrorInvalidContext;
     }else{
-      int deviceId = ctx->getDevice()->_deviceId;
-      ihipDevice_t *currentDevice = ihipGetDevice(deviceId);
-      hsa_agent_t gpuAgent = (hsa_agent_t)currentDevice->_hsaAgent;
+        int deviceId = ctx->getDevice()->_deviceId;
+        ihipDevice_t *currentDevice = ihipGetDevice(deviceId);
+        hsa_agent_t gpuAgent = (hsa_agent_t)currentDevice->_hsaAgent;
 
-      hsa_status_t status;
-      hsa_executable_symbol_t kernel_symbol;
-      hsa_executable_t executable;
-      status = hsa_executable_create(HSA_PROFILE_FULL, HSA_EXECUTABLE_STATE_UNFROZEN, NULL, &executable);
-      
-      assert(status == HSA_STATUS_SUCCESS);
-      hsa_code_object_t obj;
-      obj.handle = hmod;
-      status = hsa_executable_load_code_object(executable, gpuAgent, obj, NULL);
-      assert(status == HSA_STATUS_SUCCESS);
-      status = hsa_executable_freeze(executable, NULL);
-      assert(status == HSA_STATUS_SUCCESS);
-      status = hsa_executable_get_symbol(executable, NULL, name, gpuAgent, 0, &kernel_symbol);
-      assert(status == HSA_STATUS_SUCCESS);
-      status = hsa_executable_symbol_get_info(kernel_symbol,
+        hsa_status_t status;
+        hsa_executable_symbol_t kernel_symbol;
+        hsa_executable_t executable;
+        status = hsa_executable_create(HSA_PROFILE_FULL, HSA_EXECUTABLE_STATE_UNFROZEN, NULL, &executable);
+        if(status != HSA_STATUS_SUCCESS){
+            return hipErrorNotInitialized;
+        }
+        hsa_code_object_t obj;
+        obj.handle = hmod;
+        status = hsa_executable_load_code_object(executable, gpuAgent, obj, NULL);
+        if(status != HSA_STATUS_SUCCESS){
+            return hipErrorNotInitialized;
+        }
+        status = hsa_executable_freeze(executable, NULL);
+        status = hsa_executable_get_symbol(executable, NULL, name, gpuAgent, 0, &kernel_symbol);
+        if(status != HSA_STATUS_SUCCESS){
+            return hipErrorNotFound;
+        }
+        status = hsa_executable_symbol_get_info(kernel_symbol,
                                    HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT,
                                    func);
-      assert(status == HSA_STATUS_SUCCESS);
+        if(status != HSA_STATUS_SUCCESS){
+            return hipErrorNotFound;
+        }
+    }
+    return ret;
+}
+
+hipError_t hipLaunchModuleKernel(hipFunction f,
+            uint32_t gridDimX, uint32_t gridDimY, uint32_t gridDimZ,
+            uint32_t blockDimX, uint32_t blockDimY, uint32_t blockDimZ,
+            uint32_t sharedMemBytes, hipStream_t hStream,
+            void **kernelParams, void **extra){
+    HIP_INIT_API(f);
+    auto ctx = ihipGetTlsDefaultCtx();
+    hipError_t ret = hipSuccess;
+    if(ctx == nullptr){
+        ret = hipErrorInvalidDevice;
+    }else{
+        int deviceId = ctx->getDevice()->_deviceId;
+        ihipDevice_t *currentDevice = ihipGetDevice(deviceId);
+        hsa_agent_t gpuAgent = (hsa_agent_t)currentDevice->_hsaAgent;
+
+        void *config[5] = {0};
+        size_t kernSize;
+
+        if(extra != NULL){
+            memcpy(config, extra, sizeof(size_t)*5);
+            if(config[0] == HIP_LAUNCH_PARAM_BUFFER_POINTER && config[2] == HIP_LAUNCH_PARAM_BUFFER_SIZE && config[4] == HIP_LAUNCH_PARAM_END){
+                kernSize = *(size_t*)(config[3]);
+            }else{
+                return hipErrorNotInitialized;
+            }
+        }else{
+            return hipErrorInvalidValue;
+        }
+/*
+Kernel argument preparation.
+*/
+
+        hsa_region_t kernArg;
+        hsa_status_t status = hsa_agent_iterate_regions(gpuAgent, hipdrv::findKernArgRegions, &kernArg);
+        void *kern;
+        status = hsa_memory_allocate(kernArg, kernSize, &kern);
+        if(status != HSA_STATUS_SUCCESS){
+            return hipErrorLaunchOutOfResources;
+        }
+        memcpy(kern, config[1], kernSize);
+
+
+/*
+Pre kernel launch 
+
+    stream = ihipSyncAndResolveStream(stream);
+    stream->lockopen_preKernelCommand();
+    hc::accelerator_view av = &stream->_av;
+    hc::completion_future cf = new hc::completion_future;
+*/
+
+        hStream = ihipSyncAndResolveStream(hStream);
+        hc::accelerator_view *av = &hStream->_av;
+        hsa_queue_t *Queue = (hsa_queue_t*)av->get_hsa_queue();
+        hsa_signal_t signal;
+        status = hsa_signal_create(1, 0, NULL, &signal);
+
+/*
+Creating the packets
+*/
+
+        const uint32_t queue_mask = Queue->size-1;
+        uint32_t packet_index = hsa_queue_load_write_index_relaxed(Queue);
+        hsa_kernel_dispatch_packet_t *dispatch_packet = &(((hsa_kernel_dispatch_packet_t*)(Queue->base_address))[packet_index & queue_mask]);
+
+        dispatch_packet->completion_signal = signal;
+        dispatch_packet->workgroup_size_x = blockDimX;
+        dispatch_packet->workgroup_size_y = blockDimY;
+        dispatch_packet->workgroup_size_z = blockDimZ;
+        dispatch_packet->grid_size_x = blockDimX * gridDimX;
+        dispatch_packet->grid_size_y = blockDimY * gridDimY;
+        dispatch_packet->grid_size_z = blockDimZ * gridDimZ;
+
+        dispatch_packet->group_segment_size = 0;
+        dispatch_packet->private_segment_size = sharedMemBytes;
+        dispatch_packet->kernarg_address = kern;
+        dispatch_packet->kernel_object = f;
+        uint16_t header = (HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE) |
+        (1 << HSA_PACKET_HEADER_BARRIER) |
+        (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE) |
+        (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE);
+
+
+        uint16_t setup = 1 << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS;
+        uint32_t header32 = header | (setup << 16);
+
+        __atomic_store_n((uint32_t*)(dispatch_packet), header32, __ATOMIC_RELEASE);
+
+        hsa_queue_store_write_index_relaxed(Queue, packet_index+1);
+        hsa_signal_store_relaxed(Queue->doorbell_signal, packet_index);
+        hsa_signal_value_t value = hsa_signal_wait_acquire(signal, HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_BLOCKED);
     }
     return ret;
 }