<commit_before>#pragma once

// STL
#include <vector>
#include <numeric> // for std::iota

#ifdef __CUDACC__
// cuda::thrust
#include <thrust/remove.h>
#endif

// libcgt
#include <common/Array1DView.h>
#include <common/BasicTypes.h>
#include <math/Arithmetic.h>

// local
#include "DeviceQueue.h"

template< typename UsedListTag >
struct UsedListEntry
{
	int poolIndex;
	int isDeleted; // TODO: OPTIMIZE optimize later by packing it all into the tag
	UsedListTag tag;
};

template< typename UsedListTag >
struct KernelPool
{
	__host__
	KernelPool();

	__host__
	KernelPool( int capacity,
		int elementSizeBytes,

		KernelQueue< int > freeList,
		KernelQueue< UsedListEntry< UsedListTag > > usedList,
		KernelVector< ubyte > backingStore );

	// gives you an entry that comes with an index
	// and space for a tag
	__device__
	UsedListEntry< UsedListTag >& alloc();
	
	// marks the index at usedList()[ usedListIndex ] as deleted
	__device__
	void markAsDeleted( int usedListIndex );

	__device__
	KernelQueue< int >& freeList();

	__device__
	KernelQueue< UsedListEntry< UsedListTag > >& usedList();

	// given an index (returned from alloc()),
	// returns a pointer to the beginning of the element
	template< typename T >
	__device__
	T* getElement( int index );

	int m_capacity;
	int m_elementSizeBytes;

	KernelQueue< int > m_freeList;
	KernelQueue< UsedListEntry< UsedListTag > > m_usedList;
	KernelVector< ubyte > md_backingStore;
};

template< typename UsedListTag >
__inline__ __host__
KernelPool< UsedListTag >::KernelPool()
{

}

template< typename UsedListTag >
__inline__ __host__
KernelPool< UsedListTag >::KernelPool
(
	int capacity,
	int elementSizeBytes,

	KernelQueue< int > freeList,
	KernelQueue< UsedListEntry< UsedListTag > > usedList,
	KernelVector< ubyte > backingStore
) :

	m_capacity( capacity ),
	m_elementSizeBytes( elementSizeBytes ),
	m_freeList( freeList ),
	m_usedList( usedList ),
	md_backingStore( backingStore )

{

}

#ifdef __CUDACC__
template< typename UsedListTag >
__inline__ __device__
UsedListEntry< UsedListTag >& KernelPool< UsedListTag >::alloc()
{
	int index = m_freeList.dequeue();

	UsedListEntry< UsedListTag > entry;
	entry.poolIndex = index;
	entry.isDeleted = 0;		

	uint usedListIndex = m_usedList.enqueue( entry );
	return m_usedList.ringBuffer()[ usedListIndex ];
}
#endif

template< typename UsedListTag >
__inline__ __device__
void KernelPool< UsedListTag >::markAsDeleted( int usedListIndex )
{
	UsedListEntry< UsedListTag >& entry = m_usedList.ringBuffer()[ usedListIndex ];
	entry.isDeleted = 1;
}

template< typename UsedListTag >
__inline__ __device__<commit_msg>Added free to DevicePool, for when we have the pool index directly.<commit_after>#pragma once

// STL
#include <vector>
#include <numeric> // for std::iota

#ifdef __CUDACC__
// cuda::thrust
#include <thrust/remove.h>
#endif

// libcgt
#include <common/Array1DView.h>
#include <common/BasicTypes.h>
#include <math/Arithmetic.h>

// local
#include "DeviceQueue.h"

template< typename UsedListTag >
struct UsedListEntry
{
	int poolIndex;
	int isDeleted; // TODO: OPTIMIZE optimize later by packing it all into the tag
	UsedListTag tag;
};

template< typename UsedListTag >
struct KernelPool
{
	__host__
	KernelPool();

	__host__
	KernelPool( int capacity,
		int elementSizeBytes,

		KernelQueue< int > freeList,
		KernelQueue< UsedListEntry< UsedListTag > > usedList,
		KernelVector< ubyte > backingStore );

	// gives you an entry that comes with an index
	// and space for a tag
	__device__
	UsedListEntry< UsedListTag >& alloc();

	// returns an index back to the pool
	__device__
	void free( int index );
	
	// marks the index at usedList()[ usedListIndex ] as deleted
	__device__
	void markAsDeleted( int usedListIndex );

	__device__
	KernelQueue< int >& freeList();

	__device__
	KernelQueue< UsedListEntry< UsedListTag > >& usedList();

	// given an index (returned from alloc()),
	// returns a pointer to the beginning of the element
	template< typename T >
	__device__
	T* getElement( int index );

	int m_capacity;
	int m_elementSizeBytes;

	KernelQueue< int > m_freeList;
	KernelQueue< UsedListEntry< UsedListTag > > m_usedList;
	KernelVector< ubyte > md_backingStore;
};

template< typename UsedListTag >
__inline__ __host__
KernelPool< UsedListTag >::KernelPool()
{

}

template< typename UsedListTag >
__inline__ __host__
KernelPool< UsedListTag >::KernelPool
(
	int capacity,
	int elementSizeBytes,

	KernelQueue< int > freeList,
	KernelQueue< UsedListEntry< UsedListTag > > usedList,
	KernelVector< ubyte > backingStore
) :

	m_capacity( capacity ),
	m_elementSizeBytes( elementSizeBytes ),
	m_freeList( freeList ),
	m_usedList( usedList ),
	md_backingStore( backingStore )

{

}

#ifdef __CUDACC__
template< typename UsedListTag >
__inline__ __device__
UsedListEntry< UsedListTag >& KernelPool< UsedListTag >::alloc()
{
	int index = m_freeList.dequeue();

	UsedListEntry< UsedListTag > entry;
	entry.poolIndex = index;
	entry.isDeleted = 0;		

	uint usedListIndex = m_usedList.enqueue( entry );
	return m_usedList.ringBuffer()[ usedListIndex ];
}

template< typename UsedListTag >
__inline__ __device__
void KernelPool< UsedListTag >::free( int index )
{
	m_freeList.enqueue( index );
}
#endif

template< typename UsedListTag >
__inline__ __device__
void KernelPool< UsedListTag >::markAsDeleted( int usedListIndex )
{
	UsedListEntry< UsedListTag >& entry = m_usedList.ringBuffer()[ usedListIndex ];
	entry.isDeleted = 1;
}

template< typename UsedListTag >
__inline__ __device__
{'fraction_non_alphanumeric': 0.09692028985507246, 'fraction_numerical': 0.0010869565217391304, 'mean_word_length': 5.001086956521739, 'pattern_counts': {'":': 0, '<': 79, '<?xml version=': 0, '>': 79, 'https://': 0, 'lorem ipsum': 0, 'www.': 0, 'xml': 0}, 'pii_count': 0, 'substrings_counts': 0, 'word_list_counts': {'cursed_substrings.json': 0, 'profanity_word_list.json': 0, 'sexual_word_list.json': 0, 'zh_pornsignals.json': 0}}
{'dir': 'git-commits-cleaned', 'id': '7215664', 'n_tokens_mistral': 1875, 'n_tokens_neox': 1811, 'n_words': 515}