#include "ConcStadiumHash.cuh"
#include <stdint.h>

#define TICKET_SIZE 4
#define NUM_POS_BITS 3 // log(32/TICKET_SIZE)

// Generic Hash Function
template<typename keyT, typename valueT>
template <uint nRounds, uint rShift, uint mulOp>
__device__ uint ConcStadiumHash<keyT, valueT>::hash(const keyT key)
{
    keyT x = key;
    for (uint j = nRounds; j > 0; --j) {
        x = ((x >> rShift) ^ x) * mulOp + j;
    }
    return (uint)x;
}

// First Hash 
template<typename keyT, typename valueT>
__device__ uint ConcStadiumHash<keyT, valueT>::initHash(keyT key)
{
	uint hashed = hash<sizeof(keyT), 8, 0xE9D58A6B>(key);
	return __umulhi(hashed, tableSize);
}

// Second Hash
template<typename keyT, typename valueT>
__device__ uint ConcStadiumHash<keyT, valueT>::rehash(uint hashed, const keyT key)
{
	uint h_2 = hash<sizeof(keyT), 8, 0x6E5B9D8A>(key);
	uint dh = hashed + 1 + __umulhi(h_2, tableSize - 1);
	return (dh >= tableSize) ? (dh - tableSize) : dh;
}

// Info Hash
template<typename keyT, typename valueT>
__device__ uint ConcStadiumHash<keyT, valueT>::infoHash(keyT key)
{
	return hash<sizeof(keyT), 8, 0xCA34BE7D>(key) >> (32 - (TICKET_SIZE - 2)); // mulOp was chosen randomly
}

// First Array Hash
template<typename keyT, typename valueT>
__device__ uint ConcStadiumHash<keyT, valueT>::arrayHash(keyT* arr, uint length)
{
    uint hashed = 0;
    for (uint i = 0; i < length; i++)
    {
        hashed = hashed * 31 + hash<sizeof(keyT), 8, 0xE9D58A6B>(arr[i]);
    }

    return __umulhi(hashed, tableSize);
}

// Second Array Hash
template<typename keyT, typename valueT>
__device__ uint ConcStadiumHash<keyT, valueT>::arrayRehash(uint hashed, keyT* arr, uint length)
{
    uint h_2 = 0;
    for (uint i = 0; i < length; i++)
    {
        h_2 = h_2 * 31 + hash<sizeof(keyT), 8, 0x6E5B9D8A>(arr[i]);
    }

    uint dh = hashed + 1 + __umulhi(h_2, tableSize - 1);
    return (dh >= tableSize) ? (dh - tableSize) : dh;
}

// Info Array Hash
template<typename keyT, typename valueT>
__device__ uint ConcStadiumHash<keyT, valueT>::arrayInfoHash(keyT* arr, uint length)
{
    uint hashed = 0;
    for (uint i = 0; i < length; i++)
    {
        hashed = hashed * 31 + hash<sizeof(keyT), 8, 0xCA34BE7D>(arr[i]);
    }

    return hashed >> (32 - (TICKET_SIZE - 2));
}

// Allocation of the ticket board
template<typename keyT, typename valueT>
__host__ void ConcStadiumHash<keyT, valueT>::allocTicketBoard()
{
    const int ticketBoardSize = tableSize / (32 / TICKET_SIZE) + 1;
    cudaError_t cudaStatus;

    // Allocate memory for ticket board
    cudaStatus = cudaMalloc((void**)&ticketBoard, ticketBoardSize * sizeof(uint));
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "Ticket Board: cudaMalloc failed!\n");
    }

    // Initialize ticket board with 1s
    clearTicketBoard();
}

// Clearing of the ticket board
template<typename keyT, typename valueT>
__host__ void ConcStadiumHash<keyT, valueT>::clearTicketBoard()
{
    const int ticketBoardSize = tableSize / (32 / TICKET_SIZE) + 1;
    cudaError_t cudaStatus;

    cudaStatus = cudaMemset((void*)ticketBoard, 0xFF, ticketBoardSize * sizeof(uint));
    if (cudaStatus != cudaSuccess) {
        fprintf(stderr, "Ticket Board: cudaMemset failed!\n");
    }
}

// Deletion of the ticket board
template<typename keyT, typename valueT>
__host__ void ConcStadiumHash<keyT, valueT>::freeTicketBoard()
{
    cudaFree((void*)ticketBoard);
}

// Creates a mask containing info starting at infoStart and otherwise 1s
template<typename keyT, typename valueT>
__device__ uint ConcStadiumHash<keyT, valueT>::prepareTicket(uint info, uint infoStart)
{
    uint mask = (1 << (TICKET_SIZE - 2)) - 1;
    mask = ~(mask << infoStart);

    uint result = mask | (info << infoStart);

    return result;
}

// Extracts the info from a ticket
template<typename keyT, typename valueT>
__device__ uint ConcStadiumHash<keyT, valueT>::extractInfo(uint ticket, uint infoStart)
{
    uint mask = (1 << (TICKET_SIZE - 2)) - 1;
    mask = mask << infoStart;
    uint result = (mask & ticket) >> infoStart;

    return result;
}

// Calculates Ticket Board Index from the hashed key
template<typename keyT, typename valueT>
__device__ uint ConcStadiumHash<keyT, valueT>::getTbIndex(uint hashed)
{
    return hashed >> NUM_POS_BITS;
}

// Calculates the Position of the ticket in the Ticket Board Entry from the hashed key
template<typename keyT, typename valueT>
__device__ uint ConcStadiumHash<keyT, valueT>::getPosInInt(uint hashed)
{
    uint mask = (1 << NUM_POS_BITS) - 1;

    return (hashed & mask) << (5 - NUM_POS_BITS);
}

// Entry Reservation
template<typename keyT, typename valueT>
__device__ uint ConcStadiumHash<keyT, valueT>::tryBookASeat(uint tbIndex, uint posInInt)
{
    uint availability = 1 << posInInt;
    uint access = 1 << (posInInt + 1);
    uint permit = availability | access;

    // check if availability bit is set
    if (availability & ticketBoard[tbIndex])
    {
        // try to reserve ticket
        uint auth = atomicAnd((uint*)(ticketBoard + tbIndex), ~permit);

        if (auth & availability)
        {
            // ticket was booked succesfully
            return (~0);
        }
        else if (auth & permit == 0)
        {
            // ticket was booked first by another thread
            return 0;
        }
        else
        {
            // ticket is accessed by another thread
            // set access bit (necessary if bit was unset without obtaining the ticket)
            atomicOr((uint*)(ticketBoard + tbIndex), access);
            return 0;
        }
    }
    else
    {
        return 0;
    }
}



// Entry Search
template<typename keyT, typename valueT>
__device__ uint ConcStadiumHash<keyT, valueT>::tryFindTheSeat(keyT* key, uint hashed, uint tbIndex, uint posInInt, uint info, keyT* keys, uint keyLength)
{
    uint permit = 1 << posInInt;
    if (permit & ticketBoard[tbIndex]) {
        return KEY_NOT_INSIDE_TABLE;
    }
    else {

        // wait until ticket can be accessed
        while (isAccessed(tbIndex, posInInt)) {};

        // get and compare info from ticket
        uint retrievedInfo = extractInfo(ticketBoard[tbIndex], posInInt + 2);
        if (info != retrievedInfo) return 0;

        // compare keys
        for (int i = 0; i < keyLength; i++)
        {
            if (keys[i] != key[i]) return 0;
        }
        
        return (~0);
    }
}

// Entry Search
template<typename keyT, typename valueT>
__device__ uint ConcStadiumHash<keyT, valueT>::tryFindTheSeat(keyT key, uint hashed, uint tbIndex, uint posInInt, uint info, keyT* keys)
{
    uint permit = 1 << posInInt;
    if (permit & ticketBoard[tbIndex]) {
        return KEY_NOT_INSIDE_TABLE;
    }
    else {
        // wait until ticket can be accessed
        while (isAccessed(tbIndex, posInInt)) {};

        // get info from ticket
        uint retrievedInfo = extractInfo(ticketBoard[tbIndex], posInInt + 2); // ACCESS CHANGE
        if (info != retrievedInfo) return 0;

        // compare keys
        if (key != keys[hashed]) return 0;

        return (~0);
    }
}

// Inserts info into the ticket board
template<typename keyT, typename valueT>
__device__ void ConcStadiumHash<keyT, valueT>::insertTicketInfo(uint info, uint tbIndex, uint posInInt)
{
    uint prepTicket = prepareTicket(info, posInInt + 2);
    atomicAnd((uint*)ticketBoard + tbIndex, prepTicket);
}

// Checks if a ticket is accessed by a thread
template<typename keyT, typename valueT>
__device__ uint ConcStadiumHash<keyT, valueT>::isAccessed(uint tbIndex, uint posInInt)
{
    uint access = (1 << (posInInt + 1));

    return (ticketBoard[tbIndex] & access) ? (0) : (~0);
}

// Unlocks access to the ticket
template<typename keyT, typename valueT>
__device__ void ConcStadiumHash<keyT, valueT>::unlockAccess(uint tbIndex, uint posInInt)
{
    __threadfence();
    uint access = (1 << (posInInt + 1));
    atomicOr((uint*)ticketBoard + tbIndex, access);
}

// Constructor
template<typename keyT, typename valueT>
ConcStadiumHash<keyT, valueT>::ConcStadiumHash(uint tableSize)
{
    this->tableSize = tableSize;
}

// Destructor
template<typename keyT, typename valueT>
ConcStadiumHash<keyT, valueT>::~ConcStadiumHash()
{
    freeTicketBoard();
}

// Inserts a given key value pair into the table
template<typename keyT, typename valueT>
__device__ void ConcStadiumHash<keyT, valueT>::insert(keyT* key, valueT value, keyT* keys, valueT* values, uint keyLength)
{
    // Initial Hash
    uint hashed = arrayHash(key, keyLength);

    // Get ticket board index and the position in the integer from hashed value
    uint tbIndex = getTbIndex(hashed);
    uint posInInt = getPosInInt(hashed);

    uint tryCounter = 0;
    uint gotSeat;

    // loop until a bucket is succesfully reserved
    do
    {
        // try to reserve a free bucket
        gotSeat = tryBookASeat(tbIndex, posInInt);

        // check if bucket was already reserved
        if (!gotSeat)
        {
            // rehash the key with the second hash function
            hashed = arrayRehash(hashed, key, keyLength);

            // Get ticket board index and the position in the integer from hashed value
            tbIndex = getTbIndex(hashed);
            posInInt = getPosInInt(hashed);
        }

        // if the number of tries is bigger than the table size, the table is already full
        assert(++tryCounter < tableSize);
        //throw std::runtime_error("INSERT FAILED - Table is full!");
    } while (!gotSeat);

    // Generate Info from key
    uint info = arrayInfoHash(key, keyLength);

    // Insert info into the reserved bucket
    insertTicketInfo(info, tbIndex, posInInt);

    // Insert key value pair into the table
    for (int i = 0; i < keyLength; i++)
    {
        keys[hashed * keyLength + i] = key[i];
    }

    values[hashed] = value;

    unlockAccess(tbIndex, posInInt);
}

// searches fo a key in the table
template<typename keyT, typename valueT>
__device__ uint ConcStadiumHash<keyT, valueT>::find(keyT* key, keyT* keys, uint keyLength)
{
    // Intial Hash
    uint hashed = arrayHash(key, keyLength); // initHash(key);
    // printf("Key: %u Hashed: %u\n", key, hashed);
    // Generate Info from key
    uint info = arrayInfoHash(key, keyLength);

    uint tbIndex = getTbIndex(hashed);
    uint posInInt = getPosInInt(hashed);

    uint tryCounter = 0;
    uint seatFound;

    // loop until the entry with the searched key is found
    do {
        // search for the key in the ticket board and table
        seatFound = tryFindTheSeat(key, hashed, tbIndex, posInInt, info, keys + (hashed*keyLength), keyLength);

        // if no entry entry was found, rehash the key
        if (!seatFound)
        {
            hashed = arrayRehash(hashed, key, keyLength);
            tbIndex = getTbIndex(hashed);
            posInInt = getPosInInt(hashed);
        }

        // if the number of tries is bigger than the table size, the key isn't included in the table
        if (++tryCounter == tableSize)
            seatFound = KEY_NOT_INSIDE_TABLE;

    } while (!seatFound && seatFound != KEY_NOT_INSIDE_TABLE);


    if (seatFound == KEY_NOT_INSIDE_TABLE)
    {
        // if key wasn't found in the table return tableSize as index
        hashed = tableSize;
    }

    return hashed;
}

template class ConcStadiumHash<unsigned int, unsigned int>;
template class ConcStadiumHash<uint64_t, uint32_t>;

	.file	"tmpxft_00282fbe_00000000-6_ConcStadiumHash.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2053:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2053:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.section	.text._ZN15ConcStadiumHashIjjE8initHashEj,"axG",@progbits,_ZN15ConcStadiumHashIjjE8initHashEj,comdat
	.align 2
	.weak	_ZN15ConcStadiumHashIjjE8initHashEj
	.type	_ZN15ConcStadiumHashIjjE8initHashEj, @function
_ZN15ConcStadiumHashIjjE8initHashEj:
.LFB2122:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2122:
	.size	_ZN15ConcStadiumHashIjjE8initHashEj, .-_ZN15ConcStadiumHashIjjE8initHashEj
	.section	.text._ZN15ConcStadiumHashIjjE6rehashEjj,"axG",@progbits,_ZN15ConcStadiumHashIjjE6rehashEjj,comdat
	.align 2
	.weak	_ZN15ConcStadiumHashIjjE6rehashEjj
	.type	_ZN15ConcStadiumHashIjjE6rehashEjj, @function
_ZN15ConcStadiumHashIjjE6rehashEjj:
.LFB2123:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2123:
	.size	_ZN15ConcStadiumHashIjjE6rehashEjj, .-_ZN15ConcStadiumHashIjjE6rehashEjj
	.section	.text._ZN15ConcStadiumHashIjjE8infoHashEj,"axG",@progbits,_ZN15ConcStadiumHashIjjE8infoHashEj,comdat
	.align 2
	.weak	_ZN15ConcStadiumHashIjjE8infoHashEj
	.type	_ZN15ConcStadiumHashIjjE8infoHashEj, @function
_ZN15ConcStadiumHashIjjE8infoHashEj:
.LFB2124:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2124:
	.size	_ZN15ConcStadiumHashIjjE8infoHashEj, .-_ZN15ConcStadiumHashIjjE8infoHashEj
	.section	.text._ZN15ConcStadiumHashIjjE9arrayHashEPjj,"axG",@progbits,_ZN15ConcStadiumHashIjjE9arrayHashEPjj,comdat
	.align 2
	.weak	_ZN15ConcStadiumHashIjjE9arrayHashEPjj
	.type	_ZN15ConcStadiumHashIjjE9arrayHashEPjj, @function
_ZN15ConcStadiumHashIjjE9arrayHashEPjj:
.LFB2125:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2125:
	.size	_ZN15ConcStadiumHashIjjE9arrayHashEPjj, .-_ZN15ConcStadiumHashIjjE9arrayHashEPjj
	.section	.text._ZN15ConcStadiumHashIjjE11arrayRehashEjPjj,"axG",@progbits,_ZN15ConcStadiumHashIjjE11arrayRehashEjPjj,comdat
	.align 2
	.weak	_ZN15ConcStadiumHashIjjE11arrayRehashEjPjj
	.type	_ZN15ConcStadiumHashIjjE11arrayRehashEjPjj, @function
_ZN15ConcStadiumHashIjjE11arrayRehashEjPjj:
.LFB2126:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2126:
	.size	_ZN15ConcStadiumHashIjjE11arrayRehashEjPjj, .-_ZN15ConcStadiumHashIjjE11arrayRehashEjPjj
	.section	.text._ZN15ConcStadiumHashIjjE13arrayInfoHashEPjj,"axG",@progbits,_ZN15ConcStadiumHashIjjE13arrayInfoHashEPjj,comdat
	.align 2
	.weak	_ZN15ConcStadiumHashIjjE13arrayInfoHashEPjj
	.type	_ZN15ConcStadiumHashIjjE13arrayInfoHashEPjj, @function
_ZN15ConcStadiumHashIjjE13arrayInfoHashEPjj:
.LFB2127:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2127:
	.size	_ZN15ConcStadiumHashIjjE13arrayInfoHashEPjj, .-_ZN15ConcStadiumHashIjjE13arrayInfoHashEPjj
	.section	.rodata._ZN15ConcStadiumHashIjjE16clearTicketBoardEv.str1.1,"aMS",@progbits,1
.LC0:
	.string	"Ticket Board: cudaMemset failed!\n"
	.section	.text._ZN15ConcStadiumHashIjjE16clearTicketBoardEv,"axG",@progbits,_ZN15ConcStadiumHashIjjE16clearTicketBoardEv,comdat
	.align 2
	.weak	_ZN15ConcStadiumHashIjjE16clearTicketBoardEv
	.type	_ZN15ConcStadiumHashIjjE16clearTicketBoardEv, @function
_ZN15ConcStadiumHashIjjE16clearTicketBoardEv:
.LFB2129:
	.cfi_startproc
	endbr64
	pushq	%rsi
	.cfi_def_cfa_offset 16
	movl	8(%rdi), %eax
	movl	$255, %esi
	movq	(%rdi), %rdi
	shrl	$3, %eax
	incl	%eax
	leal	0(,%rax,4), %edx
	call	cudaMemset@PLT
	testl	%eax, %eax
	je	.L14
	movq	stderr(%rip), %rdi
	leaq	.LC0(%rip), %rdx
	movl	$2, %esi
	xorl	%eax, %eax
	popq	%rcx
	.cfi_remember_state
	.cfi_def_cfa_offset 8
	jmp	__fprintf_chk@PLT
.L14:
	.cfi_restore_state
	popq	%rax
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2129:
	.size	_ZN15ConcStadiumHashIjjE16clearTicketBoardEv, .-_ZN15ConcStadiumHashIjjE16clearTicketBoardEv
	.section	.rodata._ZN15ConcStadiumHashIjjE16allocTicketBoardEv.str1.1,"aMS",@progbits,1
.LC1:
	.string	"Ticket Board: cudaMalloc failed!\n"
	.section	.text._ZN15ConcStadiumHashIjjE16allocTicketBoardEv,"axG",@progbits,_ZN15ConcStadiumHashIjjE16allocTicketBoardEv,comdat
	.align 2
	.weak	_ZN15ConcStadiumHashIjjE16allocTicketBoardEv
	.type	_ZN15ConcStadiumHashIjjE16allocTicketBoardEv, @function
_ZN15ConcStadiumHashIjjE16allocTicketBoardEv:
.LFB2128:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movl	8(%rdi), %eax
	movq	%rdi, %rbx
	shrl	$3, %eax
	incl	%eax
	leal	0(,%rax,4), %esi
	call	cudaMalloc@PLT
	testl	%eax, %eax
	je	.L19
	movq	stderr(%rip), %rdi
	leaq	.LC1(%rip), %rdx
	movl	$2, %esi
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
.L19:
	movq	%rbx, %rdi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	_ZN15ConcStadiumHashIjjE16clearTicketBoardEv
	.cfi_endproc
.LFE2128:
	.size	_ZN15ConcStadiumHashIjjE16allocTicketBoardEv, .-_ZN15ConcStadiumHashIjjE16allocTicketBoardEv
	.section	.text._ZN15ConcStadiumHashIjjE15freeTicketBoardEv,"axG",@progbits,_ZN15ConcStadiumHashIjjE15freeTicketBoardEv,comdat
	.align 2
	.weak	_ZN15ConcStadiumHashIjjE15freeTicketBoardEv
	.type	_ZN15ConcStadiumHashIjjE15freeTicketBoardEv, @function
_ZN15ConcStadiumHashIjjE15freeTicketBoardEv:
.LFB2130:
	.cfi_startproc
	endbr64
	movq	(%rdi), %rdi
	jmp	cudaFree@PLT
	.cfi_endproc
.LFE2130:
	.size	_ZN15ConcStadiumHashIjjE15freeTicketBoardEv, .-_ZN15ConcStadiumHashIjjE15freeTicketBoardEv
	.section	.text._ZN15ConcStadiumHashIjjE13prepareTicketEjj,"axG",@progbits,_ZN15ConcStadiumHashIjjE13prepareTicketEjj,comdat
	.align 2
	.weak	_ZN15ConcStadiumHashIjjE13prepareTicketEjj
	.type	_ZN15ConcStadiumHashIjjE13prepareTicketEjj, @function
_ZN15ConcStadiumHashIjjE13prepareTicketEjj:
.LFB2131:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2131:
	.size	_ZN15ConcStadiumHashIjjE13prepareTicketEjj, .-_ZN15ConcStadiumHashIjjE13prepareTicketEjj
	.section	.text._ZN15ConcStadiumHashIjjE11extractInfoEjj,"axG",@progbits,_ZN15ConcStadiumHashIjjE11extractInfoEjj,comdat
	.align 2
	.weak	_ZN15ConcStadiumHashIjjE11extractInfoEjj
	.type	_ZN15ConcStadiumHashIjjE11extractInfoEjj, @function
_ZN15ConcStadiumHashIjjE11extractInfoEjj:
.LFB2132:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2132:
	.size	_ZN15ConcStadiumHashIjjE11extractInfoEjj, .-_ZN15ConcStadiumHashIjjE11extractInfoEjj
	.section	.text._ZN15ConcStadiumHashIjjE10getTbIndexEj,"axG",@progbits,_ZN15ConcStadiumHashIjjE10getTbIndexEj,comdat
	.align 2
	.weak	_ZN15ConcStadiumHashIjjE10getTbIndexEj
	.type	_ZN15ConcStadiumHashIjjE10getTbIndexEj, @function
_ZN15ConcStadiumHashIjjE10getTbIndexEj:
.LFB2133:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2133:
	.size	_ZN15ConcStadiumHashIjjE10getTbIndexEj, .-_ZN15ConcStadiumHashIjjE10getTbIndexEj
	.section	.text._ZN15ConcStadiumHashIjjE11getPosInIntEj,"axG",@progbits,_ZN15ConcStadiumHashIjjE11getPosInIntEj,comdat
	.align 2
	.weak	_ZN15ConcStadiumHashIjjE11getPosInIntEj
	.type	_ZN15ConcStadiumHashIjjE11getPosInIntEj, @function
_ZN15ConcStadiumHashIjjE11getPosInIntEj:
.LFB2134:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2134:
	.size	_ZN15ConcStadiumHashIjjE11getPosInIntEj, .-_ZN15ConcStadiumHashIjjE11getPosInIntEj
	.section	.text._ZN15ConcStadiumHashIjjE12tryBookASeatEjj,"axG",@progbits,_ZN15ConcStadiumHashIjjE12tryBookASeatEjj,comdat
	.align 2
	.weak	_ZN15ConcStadiumHashIjjE12tryBookASeatEjj
	.type	_ZN15ConcStadiumHashIjjE12tryBookASeatEjj, @function
_ZN15ConcStadiumHashIjjE12tryBookASeatEjj:
.LFB2135:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2135:
	.size	_ZN15ConcStadiumHashIjjE12tryBookASeatEjj, .-_ZN15ConcStadiumHashIjjE12tryBookASeatEjj
	.section	.text._ZN15ConcStadiumHashIjjE14tryFindTheSeatEPjjjjjS1_j,"axG",@progbits,_ZN15ConcStadiumHashIjjE14tryFindTheSeatEPjjjjjS1_j,comdat
	.align 2
	.weak	_ZN15ConcStadiumHashIjjE14tryFindTheSeatEPjjjjjS1_j
	.type	_ZN15ConcStadiumHashIjjE14tryFindTheSeatEPjjjjjS1_j, @function
_ZN15ConcStadiumHashIjjE14tryFindTheSeatEPjjjjjS1_j:
.LFB2136:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2136:
	.size	_ZN15ConcStadiumHashIjjE14tryFindTheSeatEPjjjjjS1_j, .-_ZN15ConcStadiumHashIjjE14tryFindTheSeatEPjjjjjS1_j
	.section	.text._ZN15ConcStadiumHashIjjE14tryFindTheSeatEjjjjjPj,"axG",@progbits,_ZN15ConcStadiumHashIjjE14tryFindTheSeatEjjjjjPj,comdat
	.align 2
	.weak	_ZN15ConcStadiumHashIjjE14tryFindTheSeatEjjjjjPj
	.type	_ZN15ConcStadiumHashIjjE14tryFindTheSeatEjjjjjPj, @function
_ZN15ConcStadiumHashIjjE14tryFindTheSeatEjjjjjPj:
.LFB2137:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2137:
	.size	_ZN15ConcStadiumHashIjjE14tryFindTheSeatEjjjjjPj, .-_ZN15ConcStadiumHashIjjE14tryFindTheSeatEjjjjjPj
	.section	.text._ZN15ConcStadiumHashIjjE16insertTicketInfoEjjj,"axG",@progbits,_ZN15ConcStadiumHashIjjE16insertTicketInfoEjjj,comdat
	.align 2
	.weak	_ZN15ConcStadiumHashIjjE16insertTicketInfoEjjj
	.type	_ZN15ConcStadiumHashIjjE16insertTicketInfoEjjj, @function
_ZN15ConcStadiumHashIjjE16insertTicketInfoEjjj:
.LFB2138:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2138:
	.size	_ZN15ConcStadiumHashIjjE16insertTicketInfoEjjj, .-_ZN15ConcStadiumHashIjjE16insertTicketInfoEjjj
	.section	.text._ZN15ConcStadiumHashIjjE10isAccessedEjj,"axG",@progbits,_ZN15ConcStadiumHashIjjE10isAccessedEjj,comdat
	.align 2
	.weak	_ZN15ConcStadiumHashIjjE10isAccessedEjj
	.type	_ZN15ConcStadiumHashIjjE10isAccessedEjj, @function
_ZN15ConcStadiumHashIjjE10isAccessedEjj:
.LFB2139:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2139:
	.size	_ZN15ConcStadiumHashIjjE10isAccessedEjj, .-_ZN15ConcStadiumHashIjjE10isAccessedEjj
	.section	.text._ZN15ConcStadiumHashIjjE12unlockAccessEjj,"axG",@progbits,_ZN15ConcStadiumHashIjjE12unlockAccessEjj,comdat
	.align 2
	.weak	_ZN15ConcStadiumHashIjjE12unlockAccessEjj
	.type	_ZN15ConcStadiumHashIjjE12unlockAccessEjj, @function
_ZN15ConcStadiumHashIjjE12unlockAccessEjj:
.LFB2140:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2140:
	.size	_ZN15ConcStadiumHashIjjE12unlockAccessEjj, .-_ZN15ConcStadiumHashIjjE12unlockAccessEjj
	.section	.text._ZN15ConcStadiumHashIjjEC2Ej,"axG",@progbits,_ZN15ConcStadiumHashIjjEC5Ej,comdat
	.align 2
	.weak	_ZN15ConcStadiumHashIjjEC2Ej
	.type	_ZN15ConcStadiumHashIjjEC2Ej, @function
_ZN15ConcStadiumHashIjjEC2Ej:
.LFB2142:
	.cfi_startproc
	endbr64
	movl	%esi, 8(%rdi)
	ret
	.cfi_endproc
.LFE2142:
	.size	_ZN15ConcStadiumHashIjjEC2Ej, .-_ZN15ConcStadiumHashIjjEC2Ej
	.weak	_ZN15ConcStadiumHashIjjEC1Ej
	.set	_ZN15ConcStadiumHashIjjEC1Ej,_ZN15ConcStadiumHashIjjEC2Ej
	.section	.text._ZN15ConcStadiumHashIjjED2Ev,"axG",@progbits,_ZN15ConcStadiumHashIjjED5Ev,comdat
	.align 2
	.weak	_ZN15ConcStadiumHashIjjED2Ev
	.type	_ZN15ConcStadiumHashIjjED2Ev, @function
_ZN15ConcStadiumHashIjjED2Ev:
.LFB2145:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA2145
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	_ZN15ConcStadiumHashIjjE15freeTicketBoardEv
	popq	%rdx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2145:
	.globl	__gxx_personality_v0
	.section	.gcc_except_table._ZN15ConcStadiumHashIjjED2Ev,"aG",@progbits,_ZN15ConcStadiumHashIjjED5Ev,comdat
.LLSDA2145:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE2145-.LLSDACSB2145
.LLSDACSB2145:
.LLSDACSE2145:
	.section	.text._ZN15ConcStadiumHashIjjED2Ev,"axG",@progbits,_ZN15ConcStadiumHashIjjED5Ev,comdat
	.size	_ZN15ConcStadiumHashIjjED2Ev, .-_ZN15ConcStadiumHashIjjED2Ev
	.weak	_ZN15ConcStadiumHashIjjED1Ev
	.set	_ZN15ConcStadiumHashIjjED1Ev,_ZN15ConcStadiumHashIjjED2Ev
	.section	.text._ZN15ConcStadiumHashIjjE6insertEPjjS1_S1_j,"axG",@progbits,_ZN15ConcStadiumHashIjjE6insertEPjjS1_S1_j,comdat
	.align 2
	.weak	_ZN15ConcStadiumHashIjjE6insertEPjjS1_S1_j
	.type	_ZN15ConcStadiumHashIjjE6insertEPjjS1_S1_j, @function
_ZN15ConcStadiumHashIjjE6insertEPjjS1_S1_j:
.LFB2147:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2147:
	.size	_ZN15ConcStadiumHashIjjE6insertEPjjS1_S1_j, .-_ZN15ConcStadiumHashIjjE6insertEPjjS1_S1_j
	.section	.text._ZN15ConcStadiumHashIjjE4findEPjS1_j,"axG",@progbits,_ZN15ConcStadiumHashIjjE4findEPjS1_j,comdat
	.align 2
	.weak	_ZN15ConcStadiumHashIjjE4findEPjS1_j
	.type	_ZN15ConcStadiumHashIjjE4findEPjS1_j, @function
_ZN15ConcStadiumHashIjjE4findEPjS1_j:
.LFB2148:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2148:
	.size	_ZN15ConcStadiumHashIjjE4findEPjS1_j, .-_ZN15ConcStadiumHashIjjE4findEPjS1_j
	.section	.text._ZN15ConcStadiumHashImjE8initHashEm,"axG",@progbits,_ZN15ConcStadiumHashImjE8initHashEm,comdat
	.align 2
	.weak	_ZN15ConcStadiumHashImjE8initHashEm
	.type	_ZN15ConcStadiumHashImjE8initHashEm, @function
_ZN15ConcStadiumHashImjE8initHashEm:
.LFB2149:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2149:
	.size	_ZN15ConcStadiumHashImjE8initHashEm, .-_ZN15ConcStadiumHashImjE8initHashEm
	.section	.text._ZN15ConcStadiumHashImjE6rehashEjm,"axG",@progbits,_ZN15ConcStadiumHashImjE6rehashEjm,comdat
	.align 2
	.weak	_ZN15ConcStadiumHashImjE6rehashEjm
	.type	_ZN15ConcStadiumHashImjE6rehashEjm, @function
_ZN15ConcStadiumHashImjE6rehashEjm:
.LFB2150:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2150:
	.size	_ZN15ConcStadiumHashImjE6rehashEjm, .-_ZN15ConcStadiumHashImjE6rehashEjm
	.section	.text._ZN15ConcStadiumHashImjE8infoHashEm,"axG",@progbits,_ZN15ConcStadiumHashImjE8infoHashEm,comdat
	.align 2
	.weak	_ZN15ConcStadiumHashImjE8infoHashEm
	.type	_ZN15ConcStadiumHashImjE8infoHashEm, @function
_ZN15ConcStadiumHashImjE8infoHashEm:
.LFB2151:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2151:
	.size	_ZN15ConcStadiumHashImjE8infoHashEm, .-_ZN15ConcStadiumHashImjE8infoHashEm
	.section	.text._ZN15ConcStadiumHashImjE9arrayHashEPmj,"axG",@progbits,_ZN15ConcStadiumHashImjE9arrayHashEPmj,comdat
	.align 2
	.weak	_ZN15ConcStadiumHashImjE9arrayHashEPmj
	.type	_ZN15ConcStadiumHashImjE9arrayHashEPmj, @function
_ZN15ConcStadiumHashImjE9arrayHashEPmj:
.LFB2152:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2152:
	.size	_ZN15ConcStadiumHashImjE9arrayHashEPmj, .-_ZN15ConcStadiumHashImjE9arrayHashEPmj
	.section	.text._ZN15ConcStadiumHashImjE11arrayRehashEjPmj,"axG",@progbits,_ZN15ConcStadiumHashImjE11arrayRehashEjPmj,comdat
	.align 2
	.weak	_ZN15ConcStadiumHashImjE11arrayRehashEjPmj
	.type	_ZN15ConcStadiumHashImjE11arrayRehashEjPmj, @function
_ZN15ConcStadiumHashImjE11arrayRehashEjPmj:
.LFB2153:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2153:
	.size	_ZN15ConcStadiumHashImjE11arrayRehashEjPmj, .-_ZN15ConcStadiumHashImjE11arrayRehashEjPmj
	.section	.text._ZN15ConcStadiumHashImjE13arrayInfoHashEPmj,"axG",@progbits,_ZN15ConcStadiumHashImjE13arrayInfoHashEPmj,comdat
	.align 2
	.weak	_ZN15ConcStadiumHashImjE13arrayInfoHashEPmj
	.type	_ZN15ConcStadiumHashImjE13arrayInfoHashEPmj, @function
_ZN15ConcStadiumHashImjE13arrayInfoHashEPmj:
.LFB2154:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2154:
	.size	_ZN15ConcStadiumHashImjE13arrayInfoHashEPmj, .-_ZN15ConcStadiumHashImjE13arrayInfoHashEPmj
	.section	.text._ZN15ConcStadiumHashImjE16clearTicketBoardEv,"axG",@progbits,_ZN15ConcStadiumHashImjE16clearTicketBoardEv,comdat
	.align 2
	.weak	_ZN15ConcStadiumHashImjE16clearTicketBoardEv
	.type	_ZN15ConcStadiumHashImjE16clearTicketBoardEv, @function
_ZN15ConcStadiumHashImjE16clearTicketBoardEv:
.LFB2156:
	.cfi_startproc
	endbr64
	pushq	%rsi
	.cfi_def_cfa_offset 16
	movl	8(%rdi), %eax
	movl	$255, %esi
	movq	(%rdi), %rdi
	shrl	$3, %eax
	incl	%eax
	leal	0(,%rax,4), %edx
	call	cudaMemset@PLT
	testl	%eax, %eax
	je	.L64
	movq	stderr(%rip), %rdi
	leaq	.LC0(%rip), %rdx
	movl	$2, %esi
	xorl	%eax, %eax
	popq	%rcx
	.cfi_remember_state
	.cfi_def_cfa_offset 8
	jmp	__fprintf_chk@PLT
.L64:
	.cfi_restore_state
	popq	%rax
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2156:
	.size	_ZN15ConcStadiumHashImjE16clearTicketBoardEv, .-_ZN15ConcStadiumHashImjE16clearTicketBoardEv
	.section	.text._ZN15ConcStadiumHashImjE16allocTicketBoardEv,"axG",@progbits,_ZN15ConcStadiumHashImjE16allocTicketBoardEv,comdat
	.align 2
	.weak	_ZN15ConcStadiumHashImjE16allocTicketBoardEv
	.type	_ZN15ConcStadiumHashImjE16allocTicketBoardEv, @function
_ZN15ConcStadiumHashImjE16allocTicketBoardEv:
.LFB2155:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movl	8(%rdi), %eax
	movq	%rdi, %rbx
	shrl	$3, %eax
	incl	%eax
	leal	0(,%rax,4), %esi
	call	cudaMalloc@PLT
	testl	%eax, %eax
	je	.L68
	movq	stderr(%rip), %rdi
	leaq	.LC1(%rip), %rdx
	movl	$2, %esi
	xorl	%eax, %eax
	call	__fprintf_chk@PLT
.L68:
	movq	%rbx, %rdi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	_ZN15ConcStadiumHashImjE16clearTicketBoardEv
	.cfi_endproc
.LFE2155:
	.size	_ZN15ConcStadiumHashImjE16allocTicketBoardEv, .-_ZN15ConcStadiumHashImjE16allocTicketBoardEv
	.section	.text._ZN15ConcStadiumHashImjE15freeTicketBoardEv,"axG",@progbits,_ZN15ConcStadiumHashImjE15freeTicketBoardEv,comdat
	.align 2
	.weak	_ZN15ConcStadiumHashImjE15freeTicketBoardEv
	.type	_ZN15ConcStadiumHashImjE15freeTicketBoardEv, @function
_ZN15ConcStadiumHashImjE15freeTicketBoardEv:
.LFB2157:
	.cfi_startproc
	endbr64
	movq	(%rdi), %rdi
	jmp	cudaFree@PLT
	.cfi_endproc
.LFE2157:
	.size	_ZN15ConcStadiumHashImjE15freeTicketBoardEv, .-_ZN15ConcStadiumHashImjE15freeTicketBoardEv
	.section	.text._ZN15ConcStadiumHashImjE13prepareTicketEjj,"axG",@progbits,_ZN15ConcStadiumHashImjE13prepareTicketEjj,comdat
	.align 2
	.weak	_ZN15ConcStadiumHashImjE13prepareTicketEjj
	.type	_ZN15ConcStadiumHashImjE13prepareTicketEjj, @function
_ZN15ConcStadiumHashImjE13prepareTicketEjj:
.LFB2158:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2158:
	.size	_ZN15ConcStadiumHashImjE13prepareTicketEjj, .-_ZN15ConcStadiumHashImjE13prepareTicketEjj
	.section	.text._ZN15ConcStadiumHashImjE11extractInfoEjj,"axG",@progbits,_ZN15ConcStadiumHashImjE11extractInfoEjj,comdat
	.align 2
	.weak	_ZN15ConcStadiumHashImjE11extractInfoEjj
	.type	_ZN15ConcStadiumHashImjE11extractInfoEjj, @function
_ZN15ConcStadiumHashImjE11extractInfoEjj:
.LFB2159:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2159:
	.size	_ZN15ConcStadiumHashImjE11extractInfoEjj, .-_ZN15ConcStadiumHashImjE11extractInfoEjj
	.section	.text._ZN15ConcStadiumHashImjE10getTbIndexEj,"axG",@progbits,_ZN15ConcStadiumHashImjE10getTbIndexEj,comdat
	.align 2
	.weak	_ZN15ConcStadiumHashImjE10getTbIndexEj
	.type	_ZN15ConcStadiumHashImjE10getTbIndexEj, @function
_ZN15ConcStadiumHashImjE10getTbIndexEj:
.LFB2160:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2160:
	.size	_ZN15ConcStadiumHashImjE10getTbIndexEj, .-_ZN15ConcStadiumHashImjE10getTbIndexEj
	.section	.text._ZN15ConcStadiumHashImjE11getPosInIntEj,"axG",@progbits,_ZN15ConcStadiumHashImjE11getPosInIntEj,comdat
	.align 2
	.weak	_ZN15ConcStadiumHashImjE11getPosInIntEj
	.type	_ZN15ConcStadiumHashImjE11getPosInIntEj, @function
_ZN15ConcStadiumHashImjE11getPosInIntEj:
.LFB2161:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2161:
	.size	_ZN15ConcStadiumHashImjE11getPosInIntEj, .-_ZN15ConcStadiumHashImjE11getPosInIntEj
	.section	.text._ZN15ConcStadiumHashImjE12tryBookASeatEjj,"axG",@progbits,_ZN15ConcStadiumHashImjE12tryBookASeatEjj,comdat
	.align 2
	.weak	_ZN15ConcStadiumHashImjE12tryBookASeatEjj
	.type	_ZN15ConcStadiumHashImjE12tryBookASeatEjj, @function
_ZN15ConcStadiumHashImjE12tryBookASeatEjj:
.LFB2162:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2162:
	.size	_ZN15ConcStadiumHashImjE12tryBookASeatEjj, .-_ZN15ConcStadiumHashImjE12tryBookASeatEjj
	.section	.text._ZN15ConcStadiumHashImjE14tryFindTheSeatEPmjjjjS1_j,"axG",@progbits,_ZN15ConcStadiumHashImjE14tryFindTheSeatEPmjjjjS1_j,comdat
	.align 2
	.weak	_ZN15ConcStadiumHashImjE14tryFindTheSeatEPmjjjjS1_j
	.type	_ZN15ConcStadiumHashImjE14tryFindTheSeatEPmjjjjS1_j, @function
_ZN15ConcStadiumHashImjE14tryFindTheSeatEPmjjjjS1_j:
.LFB2163:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2163:
	.size	_ZN15ConcStadiumHashImjE14tryFindTheSeatEPmjjjjS1_j, .-_ZN15ConcStadiumHashImjE14tryFindTheSeatEPmjjjjS1_j
	.section	.text._ZN15ConcStadiumHashImjE14tryFindTheSeatEmjjjjPm,"axG",@progbits,_ZN15ConcStadiumHashImjE14tryFindTheSeatEmjjjjPm,comdat
	.align 2
	.weak	_ZN15ConcStadiumHashImjE14tryFindTheSeatEmjjjjPm
	.type	_ZN15ConcStadiumHashImjE14tryFindTheSeatEmjjjjPm, @function
_ZN15ConcStadiumHashImjE14tryFindTheSeatEmjjjjPm:
.LFB2164:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2164:
	.size	_ZN15ConcStadiumHashImjE14tryFindTheSeatEmjjjjPm, .-_ZN15ConcStadiumHashImjE14tryFindTheSeatEmjjjjPm
	.section	.text._ZN15ConcStadiumHashImjE16insertTicketInfoEjjj,"axG",@progbits,_ZN15ConcStadiumHashImjE16insertTicketInfoEjjj,comdat
	.align 2
	.weak	_ZN15ConcStadiumHashImjE16insertTicketInfoEjjj
	.type	_ZN15ConcStadiumHashImjE16insertTicketInfoEjjj, @function
_ZN15ConcStadiumHashImjE16insertTicketInfoEjjj:
.LFB2165:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2165:
	.size	_ZN15ConcStadiumHashImjE16insertTicketInfoEjjj, .-_ZN15ConcStadiumHashImjE16insertTicketInfoEjjj
	.section	.text._ZN15ConcStadiumHashImjE10isAccessedEjj,"axG",@progbits,_ZN15ConcStadiumHashImjE10isAccessedEjj,comdat
	.align 2
	.weak	_ZN15ConcStadiumHashImjE10isAccessedEjj
	.type	_ZN15ConcStadiumHashImjE10isAccessedEjj, @function
_ZN15ConcStadiumHashImjE10isAccessedEjj:
.LFB2166:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2166:
	.size	_ZN15ConcStadiumHashImjE10isAccessedEjj, .-_ZN15ConcStadiumHashImjE10isAccessedEjj
	.section	.text._ZN15ConcStadiumHashImjE12unlockAccessEjj,"axG",@progbits,_ZN15ConcStadiumHashImjE12unlockAccessEjj,comdat
	.align 2
	.weak	_ZN15ConcStadiumHashImjE12unlockAccessEjj
	.type	_ZN15ConcStadiumHashImjE12unlockAccessEjj, @function
_ZN15ConcStadiumHashImjE12unlockAccessEjj:
.LFB2167:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2167:
	.size	_ZN15ConcStadiumHashImjE12unlockAccessEjj, .-_ZN15ConcStadiumHashImjE12unlockAccessEjj
	.section	.text._ZN15ConcStadiumHashImjEC2Ej,"axG",@progbits,_ZN15ConcStadiumHashImjEC5Ej,comdat
	.align 2
	.weak	_ZN15ConcStadiumHashImjEC2Ej
	.type	_ZN15ConcStadiumHashImjEC2Ej, @function
_ZN15ConcStadiumHashImjEC2Ej:
.LFB2169:
	.cfi_startproc
	endbr64
	movl	%esi, 8(%rdi)
	ret
	.cfi_endproc
.LFE2169:
	.size	_ZN15ConcStadiumHashImjEC2Ej, .-_ZN15ConcStadiumHashImjEC2Ej
	.weak	_ZN15ConcStadiumHashImjEC1Ej
	.set	_ZN15ConcStadiumHashImjEC1Ej,_ZN15ConcStadiumHashImjEC2Ej
	.section	.text._ZN15ConcStadiumHashImjED2Ev,"axG",@progbits,_ZN15ConcStadiumHashImjED5Ev,comdat
	.align 2
	.weak	_ZN15ConcStadiumHashImjED2Ev
	.type	_ZN15ConcStadiumHashImjED2Ev, @function
_ZN15ConcStadiumHashImjED2Ev:
.LFB2172:
	.cfi_startproc
	.cfi_personality 0x9b,DW.ref.__gxx_personality_v0
	.cfi_lsda 0x1b,.LLSDA2172
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	_ZN15ConcStadiumHashImjE15freeTicketBoardEv
	popq	%rdx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2172:
	.section	.gcc_except_table._ZN15ConcStadiumHashImjED2Ev,"aG",@progbits,_ZN15ConcStadiumHashImjED5Ev,comdat
.LLSDA2172:
	.byte	0xff
	.byte	0xff
	.byte	0x1
	.uleb128 .LLSDACSE2172-.LLSDACSB2172
.LLSDACSB2172:
.LLSDACSE2172:
	.section	.text._ZN15ConcStadiumHashImjED2Ev,"axG",@progbits,_ZN15ConcStadiumHashImjED5Ev,comdat
	.size	_ZN15ConcStadiumHashImjED2Ev, .-_ZN15ConcStadiumHashImjED2Ev
	.weak	_ZN15ConcStadiumHashImjED1Ev
	.set	_ZN15ConcStadiumHashImjED1Ev,_ZN15ConcStadiumHashImjED2Ev
	.section	.text._ZN15ConcStadiumHashImjE6insertEPmjS1_Pjj,"axG",@progbits,_ZN15ConcStadiumHashImjE6insertEPmjS1_Pjj,comdat
	.align 2
	.weak	_ZN15ConcStadiumHashImjE6insertEPmjS1_Pjj
	.type	_ZN15ConcStadiumHashImjE6insertEPmjS1_Pjj, @function
_ZN15ConcStadiumHashImjE6insertEPmjS1_Pjj:
.LFB2174:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2174:
	.size	_ZN15ConcStadiumHashImjE6insertEPmjS1_Pjj, .-_ZN15ConcStadiumHashImjE6insertEPmjS1_Pjj
	.section	.text._ZN15ConcStadiumHashImjE4findEPmS1_j,"axG",@progbits,_ZN15ConcStadiumHashImjE4findEPmS1_j,comdat
	.align 2
	.weak	_ZN15ConcStadiumHashImjE4findEPmS1_j
	.type	_ZN15ConcStadiumHashImjE4findEPmS1_j, @function
_ZN15ConcStadiumHashImjE4findEPmS1_j:
.LFB2175:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2175:
	.size	_ZN15ConcStadiumHashImjE4findEPmS1_j, .-_ZN15ConcStadiumHashImjE4findEPmS1_j
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2076:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	movq	%rax, %rdi
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rdx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2076:
	.size	_ZL24__sti____cudaRegisterAllv, .-_ZL24__sti____cudaRegisterAllv
	.section	.init_array,"aw"
	.align 8
	.quad	_ZL24__sti____cudaRegisterAllv
	.section	.nvFatBinSegment,"aw"
	.align 8
	.type	_ZL15__fatDeviceText, @object
	.size	_ZL15__fatDeviceText, 24
_ZL15__fatDeviceText:
	.long	1180844977
	.long	1
	.quad	fatbinData
	.quad	0
	.local	_ZL20__cudaFatCubinHandle
	.comm	_ZL20__cudaFatCubinHandle,8,8
	.hidden	DW.ref.__gxx_personality_v0
	.weak	DW.ref.__gxx_personality_v0
	.section	.data.rel.local.DW.ref.__gxx_personality_v0,"awG",@progbits,DW.ref.__gxx_personality_v0,comdat
	.align 8
	.type	DW.ref.__gxx_personality_v0, @object
	.size	DW.ref.__gxx_personality_v0, 8
DW.ref.__gxx_personality_v0:
	.quad	__gxx_personality_v0
	.ident	"GCC: (Ubuntu 13.3.0-6ubuntu2~24.04) 13.3.0"
	.section	.note.GNU-stack,"",@progbits
	.section	.note.gnu.property,"a"
	.align 8
	.long	1f - 0f
	.long	4f - 1f
	.long	5
0:
	.string	"GNU"
1:
	.align 8
	.long	0xc0000002
	.long	3f - 2f
2:
	.long	0x3
3:
	.align 8
4:

#pragma once
#include "hip/hip_runtime.h"
#include <assert.h>
#include <stdio.h>

typedef unsigned int uint;

#define KEY_NOT_INSIDE_TABLE 2

// Implementation of Stadium Hashing with the access bit to allow concurrent inserts and searches
template<typename keyT, typename valueT>
class ConcStadiumHash {
protected:
	volatile uint* ticketBoard;
	uint tableSize;

	// Hash Functions
	template <uint nRounds, uint rShift, uint mulOp>
	__device__ uint hash(const keyT key);

	__device__ uint initHash(const keyT key);
	__device__ uint rehash(uint hashed, const keyT key);
	__device__ uint infoHash(keyT key);

	__device__ uint arrayHash(keyT* arr, uint length);
	__device__ uint arrayRehash(uint hashed, keyT* arr, uint length);
	__device__ uint arrayInfoHash(keyT* arr, uint length);

	
	// Ticket Board
	__host__ void allocTicketBoard();
	__host__ void clearTicketBoard();
	__host__ void freeTicketBoard();
	__device__ uint prepareTicket(uint info, uint infoStart);
	__device__ uint extractInfo(uint ticket, uint infoStart);
	__device__ uint getTbIndex(uint hashed);
	__device__ uint getPosInInt(uint hashed);
	__device__ uint tryBookASeat(uint tbIndex, uint posInInt);
	__device__ uint tryFindTheSeat(keyT* key, uint hashed, uint tbIndex, uint posInInt, uint info, keyT* keys, uint keyLength);
	__device__ uint tryFindTheSeat(keyT key, uint hashed, uint tbIndex, uint posInInt, uint info, keyT* keys);
	__device__ void insertTicketInfo(uint info, uint tbIndex, uint posInInt);

	__device__ uint isAccessed(uint tbIndex, uint posInInt);
	__device__ void unlockAccess(uint tbIndex, uint posInInt);


public:
	ConcStadiumHash(uint tableSize);
	~ConcStadiumHash();
	__device__ void insert(keyT* key, valueT value, keyT* keys, valueT* values, uint keyLength);
	__device__ uint find(keyT* key, keyT* keys, uint keyLength);
};




	.text
	.file	"ConcStadiumHash.hip"
	.type	__hip_cuid_,@object             # @__hip_cuid_
	.bss
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               # 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_

	.text
	.p2alignl 7, 3214868480
	.fill 96, 4, 3214868480
	.type	__hip_cuid_,@object             ; @__hip_cuid_
	.section	.bss,"aw",@nobits
	.globl	__hip_cuid_
__hip_cuid_:
	.byte	0                               ; 0x0
	.size	__hip_cuid_, 1

	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.3.2 25012 e5bf7e55c91490b07c49d8960fa7983d864936c4)"
	.section	".note.GNU-stack","",@progbits
	.addrsig
	.addrsig_sym __hip_cuid_
	.amdgcn_target "amdgcn-amd-amdhsa--gfx1100"
	.amdgpu_metadata
---
amdhsa.kernels:  []
amdhsa.target:   amdgcn-amd-amdhsa--gfx1100
amdhsa.version:
  - 1
  - 2
...

	.end_amdgpu_metadata
