undo parallel

openvinotoolkit · Jan 15, 2025 · c1cd58d · c1cd58d
1 parent 8c13833
commit c1cd58d
Show file tree

Hide file tree

Showing 2 changed files with 49 additions and 83 deletions.
diff --git a/.../kernel_selector/cl_kernels/experimental_detectron_generate_proposals_single_image_ref.cl b/.../kernel_selector/cl_kernels/experimental_detectron_generate_proposals_single_image_ref.cl
@@ -138,97 +138,66 @@ inline void FUNC(bubbleSortIterative)(__global Box* arr, int l, int h) {
     }
 }
 
-// 1. Sort boxes by scores
-KERNEL(edgpsi_ref_stage_1)(__global OUTPUT_TYPE* proposals) {
-    __global Box* boxes = (__global Box*)proposals;
-    const int id = get_global_id(0);
-    int l = 0, h = NUM_PROPOSALS-1;
+inline void FUNC(quickSortIterative)(__global Box* arr, int l, int h) {
     // Create an auxiliary stack
     const int kStackSize = 100;
-    __local int stack[kStackSize];
-    // initialize top of stack
-    __local int top; 
-    if(id==0) {      
-        top = -1;
-        // push initial values of l and h to stack
-        stack[++top] = l;
-        stack[++top] = h;
-
-        // Keep popping from stack while is not empty
-        int run_counter = 0;
-        const int maximal_counter = 4; //log2(16) , 16 is worker num
-        while (top >= 0 && run_counter < maximal_counter) {
-            run_counter++;
-            // Pop h and l
-            h = stack[top--];
-            l = stack[top--];
-            // Set pivot element at its correct position
-            // in sorted array
-            int p = FUNC_CALL(partition)(boxes, l, h);
-
-            // If there are elements on left side of pivot,
-            // then push left side to stack
-            if (p - 1 > l && l < PRE_NMS_TOPN) {
-                if (top >= (kStackSize - 1)) {
-                    FUNC_CALL(bubbleSortIterative)(boxes, l, p - 1);
-                } else {
-                    stack[++top] = l;
-                    stack[++top] = p - 1;
-                }
-            }
+    int stack[kStackSize];
 
-            // If there are elements on right side of pivot,
-            // then push right side to stack
-            if (p + 1 < h && l < PRE_NMS_TOPN) {
-                if (top >= (kStackSize - 1)) {
-                    FUNC_CALL(bubbleSortIterative)(boxes, p + 1, h);
-                } else {
-                    stack[++top] = p + 1;
-                    stack[++top] = h;
-                }
+    // initialize top of stack
+    int top = -1;
+
+    // push initial values of l and h to stack
+    stack[++top] = l;
+    stack[++top] = h;
+
+    // Keep popping from stack while is not empty
+    while (top >= 0) {
+        // Pop h and l
+        h = stack[top--];
+        l = stack[top--];
+        bool all_zeroes = true; //when all zeroes algorithm stuck
+        for(int i=l;i<h;i++) {
+            if(arr[i].score != 0.0f) {
+                all_zeroes = false;
+                break;
             }
         }
-    }
-    barrier(CLK_LOCAL_MEM_FENCE);
-    if(id<top/2) {
-        l = stack[id*2];
-        h = stack[id*2+1];
-        int private_stack[kStackSize];
-        int private_top = -1;
-        private_stack[++private_top] = l;
-        private_stack[++private_top] = h;
-        while (private_top >= 0) {
-            // Pop h and l
-            h = private_stack[private_top--];
-            l = private_stack[private_top--];
-            // Set pivot element at its correct position
-            // in sorted array
-            int p = FUNC_CALL(partition)(boxes, l, h);
-
-            // If there are elements on left side of pivot,
-            // then push left side to stack
-            if (p - 1 > l && l<PRE_NMS_TOPN) {
-                if (top >= (kStackSize - 1)) {
-                    FUNC_CALL(bubbleSortIterative)(boxes, l, p - 1);
-                } else {
-                    private_stack[++private_top] = l;
-                    private_stack[++private_top] = p - 1;
-                }
+        if(all_zeroes) {
+            continue;
+        }
+        // Set pivot element at its correct position
+        // in sorted array
+        int p = FUNC_CALL(partition)(arr, l, h);
+
+        // If there are elements on left side of pivot,
+        // then push left side to stack
+        if (p - 1 > l && l < PRE_NMS_TOPN) {
+            if (top >= (kStackSize - 1)) {
+                FUNC_CALL(bubbleSortIterative)(arr, l, p - 1);
+            } else {
+                stack[++top] = l;
+                stack[++top] = p - 1;
             }
+        }
 
-            // If there are elements on right side of pivot,
-            // then push right side to stack
-            if (p + 1 < h && l < PRE_NMS_TOPN) {
-                if (top >= (kStackSize - 1)) {
-                    FUNC_CALL(bubbleSortIterative)(boxes, p + 1, h);
-                } else {
-                    private_stack[++private_top] = p + 1;
-                    private_stack[++private_top] = h;
-                }
+        // If there are elements on right side of pivot,
+        // then push right side to stack
+        if (p + 1 < h && l < PRE_NMS_TOPN) {
+            if (top >= (kStackSize - 1)) {
+                FUNC_CALL(bubbleSortIterative)(arr, p + 1, h);
+            } else {
+                stack[++top] = p + 1;
+                stack[++top] = h;
             }
         }
     }
 }
+
+// 1. Sort boxes by scores
+KERNEL(edgpsi_ref_stage_1)(__global OUTPUT_TYPE* proposals) {
+    __global Box* boxes = (__global Box*)proposals;
+    FUNC_CALL(quickSortIterative)(boxes, 0, NUM_PROPOSALS-1);
+}
 #undef Box
 #endif /* EDGPSI_STAGE_1 */
 

diff --git a/...el_gpu/src/kernel_selector/kernels/ed_gpsi/generate_proposals_single_image_kernel_ref.cpp b/...el_gpu/src/kernel_selector/kernels/ed_gpsi/generate_proposals_single_image_kernel_ref.cpp
@@ -51,9 +51,6 @@ ExperimentalDetectronGenerateProposalsSingleImageRef::DispatchData SetDefault(
         dispatch_data.gws = {bottom_H, bottom_W, anchors_num};
     } else if (idx == 3) {
         dispatch_data.gws = {params.post_nms_count, 1, 1};
-    } else if (idx == 1) {
-        dispatch_data.gws = {16, 1, 1};
-        dispatch_data.lws = {16, 1, 1};
     } else {
         dispatch_data.gws = {1, 1, 1};
     }