__kernel void slotCompactor(
	__global uchar* assembly,
	uint numSlots,
	uint slotStride,
	uint slotSize,
	uint firstSlotOffset,
	uint nSucceeded)
{
	// Sequential processing: single work item processes all slots
	// Compact non-dummy slots to the beginning (lowest addresses)
	// Dummy slots will remain at the end (highest addresses)
	// Optimizations:
	// 1. For each dummy, find rightmost non-dummy and copy it there
	// 2. Exit early once we've seen nSucceeded non-dummy slots
	// 3. Exit early once we've moved nFailed dummy slots

	uint nFailed = numSlots - nSucceeded;  // Calculate number of failed slots
	uint nonDummiesSeen = 0;  // Track how many non-dummy slots we've seen
	uint dummiesMoved = 0;     // Track how many dummy slots we've moved

	// Initialize rightmostNonDummy to start from the end
	// We'll decrement it each time we use it to avoid re-selecting the same slot
	uint rightmostNonDummy = numSlots - 1;

	// Process slots from beginning to end
	for (uint i = 0; i < numSlots; ++i)
	{
		// Optimization 2: Exit early once we've seen nSucceeded non-dummy slots
		if (nonDummiesSeen >= nSucceeded) { break; }
		// Optimization 3: Exit early once we've moved nFailed dummy slots
		if (dummiesMoved >= nFailed) { break; }

		// Calculate slot address
		__global uchar* slotAddr = assembly + firstSlotOffset
			+ (i * slotStride);

		// Check if slot is dummy: first 4 bytes should all be 0xFF
		bool isDummy = (slotAddr[0] == 0xFF) && (slotAddr[1] == 0xFF)
			&& (slotAddr[2] == 0xFF) && (slotAddr[3] == 0xFF);

		// Early continue for non-dummy slots (already in the right place)
		if (!isDummy)
		{
			++nonDummiesSeen;
			continue;
		}

		// Optimization 1: Find rightmost non-dummy slot starting from where we left off
		// Search backwards from rightmostNonDummy until we find a non-dummy slot
		// or reach the current position
		bool foundNonDummy = false;
		for (int j = (int)rightmostNonDummy; j > (int)i; --j)
		{
			__global uchar* checkSlotAddr = assembly + firstSlotOffset
				+ (j * slotStride);

			bool checkIsDummy = (checkSlotAddr[0] == 0xFF)
				&& (checkSlotAddr[1] == 0xFF) && (checkSlotAddr[2] == 0xFF)
				&& (checkSlotAddr[3] == 0xFF);

			// Early continue for dummy slots
			if (checkIsDummy)
				{ continue; }

			// Found a non-dummy slot
			rightmostNonDummy = (uint)j;
			foundNonDummy = true;
			break;
		}

		// If we found a non-dummy slot to the right, copy it here
		if (foundNonDummy)
		{
			__global uchar* srcAddr = assembly + firstSlotOffset
				+ (rightmostNonDummy * slotStride);

			// Copy slot data (byte-by-byte copy)
			for (uint j = 0; j < slotSize; ++j) {
				slotAddr[j] = srcAddr[j];
			}

			// Mark the source slot as dummy (move it to the end)
			for (uint j = 0; j < 4; ++j) {
				srcAddr[j] = 0xFF;
			}

			// Decrement rightmostNonDummy to avoid re-selecting the same slot
			--rightmostNonDummy;

			++dummiesMoved;
			++nonDummiesSeen;  // We just moved a non-dummy to this position
		}
	}
}