 // reclaim all outstanding global memory slots since it is
 // likely we have reached the end of the kernel.
 EXTERN void __kmpc_data_sharing_pop_stack(void *FrameStart) {
-  if (IsMasterThread()) {
+  if (IsWarpMasterActiveThread()) {
     unsigned WID = getWarpId();
 
     __kmpc_data_sharing_slot *&SlotP = DataSharingState.SlotPtr[WID];
     void *&StackP = DataSharingState.StackPtr[WID];
+
+    // Pop current frame from slot.
+    StackP = FrameStart;
 
     // If we try to pop the last frame of the current slot we need to
     // move to the previous slot if there is one.

__device__ static unsigned getThreadId() { return threadIdx.x; }
// Warp ID in the CUDA block
__device__ static unsigned getWarpId() { return threadIdx.x / WARPSIZE; }

// The CUDA thread ID of the master thread.
__device__ static unsigned getMasterThreadId() {

      // Reset stack pointer to the requested address.
      StackP = (void *)RequestedEndAddress;
    }
  }

  __threadfence_block();

  // Compute the start address of the frame of each thread in the warp.
  uintptr_t FrameStartAddress = (uintptr_t)FrameP;
  FrameStartAddress += (uintptr_t) (getLaneId() * DataSize);
  return (void *)FrameStartAddress;
}

// Pop the stack and free any memory which can be reclaimed.
