             do
             {
                 ref_item_idx++;
-            } while (ref_item_idx < max_idx && !kept_bboxinfo_flag[ref_item_idx - offset]);
-
-            ref_bbox_idx = share_location ? (beforeNMS_index_array[ref_item_idx] % num_preds_per_class + bbox_idx_offset) : beforeNMS_index_array[ref_item_idx];
+            } while (ref_item_idx < max_read_idx && !kept_bboxinfo_flag[ref_item_idx - offset]);
+
+            // Move to next valid point
+            if (ref_item_idx < max_read_idx)
+            {
+                ref_bbox_idx = share_location
+                    ? (beforeNMS_index_array[ref_item_idx] % num_preds_per_class + bbox_idx_offset)
+                    : beforeNMS_index_array[ref_item_idx];
+            }
         }
 
         // store data

        {
            const int cur_idx = threadIdx.x + blockDim.x * t;
            const int item_idx = offset + cur_idx;

            if (item_idx < max_idx)
            {
                loc_bboxIndex[t] = beforeNMS_index_array[item_idx];

                if (loc_bboxIndex[t] != -1)
                {


    for (int i = 0; i < num; i++)
    {
        int32_t const offset = i * num_classes * num_preds_per_class + blockIdx.x * num_preds_per_class;
        // Should not write data beyond [offset, top_k).
        int32_t const max_idx = offset + top_k;
        // Should not read beyond [offset, num_preds_per_class).
        int32_t const max_read_idx = offset + min(top_k, num_preds_per_class);
        int32_t const bbox_idx_offset = i * num_preds_per_class * (share_location ? 1 : num_classes);

        // local thread data
        int loc_bboxIndex[TSIZE];
        Bbox<T_BBOX> loc_bbox[TSIZE];

        // initialize Bbox, Bboxinfo, kept_bboxinfo_flag
        // Eliminate shared memory RAW hazard
        __syncthreads();
#pragma unroll
