Remove USE_DISPATCH_HSA_KERNEL=0 path.

hipamd/src/hip_hcc.cpp

@@ -387,57 +387,6 @@ void ihipStream_t::lockclose_postKernelCommand(const char * kernelName, hc::acce
 };
 
 
-
-
-#if USE_DISPATCH_HSA_KERNEL==0
-// Precursor: the stream is already locked,specifically so this routine can enqueue work into the specified av.
-void ihipStream_t::launchModuleKernel(
-                        hc::accelerator_view av,
-                        hsa_signal_t signal,
-                        uint32_t blockDimX,
-                        uint32_t blockDimY,
-                        uint32_t blockDimZ,
-                        uint32_t gridDimX,
-                        uint32_t gridDimY,
-                        uint32_t gridDimZ,
-                        uint32_t groupSegmentSize,
-			uint32_t privateSegmentSize,
-                        void *kernarg,
-                        size_t kernSize,
-                        uint64_t kernel){
-    hsa_status_t status;
-    hsa_queue_t *Queue = (hsa_queue_t*)av.get_hsa_queue();
-    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 = groupSegmentSize;
-    dispatch_packet->private_segment_size = privateSegmentSize;
-    dispatch_packet->kernarg_address = kernarg;
-    dispatch_packet->kernel_object = kernel;
-    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 = 3 << 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);
-}
-#endif
-
-
 //=============================================================================
 // Recompute the peercnt and the packed _peerAgents whenever a peer is added or deleted.
 // The packed _peerAgents can efficiently be used on each memory allocation.

hipamd/src/hip_hcc.h

@@ -33,9 +33,7 @@ THE SOFTWARE.
 #error("This version of HIP requires a newer version of HCC.");
 #endif
 
-#define USE_DISPATCH_HSA_KERNEL 1
 #define USE_IPC 0
-//
 
 
 //---

hipamd/src/hip_module.cpp

@@ -300,7 +300,6 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f,
         grid_launch_parm lp;
         hStream = ihipPreLaunchKernel(hStream, 0, 0, &lp, "ModuleKernel");
 
-#if USE_DISPATCH_HSA_KERNEL
 
         hsa_kernel_dispatch_packet_t aql;
 
@@ -325,65 +324,6 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f,
                                     (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE);
 
         lp.av->dispatch_hsa_kernel(&aql, config[1] /* kernarg*/, kernArgSize);
-#else
-
-        /*
-          Create signal
-        */
-
-        hsa_signal_t signal;
-        status = hsa_signal_create(1, 0, NULL, &signal);
-        if(status != HSA_STATUS_SUCCESS){
-            return ihipLogStatus(hipErrorNotFound);
-        }
-
-        /*
-           Allocate kernarg
-        */
-        void *kern = nullptr;
-
-        hsa_amd_memory_pool_t *pool = reinterpret_cast<hsa_amd_memory_pool_t*>(lp.av->get_hsa_kernarg_region());
-        status = hsa_amd_memory_pool_allocate(*pool, kernArgSize, 0, &kern);
-        if(status != HSA_STATUS_SUCCESS){
-            return ihipLogStatus(hipErrorNotFound);
-        }
-        status = hsa_amd_agents_allow_access(1, (hsa_agent_t*)lp.av->get_hsa_agent(), 0, kern);
-        if(status != HSA_STATUS_SUCCESS){
-            return ihipLogStatus(hipErrorNotFound);
-        }
-        memcpy(kern, config[1], kernArgSize);
-
-
-        /*
-          Launch AQL packet
-        */
-        hStream->launchModuleKernel(*lp.av, signal, blockDimX, blockDimY, blockDimZ,
-                  gridDimX, gridDimY, gridDimZ, groupSegmentSize, privateSegmentSize, kern, kernArgSize, f->_object);
-
-
-        /*
-          Wait for signal
-        */
-
-        hsa_signal_value_t value = hsa_signal_wait_acquire(signal, HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_BLOCKED);
-
-        /*
-           Destroy kernarg
-        */
-        status = hsa_amd_memory_pool_free(kern);
-        if(status != HSA_STATUS_SUCCESS){
-            return ihipLogStatus(hipErrorNotFound);
-        }
-
-        /*
-          Destroy the signal
-        */
-        status = hsa_signal_destroy(signal);
-        if(status != HSA_STATUS_SUCCESS){
-            return ihipLogStatus(hipErrorNotFound);
-        }
-
-#endif // USE_DISPATCH_HSA_KERNEL
 
 
         ihipPostLaunchKernel("ModuleKernel", hStream, lp);