From 3043fb56751f689cd7009630abf0da5b0bcfc7e7 Mon Sep 17 00:00:00 2001
From: Brian Sumner <brian.sumner@amd.com>
Date: Tue, 7 May 2024 14:50:39 -0700
Subject: [PATCH] Remove old convergence workaround

Change-Id: I996e20e4e189385fe998a4628b35163c0da559ae
---
 amd/device-libs/ockl/src/hostcall_impl.cl | 9 ---------
 1 file changed, 9 deletions(-)

diff --git a/amd/device-libs/ockl/src/hostcall_impl.cl b/amd/device-libs/ockl/src/hostcall_impl.cl
index 1a70538fda38a1..325ee8e45c4145 100644
--- a/amd/device-libs/ockl/src/hostcall_impl.cl
+++ b/amd/device-libs/ockl/src/hostcall_impl.cl
@@ -96,14 +96,6 @@ set_ready_flag(uint control)
                              CONTROL_WIDTH_READY_FLAG, 1);
 }
 
-static uint
-optimizationBarrierHack(uint in_val)
-{
-    uint out_val;
-    __asm__ volatile("" : "=v"(out_val) : "0"(in_val));
-    return out_val;
-}
-
 static ulong
 pop(__global ulong *top, __global buffer_t *buffer)
 {
@@ -290,7 +282,6 @@ __ockl_hostcall_internal(void *_buffer, uint service_id, ulong arg0, ulong arg1,
                          ulong arg6, ulong arg7)
 {
     uint me = __ockl_lane_u32();
-    me = optimizationBarrierHack(me);
     uint low = __builtin_amdgcn_readfirstlane(me);
 
     __global buffer_t *buffer = (__global buffer_t *)_buffer;