
/**
 * Copyright 1993-2012 NVIDIA Corporation.  All rights reserved.
 *
 * Please refer to the NVIDIA end user license agreement (EULA) associated
 * with this source code for terms and conditions that govern your use of
 * this software. Any use, reproduction, disclosure, or distribution of
 * this software and related documentation outside the terms of the EULA
 * is strictly prohibited.
 */
#include <stdio.h>
#include <stdlib.h>

#include <cuda.h>
#include <assert.h>
#include "tree.cuh"

__device__ node* tempData[INPUTSIZE];
__device__ node* tempData1[INPUTSIZE];
__device__ node* root1;
__device__ node* globalCurr;
__device__ node* globalCurrs[ORDER];
__device__ node* newNode;
__device__ int globalIdx;
__device__ int tempKeys[ORDER];
__device__ node* tempPointers[ORDER];
__device__ int globalPointerIdx;
__device__ node* globalCurr1, *globalCurr2;
__device__ node* foundChild;


__device__ void make_node(node*& new_node)
{
	new_node = (node*)malloc(sizeof(node));
	new_node->keys = (int*)malloc( (ORDER - 1) * sizeof(int) );
	new_node->pointers = (node**)malloc( ORDER * sizeof(node *) );
	new_node->is_leaf = false;
	new_node->num_keys = 0;
	new_node->parent = NULL;
	new_node->next = NULL;
}

__device__ void make_leaf(node*& new_node)
{
	make_node(new_node);
	new_node->is_leaf = true;
}

__global__ void buildLeaves(node*& root, int* input, int* result)
{
	unsigned inWholeIdx = blockIdx.x*blockDim.x+threadIdx.x;
	unsigned noOfNodes = INPUTSIZE / ((ORDER / 2) - 1);
	node* newNode;
	if(inWholeIdx < noOfNodes)
	{
		make_leaf(newNode);
		tempData[inWholeIdx] = newNode;
		assert(tempData[inWholeIdx]);
	}
}

__global__ void buildRoot(node*& root, int* input, int* result)
{
	unsigned inWholeIdx = blockIdx.x*blockDim.x+threadIdx.x;
	if(inWholeIdx == 0)
	{
		root1 = (node*)malloc(sizeof(node));
		root1->keys = (int*)malloc( (ORDER - 1) * sizeof(int) );
		root1->pointers = (node**)malloc( ORDER * sizeof(node *) );
		root1->is_leaf = false;
		root1->num_keys = 0;
		root1->parent = NULL;
		root1->next = NULL;
		root1->keys[0] = 5;
	}
}

__global__ void buildLevel(node*& root, int* input, int* result, int size, int x)
{
	node** arr;
	unsigned inWholeIdx = blockIdx.x*blockDim.x+threadIdx.x;
	unsigned noOfNodes = size / (ORDER / 2);
	if(x)
		arr = tempData1;
	else
		arr = tempData;
	if(inWholeIdx < noOfNodes)
	{
		node* newNode;
		make_node(newNode);
		newNode->keys[0] = inWholeIdx;
		arr[inWholeIdx] = newNode;
	}
}

__global__ void fillLevel(node*& root, int* input, int* result, int size, int x)
{
	node** parent;
	node** children;
	if(x)
	{
		parent = tempData1;
		children = tempData;
	}
	else
	{
		parent = tempData;
		children = tempData1;
	}
	unsigned inWholeIdx = blockIdx.x*blockDim.x+threadIdx.x;
	unsigned noOfNodes = size / (ORDER / 2);
	unsigned inNodeIdx = inWholeIdx % (ORDER / 2);
	unsigned nodeNo = inWholeIdx / (ORDER / 2 );
	if(nodeNo == noOfNodes)
	{
		nodeNo--;
		inNodeIdx = ((ORDER/2)) + inNodeIdx;
	}
	if(inWholeIdx < size)
	{
		assert(parent[nodeNo]);
		assert(parent[nodeNo]->keys);
		assert(children[inWholeIdx]);
		parent[nodeNo]->pointers[inNodeIdx] = children[inWholeIdx];
		children[inWholeIdx]->parent = parent[nodeNo];
		if(inNodeIdx < (ORDER/2) - 1 || (nodeNo == noOfNodes -1 && inWholeIdx != size - 1))
		{
			assert(children[inWholeIdx]);
			assert(children[inWholeIdx]->num_keys);
			assert(children[inWholeIdx]->keys[children[inWholeIdx]->num_keys-1]);
			parent[nodeNo]->keys[inNodeIdx] = children[inWholeIdx]->keys[children[inWholeIdx]->num_keys-1];
			assert(parent[nodeNo]->keys[inNodeIdx]);
		}
	}
	if(inNodeIdx == 0)
	{
		if(nodeNo < noOfNodes -1)
		{
			parent[nodeNo]->num_keys = (ORDER / 2) - 1;
		}
		else if(nodeNo == noOfNodes - 1)
			parent[nodeNo]->num_keys = (size % (ORDER / 2)) + (ORDER / 2) - 1;
	}
}

__global__ void fillRoot(node*& root, int* input, int* result, int size, int x)
{
	node** children;
	if(x)
	{
		children = tempData;
	}
	else
	{
		children = tempData1;
	}
	unsigned inWholeIdx = blockIdx.x*blockDim.x+threadIdx.x;
	unsigned inNodeIdx = inWholeIdx % size;
	if(inWholeIdx < size)
	{
		assert(children[inWholeIdx]);
		root1->pointers[inNodeIdx] = children[inWholeIdx];
		children[inWholeIdx]->parent = root1;
		if(inNodeIdx < size -1 )
		{
			assert(children[inWholeIdx]);
			assert(children[inWholeIdx]->num_keys);
			assert(children[inWholeIdx]->keys[children[inWholeIdx]->num_keys-1]);
			root1->keys[inWholeIdx] = children[inWholeIdx]->keys[children[inWholeIdx]->num_keys-1];
			assert(root1->keys[inNodeIdx]);
		}
	}
	if(inWholeIdx == 0)
	{
		root1->num_keys = size - 1;
	}
}

__global__ void fillLeaves(node*& root, int* input, int* result)
{
	unsigned inWholeIdx = blockIdx.x*blockDim.x+threadIdx.x;
	unsigned noOfNodes = INPUTSIZE / ((ORDER / 2) - 1);
	unsigned inNodeIdx = inWholeIdx % ((ORDER / 2) - 1);
	unsigned nodeNo = inWholeIdx / ((ORDER / 2 ) -1);
	if(nodeNo == noOfNodes)
	{
		nodeNo--;
		inNodeIdx = ((ORDER/2) - 1) + inNodeIdx;
	}
	if(inWholeIdx < INPUTSIZE)
	{
		assert(tempData[nodeNo]);
		assert(tempData[nodeNo]->keys);
		assert(input[inWholeIdx]);
		tempData[nodeNo]->keys[inNodeIdx] = input[inWholeIdx];
	}
	if(inNodeIdx == 0)
	{
		if(nodeNo < noOfNodes -1)
		{
			tempData[nodeNo]->next = tempData[nodeNo + 1];
			tempData[nodeNo]->num_keys = ((ORDER / 2) - 1);
		}
		else if(nodeNo == noOfNodes - 1)
			tempData[nodeNo]->num_keys = (INPUTSIZE % ((ORDER / 2) - 1)) + ((ORDER / 2) - 1);
	}
}

__global__ void bulkLoad(node*& root, int* input, int* result)
{
	unsigned inWholeIdx = blockIdx.x*blockDim.x+threadIdx.x;
	unsigned noOfNodes = INPUTSIZE / ((ORDER / 2) - 1);
	unsigned inNodeIdx = inWholeIdx % ((ORDER / 2) - 1);
	unsigned nodeNo = inWholeIdx / ((ORDER / 2 ) -1);
	if(inWholeIdx == 0)
	{
		root->keys = (int*)malloc( (ORDER - 1) * sizeof(int) );
		root->pointers = (node**)malloc( ORDER * sizeof(node *) );
		root->is_leaf = false;
		root->num_keys = 0;
		root->parent = NULL;
		root->next = NULL;
	}
	node* newNode;
	if(inNodeIdx == 0 && nodeNo < noOfNodes)
	{
		make_leaf(newNode);
		tempData[nodeNo] = newNode;
		assert(tempData[nodeNo]);
	}
	if(nodeNo == noOfNodes)
	{
		nodeNo--;
		inNodeIdx = ((ORDER/2) - 1) + inNodeIdx;
	}
	__syncthreads();
	if(inWholeIdx < INPUTSIZE && nodeNo < noOfNodes)
	{
		tempData[nodeNo]->keys[inNodeIdx] = input[inWholeIdx];
	}
}

__device__ void addKey(node* curr, node* child)
{
	unsigned inWholeIdx = blockIdx.x*blockDim.x+threadIdx.x;
	int val = child->keys[0];
	if(contains(curr, val))
		return;
	if(inWholeIdx <= curr->num_keys)
	{
		if(inWholeIdx < curr->num_keys)
			tempKeys[inWholeIdx] = curr->keys[inWholeIdx];
		if(!curr->is_leaf)
			tempPointers[inWholeIdx] = curr->pointers[inWholeIdx];
	}
	if(inWholeIdx <= curr->num_keys)
	{
		if(inWholeIdx == 0)
		{
			if(val <= curr->keys[0])
			{
				globalIdx = 0;
			}
		}
		else if(inWholeIdx < curr->num_keys && inWholeIdx > 0)
		{
			if(curr->keys[inWholeIdx-1] < val && val <= curr->keys[inWholeIdx])
			{
				globalIdx = inWholeIdx;
			}
		}
		else if(inWholeIdx == curr->num_keys)
		{
			if(val > curr->keys[curr->num_keys - 1])
			{
				globalIdx = curr->num_keys;
			}
		}
	}
	__syncthreads();
	if(inWholeIdx >= globalIdx && inWholeIdx <= curr->num_keys)
	{
		if(inWholeIdx < curr->num_keys)
			curr->keys[inWholeIdx+1] = tempKeys[inWholeIdx];
		if(!curr->is_leaf)
			curr->pointers[inWholeIdx+1] = tempPointers[inWholeIdx];
	}
	__syncthreads();
	if(inWholeIdx == globalIdx)
	{
		if(inWholeIdx > 0)
			curr->keys[globalIdx] = val;
		else
			curr->keys[globalIdx] = child->keys[child->num_keys]+1;
		if(!curr->is_leaf)
			curr->pointers[globalIdx] = child;
	}
	__syncthreads();
	if(inWholeIdx == 0)
		curr->num_keys++;
}

__device__ void split(node* curr, node* child)
{
	unsigned inWholeIdx = blockIdx.x*blockDim.x+threadIdx.x;
	node* newNodeLocal;
	if(inWholeIdx == 0)
	{
		newNode = (node*)malloc(sizeof(node));
		newNode->keys = (int*)malloc( (ORDER - 1) * sizeof(int) );
		newNode->pointers = (node**)malloc( ORDER * sizeof(node *) );
		newNode->is_leaf = curr->is_leaf;
		newNode->num_keys = ORDER/2;
		newNode->parent = curr->parent;
		newNode->next = curr->next;
		curr->num_keys = ORDER/2;
		curr->next = newNode;
		globalPointerIdx = 0;
	}
	__syncthreads();
	newNodeLocal = newNode;
	__syncthreads();
	if(inWholeIdx < (ORDER /2))
	{
		newNode->keys[inWholeIdx] = curr->keys[ORDER/2 + inWholeIdx];
	}
	if(!curr->is_leaf && inWholeIdx <= (ORDER /2))
	{
		newNode->pointers[inWholeIdx] = curr->pointers[ORDER/2 + inWholeIdx];
	}
	if(curr->parent->num_keys >= ORDER)
		split(curr, newNode);
	else
		addKey(curr->parent, newNodeLocal);
}

__global__ void createNewNode(node*& root)
{
	unsigned inWholeIdx = blockIdx.x*blockDim.x+threadIdx.x;
	node* curr = foundChild;
	if(inWholeIdx == 0)
	{
		root = (node*)malloc(sizeof(node));
		root->keys = (int*)malloc( (ORDER - 1) * sizeof(int) );
		root->pointers = (node**)malloc( ORDER * sizeof(node *) );
		root->is_leaf = curr->is_leaf;
		root->num_keys = ORDER/2;
		root->parent = curr->parent;
		root->next = curr->next;
		curr->num_keys = ORDER/2;
		curr->next = newNode;
		globalPointerIdx = 0;
	}
}

__device__ void split(node* curr, int val)
{
	unsigned inWholeIdx = blockIdx.x*blockDim.x+threadIdx.x;
	node* newNodeLocal;
	if(inWholeIdx == 0)
	{
		newNode = (node*)malloc(sizeof(node));
		newNode->keys = (int*)malloc( (ORDER - 1) * sizeof(int) );
		newNode->pointers = (node**)malloc( ORDER * sizeof(node *) );
		newNode->is_leaf = curr->is_leaf;
		newNode->num_keys = ORDER/2;
		newNode->parent = curr->parent;
		newNode->next = curr->next;
		curr->num_keys = ORDER/2;
		curr->next = newNode;
		globalPointerIdx = 0;
	}
	__syncthreads();
	newNodeLocal = newNode;
	__syncthreads();
	if(inWholeIdx < (ORDER /2))
	{
		newNode->keys[inWholeIdx] = curr->keys[ORDER/2 + inWholeIdx];
	}
	if(!curr->is_leaf && inWholeIdx <= (ORDER /2))
	{
		newNode->pointers[inWholeIdx] = curr->pointers[ORDER/2 + inWholeIdx];
	}
	if(curr->parent->num_keys >= ORDER)
		split(curr, newNode);
	else
		addKey(curr->parent, newNodeLocal);
}

__device__ void addKey(node* curr, int val)
{
	unsigned inWholeIdx = blockIdx.x*blockDim.x+threadIdx.x;
	if(contains(curr, val))
		return;
	if(inWholeIdx < curr->num_keys)
		tempKeys[inWholeIdx] = curr->keys[inWholeIdx];
	if(inWholeIdx <= curr->num_keys)
	{
		if(inWholeIdx == 0)
		{
			if(val <= curr->keys[0])
			{
				globalIdx = 0;
			}
		}
		else if(inWholeIdx < curr->num_keys && inWholeIdx > 0)
		{
			if(curr->keys[inWholeIdx-1] < val && val <= curr->keys[inWholeIdx])
			{
				globalIdx = inWholeIdx;
			}
		}
		else if(inWholeIdx == curr->num_keys)
		{
			if(val > curr->keys[curr->num_keys - 1])
			{
				globalIdx = curr->num_keys;
			}
		}
	}
	__syncthreads();
	if(inWholeIdx >= globalIdx && inWholeIdx < curr->num_keys)
		curr->keys[inWholeIdx+1] = tempKeys[inWholeIdx];
	__syncthreads();
	if(inWholeIdx == globalIdx)
		curr->keys[globalIdx] = val;
	__syncthreads();
	if(inWholeIdx == 0)
		curr->num_keys++;
}

__global__ void insertVal(int val)
{
	node* curr = find(val);
	__syncthreads();
	assert(curr->num_keys < ORDER -1);
	if(curr->num_keys < ORDER -1)
		addKey(curr, val);
	else
		split(curr, val);
}

__global__ void copyNode(node* node1, int* full)
{
	unsigned inWholeIdx = blockIdx.x*blockDim.x+threadIdx.x;
	foundChild = foundChild->parent;
	node* curr = foundChild;
	int val = node1->keys[0];
	if(inWholeIdx == 0)
	{
		if(foundChild->num_keys == ORDER - 1)
			full[0] = 1;
		else
			full[0] = 0;

	}
	if(inWholeIdx <= curr->num_keys)
	{
		if(inWholeIdx < curr->num_keys)
			tempKeys[inWholeIdx] = curr->keys[inWholeIdx];
		if(!curr->is_leaf)
			tempPointers[inWholeIdx] = curr->pointers[inWholeIdx];
	}
	if(inWholeIdx <= curr->num_keys)
	{
		if(inWholeIdx == 0)
		{
			if(val <= curr->keys[0])
			{
				globalIdx = 0;
			}
		}
		else if(inWholeIdx < curr->num_keys && inWholeIdx > 0)
		{
			if(curr->keys[inWholeIdx-1] < val && val <= curr->keys[inWholeIdx])
			{
				globalIdx = inWholeIdx;
			}
		}
		else if(inWholeIdx == curr->num_keys)
		{
			if(val > curr->keys[curr->num_keys - 1])
			{
				globalIdx = curr->num_keys;
			}
		}
	}
	if(inWholeIdx == 0 && foundChild->num_keys == ORDER - 1)
	{
		foundChild->num_keys = ORDER/2;
		foundChild = foundChild->parent;
	}
}

__global__ void copyNode(int val, int* full)
{
	unsigned inWholeIdx = blockIdx.x*blockDim.x+threadIdx.x;
	node* curr = foundChild;
	if(inWholeIdx == 0)
	{
		if(foundChild->num_keys == ORDER - 1)
			full[0] = 1;
		else
			full[0] = 0;

	}
	if(inWholeIdx <= curr->num_keys)
	{
		if(inWholeIdx < curr->num_keys)
			tempKeys[inWholeIdx] = curr->keys[inWholeIdx];
		if(!curr->is_leaf)
			tempPointers[inWholeIdx] = curr->pointers[inWholeIdx];
	}
	if(inWholeIdx <= curr->num_keys)
	{
		if(inWholeIdx == 0)
		{
			if(val <= curr->keys[0])
			{
				globalIdx = 0;
			}
		}
		else if(inWholeIdx < curr->num_keys && inWholeIdx > 0)
		{
			if(curr->keys[inWholeIdx-1] < val && val <= curr->keys[inWholeIdx])
			{
				globalIdx = inWholeIdx;
			}
		}
		else if(inWholeIdx == curr->num_keys)
		{
			if(val > curr->keys[curr->num_keys - 1])
			{
				globalIdx = curr->num_keys;
			}
		}
	}
	if(inWholeIdx == 0 && foundChild->num_keys == ORDER - 1)
	{
		foundChild->num_keys = ORDER/2;
	}
}

__global__ void addValue(int val)
{
	node* curr = foundChild;
	unsigned inWholeIdx = blockIdx.x*blockDim.x+threadIdx.x;
	if(inWholeIdx > globalIdx && inWholeIdx < curr->num_keys)
			curr->keys[inWholeIdx] = tempKeys[inWholeIdx - 1];
	if(inWholeIdx == globalIdx)
	{
		curr->keys[globalIdx] = val;
		curr->num_keys++;
	}
}

__global__ void addValue(node* nn)
{
	node* curr = foundChild;
	unsigned inWholeIdx = blockIdx.x*blockDim.x+threadIdx.x;
	if(inWholeIdx > globalIdx && inWholeIdx < curr->num_keys)
			curr->keys[inWholeIdx] = tempKeys[inWholeIdx - 1];
	if(inWholeIdx == globalIdx)
	{
		curr->pointers[globalIdx] = nn;
		curr->num_keys++;
		if(globalIdx == 0)
			curr->keys[0] = nn->keys[nn->num_keys-1];
		else
			curr->keys[globalIdx] = nn->keys[0];
	}
}

__global__ void copyToNewNode(node*& nnode)
{
	unsigned inWholeIdx = blockIdx.x*blockDim.x+threadIdx.x;
	if(inWholeIdx < (ORDER /2))
	{
		nnode->keys[inWholeIdx] = tempKeys[ORDER/2 + inWholeIdx];
	}
	if(!nnode->is_leaf && inWholeIdx <= (ORDER /2))
	{
		nnode->pointers[inWholeIdx] = tempPointers[ORDER/2 + inWholeIdx];
	}
}


__device__ int contains(node* curr, int val)
{
	unsigned inWholeIdx = blockIdx.x*blockDim.x+threadIdx.x;
	if(inWholeIdx < curr->num_keys)
	{
		if(curr->keys[inWholeIdx] == val)
			globalIdx = 1;
	}
	__syncthreads();
	return globalIdx;
}

__device__ node* find(int val)
{
	unsigned inWholeIdx = blockIdx.x*blockDim.x+threadIdx.x;
	assert(root1);
	node* curr = root1;
	assert(curr);
	assert(!curr->is_leaf);
	__syncthreads();
	while(!curr->is_leaf)
	{
		if(inWholeIdx <= curr->num_keys)
		{
			if(inWholeIdx == 0)
			{
				assert(curr->keys[0]);
				if(val <= curr->keys[0])
				{
					assert(curr->pointers[0]);
					globalCurr = curr->pointers[0];
				}
			}
			else if(inWholeIdx < curr->num_keys && inWholeIdx > 0)
			{
				assert(curr->keys[inWholeIdx-1]);
				assert(curr->keys[inWholeIdx]);
				if(curr->keys[inWholeIdx-1] < val && val <= curr->keys[inWholeIdx])
				{
					assert(curr->pointers[inWholeIdx]);
					globalCurr = curr->pointers[inWholeIdx];
				}
			}
			else if(inWholeIdx == curr->num_keys)
			{
				assert(curr->keys[curr->num_keys - 1]);
				if(val > curr->keys[curr->num_keys - 1])
				{
					assert(curr->pointers[inWholeIdx]);
					globalCurr = curr->pointers[inWholeIdx];
				}
			}
		}
		__syncthreads();
		curr = globalCurr;
		__syncthreads();
	}
	return curr;
}

__device__ node* find(int* values, int len)
{
	//unsigned inWholeIdx = blockIdx.x*blockDim.x+threadIdx.x;
	unsigned inNodeIdx = threadIdx.x;
	unsigned nodeNo = blockIdx.x;
	int val;
	if(nodeNo < len)
		val = values[nodeNo];
	assert(root1);
	node* curr = root1;
	assert(curr);
	assert(!curr->is_leaf);
	__syncthreads();
	while(!curr->is_leaf)
	{
		if(inNodeIdx <= curr->num_keys && nodeNo < len)
		{
			if(inNodeIdx == 0)
			{
				assert(curr->keys[0]);
				if(val <= curr->keys[0])
				{
					assert(curr->pointers[0]);
					globalCurrs[nodeNo] = curr->pointers[0];
				}
			}
			else if(inNodeIdx < curr->num_keys && inNodeIdx > 0)
			{
				assert(curr->keys[inNodeIdx-1]);
				assert(curr->keys[inNodeIdx]);
				if(curr->keys[inNodeIdx-1] < val && val <= curr->keys[inNodeIdx])
				{
					assert(curr->pointers[inNodeIdx]);
					globalCurrs[nodeNo] = curr->pointers[inNodeIdx];
				}
			}
			else if(inNodeIdx == curr->num_keys)
			{
				assert(curr->keys[curr->num_keys - 1]);
				if(val > curr->keys[curr->num_keys - 1])
				{
					assert(curr->pointers[inNodeIdx]);
					globalCurrs[nodeNo] = curr->pointers[inNodeIdx];
				}
			}
		}
		__syncthreads();
		assert(globalCurrs[nodeNo]);
		curr = globalCurrs[nodeNo];
		__syncthreads();
	}
	return curr;
}

__global__ void searchBetter(int val, int* result)
{
	unsigned inWholeIdx = blockIdx.x*blockDim.x+threadIdx.x;
	node* curr;
	switch(result[0])
	{
		case 2:
			curr = root1;
			break;
		case 3:
			curr = globalCurr1;
			break;
		case 4:
			curr = globalCurr2;
			break;
		default:
			return;
	}
	assert(curr);
	assert(!curr->is_leaf);
	node* found = NULL;
	if(inWholeIdx <= curr->num_keys)
	{
		if(inWholeIdx == 0)
		{
			assert(curr->keys[0]);
			if(val <= curr->keys[0])
			{
				assert(curr->pointers[0]);
				found = curr->pointers[0];
			}
		}
		else if(inWholeIdx < curr->num_keys && inWholeIdx > 0)
		{
			assert(curr->keys[inWholeIdx-1]);
			assert(curr->keys[inWholeIdx]);
			if(curr->keys[inWholeIdx-1] < val && val <= curr->keys[inWholeIdx])
			{
				assert(curr->pointers[inWholeIdx]);
				found = curr->pointers[inWholeIdx];
			}
		}
		else if(inWholeIdx == curr->num_keys)
		{
			assert(curr->keys[curr->num_keys - 1]);
			if(val > curr->keys[curr->num_keys - 1])
			{
				assert(curr->pointers[inWholeIdx]);
				found = curr->pointers[inWholeIdx];
			}
		}
	}
	if(found != NULL)
	{
		assert(found);
		if(result[0] == 2 || result[0] == 3)
		{
			globalCurr2 = found;
			result[0] = 4;
		}
		else if(result[0] == 4)
		{
			globalCurr1 = found;
			result[0] = 3;
		}
		if(found->is_leaf)
			result[0] = result[0] * 2;
	}
}

__global__ void containsBetter(int val, int* result)
{
	unsigned inWholeIdx = blockIdx.x*blockDim.x+threadIdx.x;
	node* curr;
	if(result[0] == 6)
		curr = globalCurr1;
	else if(result[0] == 8)
		curr = globalCurr2;
	assert(curr->is_leaf);
	if(inWholeIdx == 0)
		foundChild = curr;
	if(inWholeIdx < curr->num_keys)
	{
		if(curr->keys[inWholeIdx] == val)
			result[0] = 1;
	}
}

__global__ void search(int val, int* result)
{
	unsigned inWholeIdx = blockIdx.x*blockDim.x+threadIdx.x;
	node* curr = find(val);
	result[0] = 0;
	if(inWholeIdx < curr->num_keys)
	{
		if(curr->keys[inWholeIdx] == val)
			result[0] = 1;
	}
}

__global__ void search(int* vals, int* results, int len)
{
	//unsigned inWholeIdx = blockIdx.x*blockDim.x+threadIdx.x;
	unsigned inNodeIdx = threadIdx.x;
	unsigned nodeNo = blockIdx.x;
	node* curr = find(vals, len);
	if(nodeNo < len)
		results[nodeNo] = 0;
	if(nodeNo < len && inNodeIdx < curr->num_keys)
	{
		if(curr->keys[inNodeIdx] == vals[nodeNo])
			results[nodeNo] = 1;
	}
}

__global__ void test(node*& root, int* input, int* result)
{
	unsigned inWholeIdx = blockIdx.x*blockDim.x+threadIdx.x;
	if(inWholeIdx== 0)
	{
		//node* curr = root1;
		/*while(!curr->is_leaf)
		{
			curr = curr->pointers[2];
		}*/
		result[0] = root1->pointers[root1->num_keys]->keys[0];
	}
}

	.file	"tmpxft_00224fbc_00000000-6_tree.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2038:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2038:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.type	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0, @function
_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0:
.LFB2222:
	.cfi_startproc
	jmp	cudaLaunchKernel@PLT
	.cfi_endproc
.LFE2222:
	.size	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0, .-_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	.globl	_Z9make_nodeRP4node
	.type	_Z9make_nodeRP4node, @function
_Z9make_nodeRP4node:
.LFB2027:
	.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
.LFE2027:
	.size	_Z9make_nodeRP4node, .-_Z9make_nodeRP4node
	.globl	_Z9make_leafRP4node
	.type	_Z9make_leafRP4node, @function
_Z9make_leafRP4node:
.LFB2217:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	_Z9make_nodeRP4node
	.cfi_endproc
.LFE2217:
	.size	_Z9make_leafRP4node, .-_Z9make_leafRP4node
	.globl	_Z6addKeyP4nodeS0_
	.type	_Z6addKeyP4nodeS0_, @function
_Z6addKeyP4nodeS0_:
.LFB2029:
	.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
.LFE2029:
	.size	_Z6addKeyP4nodeS0_, .-_Z6addKeyP4nodeS0_
	.globl	_Z5splitP4nodeS0_
	.type	_Z5splitP4nodeS0_, @function
_Z5splitP4nodeS0_:
.LFB2221:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	_Z6addKeyP4nodeS0_
	.cfi_endproc
.LFE2221:
	.size	_Z5splitP4nodeS0_, .-_Z5splitP4nodeS0_
	.globl	_Z6addKeyP4nodei
	.type	_Z6addKeyP4nodei, @function
_Z6addKeyP4nodei:
.LFB2032:
	.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
.LFE2032:
	.size	_Z6addKeyP4nodei, .-_Z6addKeyP4nodei
	.globl	_Z5splitP4nodei
	.type	_Z5splitP4nodei, @function
_Z5splitP4nodei:
.LFB2219:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	pushq	%rax
	.cfi_def_cfa_offset 16
	call	_Z6addKeyP4nodei
	.cfi_endproc
.LFE2219:
	.size	_Z5splitP4nodei, .-_Z5splitP4nodei
	.globl	_Z8containsP4nodei
	.type	_Z8containsP4nodei, @function
_Z8containsP4nodei:
.LFB2033:
	.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
.LFE2033:
	.size	_Z8containsP4nodei, .-_Z8containsP4nodei
	.globl	_Z4findi
	.type	_Z4findi, @function
_Z4findi:
.LFB2034:
	.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
.LFE2034:
	.size	_Z4findi, .-_Z4findi
	.globl	_Z4findPii
	.type	_Z4findPii, @function
_Z4findPii:
.LFB2035:
	.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
.LFE2035:
	.size	_Z4findPii, .-_Z4findPii
	.globl	_Z41__device_stub__Z11buildLeavesRP4nodePiS2_PP4nodePiS2_
	.type	_Z41__device_stub__Z11buildLeavesRP4nodePiS2_PP4nodePiS2_, @function
_Z41__device_stub__Z11buildLeavesRP4nodePiS2_PP4nodePiS2_:
.LFB2060:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 24(%rsp)
	leaq	48(%rsp), %rcx
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	40(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 80(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 88(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L21
	pushq	48(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z11buildLeavesRP4nodePiS2_(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 144
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	96(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L21:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L23
	call	__stack_chk_fail@PLT
.L23:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2060:
	.size	_Z41__device_stub__Z11buildLeavesRP4nodePiS2_PP4nodePiS2_, .-_Z41__device_stub__Z11buildLeavesRP4nodePiS2_PP4nodePiS2_
	.globl	_Z11buildLeavesRP4nodePiS2_
	.type	_Z11buildLeavesRP4nodePiS2_, @function
_Z11buildLeavesRP4nodePiS2_:
.LFB2061:
	.cfi_startproc
	endbr64
	jmp	_Z41__device_stub__Z11buildLeavesRP4nodePiS2_PP4nodePiS2_
	.cfi_endproc
.LFE2061:
	.size	_Z11buildLeavesRP4nodePiS2_, .-_Z11buildLeavesRP4nodePiS2_
	.globl	_Z38__device_stub__Z9buildRootRP4nodePiS2_PP4nodePiS2_
	.type	_Z38__device_stub__Z9buildRootRP4nodePiS2_PP4nodePiS2_, @function
_Z38__device_stub__Z9buildRootRP4nodePiS2_PP4nodePiS2_:
.LFB2062:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 24(%rsp)
	leaq	48(%rsp), %rcx
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	40(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 80(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 88(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L27
	pushq	48(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z9buildRootRP4nodePiS2_(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 144
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	96(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L27:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L29
	call	__stack_chk_fail@PLT
.L29:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2062:
	.size	_Z38__device_stub__Z9buildRootRP4nodePiS2_PP4nodePiS2_, .-_Z38__device_stub__Z9buildRootRP4nodePiS2_PP4nodePiS2_
	.globl	_Z9buildRootRP4nodePiS2_
	.type	_Z9buildRootRP4nodePiS2_, @function
_Z9buildRootRP4nodePiS2_:
.LFB2063:
	.cfi_startproc
	endbr64
	jmp	_Z38__device_stub__Z9buildRootRP4nodePiS2_PP4nodePiS2_
	.cfi_endproc
.LFE2063:
	.size	_Z9buildRootRP4nodePiS2_, .-_Z9buildRootRP4nodePiS2_
	.globl	_Z42__device_stub__Z10buildLevelRP4nodePiS2_iiPP4nodePiS2_ii
	.type	_Z42__device_stub__Z10buildLevelRP4nodePiS2_iiPP4nodePiS2_ii, @function
_Z42__device_stub__Z10buildLevelRP4nodePiS2_iiPP4nodePiS2_ii:
.LFB2064:
	.cfi_startproc
	endbr64
	subq	$136, %rsp
	.cfi_def_cfa_offset 144
	movq	%rdi, 24(%rsp)
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	40(%rsp), %rdx
	movl	%ecx, 4(%rsp)
	leaq	48(%rsp), %rcx
	movl	%r8d, (%rsp)
	movq	%fs:40, %rax
	movq	%rax, 120(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 80(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 88(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 96(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 104(%rsp)
	movq	%rsp, %rax
	movq	%rax, 112(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L32
	pushq	48(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z10buildLevelRP4nodePiS2_ii(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 160
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	96(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 152
	popq	%rdx
	.cfi_def_cfa_offset 144
.L32:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L34
	call	__stack_chk_fail@PLT
.L34:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2064:
	.size	_Z42__device_stub__Z10buildLevelRP4nodePiS2_iiPP4nodePiS2_ii, .-_Z42__device_stub__Z10buildLevelRP4nodePiS2_iiPP4nodePiS2_ii
	.globl	_Z10buildLevelRP4nodePiS2_ii
	.type	_Z10buildLevelRP4nodePiS2_ii, @function
_Z10buildLevelRP4nodePiS2_ii:
.LFB2065:
	.cfi_startproc
	endbr64
	jmp	_Z42__device_stub__Z10buildLevelRP4nodePiS2_iiPP4nodePiS2_ii
	.cfi_endproc
.LFE2065:
	.size	_Z10buildLevelRP4nodePiS2_ii, .-_Z10buildLevelRP4nodePiS2_ii
	.globl	_Z40__device_stub__Z9fillLevelRP4nodePiS2_iiPP4nodePiS2_ii
	.type	_Z40__device_stub__Z9fillLevelRP4nodePiS2_iiPP4nodePiS2_ii, @function
_Z40__device_stub__Z9fillLevelRP4nodePiS2_iiPP4nodePiS2_ii:
.LFB2066:
	.cfi_startproc
	endbr64
	subq	$136, %rsp
	.cfi_def_cfa_offset 144
	movq	%rdi, 24(%rsp)
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	40(%rsp), %rdx
	movl	%ecx, 4(%rsp)
	leaq	48(%rsp), %rcx
	movl	%r8d, (%rsp)
	movq	%fs:40, %rax
	movq	%rax, 120(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 80(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 88(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 96(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 104(%rsp)
	movq	%rsp, %rax
	movq	%rax, 112(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L37
	pushq	48(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z9fillLevelRP4nodePiS2_ii(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 160
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	96(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 152
	popq	%rdx
	.cfi_def_cfa_offset 144
.L37:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L39
	call	__stack_chk_fail@PLT
.L39:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2066:
	.size	_Z40__device_stub__Z9fillLevelRP4nodePiS2_iiPP4nodePiS2_ii, .-_Z40__device_stub__Z9fillLevelRP4nodePiS2_iiPP4nodePiS2_ii
	.globl	_Z9fillLevelRP4nodePiS2_ii
	.type	_Z9fillLevelRP4nodePiS2_ii, @function
_Z9fillLevelRP4nodePiS2_ii:
.LFB2067:
	.cfi_startproc
	endbr64
	jmp	_Z40__device_stub__Z9fillLevelRP4nodePiS2_iiPP4nodePiS2_ii
	.cfi_endproc
.LFE2067:
	.size	_Z9fillLevelRP4nodePiS2_ii, .-_Z9fillLevelRP4nodePiS2_ii
	.globl	_Z39__device_stub__Z8fillRootRP4nodePiS2_iiPP4nodePiS2_ii
	.type	_Z39__device_stub__Z8fillRootRP4nodePiS2_iiPP4nodePiS2_ii, @function
_Z39__device_stub__Z8fillRootRP4nodePiS2_iiPP4nodePiS2_ii:
.LFB2068:
	.cfi_startproc
	endbr64
	subq	$136, %rsp
	.cfi_def_cfa_offset 144
	movq	%rdi, 24(%rsp)
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	40(%rsp), %rdx
	movl	%ecx, 4(%rsp)
	leaq	48(%rsp), %rcx
	movl	%r8d, (%rsp)
	movq	%fs:40, %rax
	movq	%rax, 120(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 80(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 88(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 96(%rsp)
	leaq	4(%rsp), %rax
	movq	%rax, 104(%rsp)
	movq	%rsp, %rax
	movq	%rax, 112(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L42
	pushq	48(%rsp)
	.cfi_def_cfa_offset 152
	leaq	_Z8fillRootRP4nodePiS2_ii(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 160
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	96(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 152
	popq	%rdx
	.cfi_def_cfa_offset 144
.L42:
	movq	120(%rsp), %rax
	subq	%fs:40, %rax
	je	.L44
	call	__stack_chk_fail@PLT
.L44:
	addq	$136, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2068:
	.size	_Z39__device_stub__Z8fillRootRP4nodePiS2_iiPP4nodePiS2_ii, .-_Z39__device_stub__Z8fillRootRP4nodePiS2_iiPP4nodePiS2_ii
	.globl	_Z8fillRootRP4nodePiS2_ii
	.type	_Z8fillRootRP4nodePiS2_ii, @function
_Z8fillRootRP4nodePiS2_ii:
.LFB2069:
	.cfi_startproc
	endbr64
	jmp	_Z39__device_stub__Z8fillRootRP4nodePiS2_iiPP4nodePiS2_ii
	.cfi_endproc
.LFE2069:
	.size	_Z8fillRootRP4nodePiS2_ii, .-_Z8fillRootRP4nodePiS2_ii
	.globl	_Z40__device_stub__Z10fillLeavesRP4nodePiS2_PP4nodePiS2_
	.type	_Z40__device_stub__Z10fillLeavesRP4nodePiS2_PP4nodePiS2_, @function
_Z40__device_stub__Z10fillLeavesRP4nodePiS2_PP4nodePiS2_:
.LFB2070:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 24(%rsp)
	leaq	48(%rsp), %rcx
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	40(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 80(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 88(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L47
	pushq	48(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z10fillLeavesRP4nodePiS2_(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 144
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	96(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L47:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L49
	call	__stack_chk_fail@PLT
.L49:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2070:
	.size	_Z40__device_stub__Z10fillLeavesRP4nodePiS2_PP4nodePiS2_, .-_Z40__device_stub__Z10fillLeavesRP4nodePiS2_PP4nodePiS2_
	.globl	_Z10fillLeavesRP4nodePiS2_
	.type	_Z10fillLeavesRP4nodePiS2_, @function
_Z10fillLeavesRP4nodePiS2_:
.LFB2071:
	.cfi_startproc
	endbr64
	jmp	_Z40__device_stub__Z10fillLeavesRP4nodePiS2_PP4nodePiS2_
	.cfi_endproc
.LFE2071:
	.size	_Z10fillLeavesRP4nodePiS2_, .-_Z10fillLeavesRP4nodePiS2_
	.globl	_Z37__device_stub__Z8bulkLoadRP4nodePiS2_PP4nodePiS2_
	.type	_Z37__device_stub__Z8bulkLoadRP4nodePiS2_PP4nodePiS2_, @function
_Z37__device_stub__Z8bulkLoadRP4nodePiS2_PP4nodePiS2_:
.LFB2072:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 24(%rsp)
	leaq	48(%rsp), %rcx
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	40(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 80(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 88(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L52
	pushq	48(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z8bulkLoadRP4nodePiS2_(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 144
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	96(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L52:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L54
	call	__stack_chk_fail@PLT
.L54:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2072:
	.size	_Z37__device_stub__Z8bulkLoadRP4nodePiS2_PP4nodePiS2_, .-_Z37__device_stub__Z8bulkLoadRP4nodePiS2_PP4nodePiS2_
	.globl	_Z8bulkLoadRP4nodePiS2_
	.type	_Z8bulkLoadRP4nodePiS2_, @function
_Z8bulkLoadRP4nodePiS2_:
.LFB2073:
	.cfi_startproc
	endbr64
	jmp	_Z37__device_stub__Z8bulkLoadRP4nodePiS2_PP4nodePiS2_
	.cfi_endproc
.LFE2073:
	.size	_Z8bulkLoadRP4nodePiS2_, .-_Z8bulkLoadRP4nodePiS2_
	.globl	_Z38__device_stub__Z13createNewNodeRP4nodePP4node
	.type	_Z38__device_stub__Z13createNewNodeRP4nodePP4node, @function
_Z38__device_stub__Z13createNewNodeRP4nodePP4node:
.LFB2074:
	.cfi_startproc
	endbr64
	subq	$88, %rsp
	.cfi_def_cfa_offset 96
	movq	%rdi, 8(%rsp)
	leaq	32(%rsp), %rcx
	leaq	24(%rsp), %rdx
	leaq	52(%rsp), %rsi
	leaq	40(%rsp), %rdi
	movq	%fs:40, %rax
	movq	%rax, 72(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 48(%rsp)
	movq	%rax, 64(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 40(%rsp)
	movq	%rax, 52(%rsp)
	movl	$1, 60(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L57
	pushq	32(%rsp)
	.cfi_def_cfa_offset 104
	leaq	_Z13createNewNodeRP4node(%rip), %rdi
	pushq	32(%rsp)
	.cfi_def_cfa_offset 112
	movq	68(%rsp), %rcx
	movl	76(%rsp), %r8d
	movq	56(%rsp), %rsi
	movl	64(%rsp), %edx
	leaq	80(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 104
	popq	%rdx
	.cfi_def_cfa_offset 96
.L57:
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	je	.L59
	call	__stack_chk_fail@PLT
.L59:
	addq	$88, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2074:
	.size	_Z38__device_stub__Z13createNewNodeRP4nodePP4node, .-_Z38__device_stub__Z13createNewNodeRP4nodePP4node
	.globl	_Z13createNewNodeRP4node
	.type	_Z13createNewNodeRP4node, @function
_Z13createNewNodeRP4node:
.LFB2075:
	.cfi_startproc
	endbr64
	jmp	_Z38__device_stub__Z13createNewNodeRP4nodePP4node
	.cfi_endproc
.LFE2075:
	.size	_Z13createNewNodeRP4node, .-_Z13createNewNodeRP4node
	.globl	_Z27__device_stub__Z9insertValii
	.type	_Z27__device_stub__Z9insertValii, @function
_Z27__device_stub__Z9insertValii:
.LFB2076:
	.cfi_startproc
	endbr64
	subq	$88, %rsp
	.cfi_def_cfa_offset 96
	movl	%edi, 12(%rsp)
	leaq	32(%rsp), %rcx
	leaq	24(%rsp), %rdx
	leaq	52(%rsp), %rsi
	leaq	40(%rsp), %rdi
	movq	%fs:40, %rax
	movq	%rax, 72(%rsp)
	xorl	%eax, %eax
	leaq	12(%rsp), %rax
	movl	$1, 48(%rsp)
	movq	%rax, 64(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 40(%rsp)
	movq	%rax, 52(%rsp)
	movl	$1, 60(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L62
	pushq	32(%rsp)
	.cfi_def_cfa_offset 104
	leaq	_Z9insertVali(%rip), %rdi
	pushq	32(%rsp)
	.cfi_def_cfa_offset 112
	movq	68(%rsp), %rcx
	movl	76(%rsp), %r8d
	movq	56(%rsp), %rsi
	movl	64(%rsp), %edx
	leaq	80(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 104
	popq	%rdx
	.cfi_def_cfa_offset 96
.L62:
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	je	.L64
	call	__stack_chk_fail@PLT
.L64:
	addq	$88, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2076:
	.size	_Z27__device_stub__Z9insertValii, .-_Z27__device_stub__Z9insertValii
	.globl	_Z9insertVali
	.type	_Z9insertVali, @function
_Z9insertVali:
.LFB2077:
	.cfi_startproc
	endbr64
	jmp	_Z27__device_stub__Z9insertValii
	.cfi_endproc
.LFE2077:
	.size	_Z9insertVali, .-_Z9insertVali
	.globl	_Z33__device_stub__Z8copyNodeP4nodePiP4nodePi
	.type	_Z33__device_stub__Z8copyNodeP4nodePiP4nodePi, @function
_Z33__device_stub__Z8copyNodeP4nodePiP4nodePi:
.LFB2078:
	.cfi_startproc
	endbr64
	subq	$88, %rsp
	.cfi_def_cfa_offset 96
	movq	%rdi, 8(%rsp)
	leaq	24(%rsp), %rcx
	leaq	16(%rsp), %rdx
	movq	%rsi, (%rsp)
	leaq	32(%rsp), %rdi
	leaq	44(%rsp), %rsi
	movq	%fs:40, %rax
	movq	%rax, 72(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 40(%rsp)
	movq	%rax, 56(%rsp)
	movq	%rsp, %rax
	movq	%rax, 64(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 32(%rsp)
	movq	%rax, 44(%rsp)
	movl	$1, 52(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L67
	pushq	24(%rsp)
	.cfi_def_cfa_offset 104
	leaq	_Z8copyNodeP4nodePi(%rip), %rdi
	pushq	24(%rsp)
	.cfi_def_cfa_offset 112
	movq	60(%rsp), %rcx
	movl	68(%rsp), %r8d
	movq	48(%rsp), %rsi
	movl	56(%rsp), %edx
	leaq	72(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 104
	popq	%rdx
	.cfi_def_cfa_offset 96
.L67:
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	je	.L69
	call	__stack_chk_fail@PLT
.L69:
	addq	$88, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2078:
	.size	_Z33__device_stub__Z8copyNodeP4nodePiP4nodePi, .-_Z33__device_stub__Z8copyNodeP4nodePiP4nodePi
	.globl	_Z8copyNodeP4nodePi
	.type	_Z8copyNodeP4nodePi, @function
_Z8copyNodeP4nodePi:
.LFB2079:
	.cfi_startproc
	endbr64
	jmp	_Z33__device_stub__Z8copyNodeP4nodePiP4nodePi
	.cfi_endproc
.LFE2079:
	.size	_Z8copyNodeP4nodePi, .-_Z8copyNodeP4nodePi
	.globl	_Z28__device_stub__Z8copyNodeiPiiPi
	.type	_Z28__device_stub__Z8copyNodeiPiiPi, @function
_Z28__device_stub__Z8copyNodeiPiiPi:
.LFB2080:
	.cfi_startproc
	endbr64
	subq	$88, %rsp
	.cfi_def_cfa_offset 96
	movl	%edi, 12(%rsp)
	leaq	24(%rsp), %rcx
	leaq	16(%rsp), %rdx
	movq	%rsi, (%rsp)
	leaq	32(%rsp), %rdi
	leaq	44(%rsp), %rsi
	movq	%fs:40, %rax
	movq	%rax, 72(%rsp)
	xorl	%eax, %eax
	leaq	12(%rsp), %rax
	movl	$1, 40(%rsp)
	movq	%rax, 56(%rsp)
	movq	%rsp, %rax
	movq	%rax, 64(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 32(%rsp)
	movq	%rax, 44(%rsp)
	movl	$1, 52(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L72
	pushq	24(%rsp)
	.cfi_def_cfa_offset 104
	leaq	_Z8copyNodeiPi(%rip), %rdi
	pushq	24(%rsp)
	.cfi_def_cfa_offset 112
	movq	60(%rsp), %rcx
	movl	68(%rsp), %r8d
	movq	48(%rsp), %rsi
	movl	56(%rsp), %edx
	leaq	72(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 104
	popq	%rdx
	.cfi_def_cfa_offset 96
.L72:
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	je	.L74
	call	__stack_chk_fail@PLT
.L74:
	addq	$88, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2080:
	.size	_Z28__device_stub__Z8copyNodeiPiiPi, .-_Z28__device_stub__Z8copyNodeiPiiPi
	.globl	_Z8copyNodeiPi
	.type	_Z8copyNodeiPi, @function
_Z8copyNodeiPi:
.LFB2081:
	.cfi_startproc
	endbr64
	jmp	_Z28__device_stub__Z8copyNodeiPiiPi
	.cfi_endproc
.LFE2081:
	.size	_Z8copyNodeiPi, .-_Z8copyNodeiPi
	.globl	_Z26__device_stub__Z8addValueii
	.type	_Z26__device_stub__Z8addValueii, @function
_Z26__device_stub__Z8addValueii:
.LFB2082:
	.cfi_startproc
	endbr64
	subq	$88, %rsp
	.cfi_def_cfa_offset 96
	movl	%edi, 12(%rsp)
	leaq	32(%rsp), %rcx
	leaq	24(%rsp), %rdx
	leaq	52(%rsp), %rsi
	leaq	40(%rsp), %rdi
	movq	%fs:40, %rax
	movq	%rax, 72(%rsp)
	xorl	%eax, %eax
	leaq	12(%rsp), %rax
	movl	$1, 48(%rsp)
	movq	%rax, 64(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 40(%rsp)
	movq	%rax, 52(%rsp)
	movl	$1, 60(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L77
	pushq	32(%rsp)
	.cfi_def_cfa_offset 104
	leaq	_Z8addValuei(%rip), %rdi
	pushq	32(%rsp)
	.cfi_def_cfa_offset 112
	movq	68(%rsp), %rcx
	movl	76(%rsp), %r8d
	movq	56(%rsp), %rsi
	movl	64(%rsp), %edx
	leaq	80(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 104
	popq	%rdx
	.cfi_def_cfa_offset 96
.L77:
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	je	.L79
	call	__stack_chk_fail@PLT
.L79:
	addq	$88, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2082:
	.size	_Z26__device_stub__Z8addValueii, .-_Z26__device_stub__Z8addValueii
	.globl	_Z8addValuei
	.type	_Z8addValuei, @function
_Z8addValuei:
.LFB2083:
	.cfi_startproc
	endbr64
	jmp	_Z26__device_stub__Z8addValueii
	.cfi_endproc
.LFE2083:
	.size	_Z8addValuei, .-_Z8addValuei
	.globl	_Z31__device_stub__Z8addValueP4nodeP4node
	.type	_Z31__device_stub__Z8addValueP4nodeP4node, @function
_Z31__device_stub__Z8addValueP4nodeP4node:
.LFB2084:
	.cfi_startproc
	endbr64
	subq	$88, %rsp
	.cfi_def_cfa_offset 96
	movq	%rdi, 8(%rsp)
	leaq	32(%rsp), %rcx
	leaq	24(%rsp), %rdx
	leaq	52(%rsp), %rsi
	leaq	40(%rsp), %rdi
	movq	%fs:40, %rax
	movq	%rax, 72(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 48(%rsp)
	movq	%rax, 64(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 40(%rsp)
	movq	%rax, 52(%rsp)
	movl	$1, 60(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L82
	pushq	32(%rsp)
	.cfi_def_cfa_offset 104
	leaq	_Z8addValueP4node(%rip), %rdi
	pushq	32(%rsp)
	.cfi_def_cfa_offset 112
	movq	68(%rsp), %rcx
	movl	76(%rsp), %r8d
	movq	56(%rsp), %rsi
	movl	64(%rsp), %edx
	leaq	80(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 104
	popq	%rdx
	.cfi_def_cfa_offset 96
.L82:
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	je	.L84
	call	__stack_chk_fail@PLT
.L84:
	addq	$88, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2084:
	.size	_Z31__device_stub__Z8addValueP4nodeP4node, .-_Z31__device_stub__Z8addValueP4nodeP4node
	.globl	_Z8addValueP4node
	.type	_Z8addValueP4node, @function
_Z8addValueP4node:
.LFB2085:
	.cfi_startproc
	endbr64
	jmp	_Z31__device_stub__Z8addValueP4nodeP4node
	.cfi_endproc
.LFE2085:
	.size	_Z8addValueP4node, .-_Z8addValueP4node
	.globl	_Z38__device_stub__Z13copyToNewNodeRP4nodePP4node
	.type	_Z38__device_stub__Z13copyToNewNodeRP4nodePP4node, @function
_Z38__device_stub__Z13copyToNewNodeRP4nodePP4node:
.LFB2086:
	.cfi_startproc
	endbr64
	subq	$88, %rsp
	.cfi_def_cfa_offset 96
	movq	%rdi, 8(%rsp)
	leaq	32(%rsp), %rcx
	leaq	24(%rsp), %rdx
	leaq	52(%rsp), %rsi
	leaq	40(%rsp), %rdi
	movq	%fs:40, %rax
	movq	%rax, 72(%rsp)
	xorl	%eax, %eax
	leaq	8(%rsp), %rax
	movl	$1, 48(%rsp)
	movq	%rax, 64(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 40(%rsp)
	movq	%rax, 52(%rsp)
	movl	$1, 60(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L87
	pushq	32(%rsp)
	.cfi_def_cfa_offset 104
	leaq	_Z13copyToNewNodeRP4node(%rip), %rdi
	pushq	32(%rsp)
	.cfi_def_cfa_offset 112
	movq	68(%rsp), %rcx
	movl	76(%rsp), %r8d
	movq	56(%rsp), %rsi
	movl	64(%rsp), %edx
	leaq	80(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 104
	popq	%rdx
	.cfi_def_cfa_offset 96
.L87:
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	je	.L89
	call	__stack_chk_fail@PLT
.L89:
	addq	$88, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2086:
	.size	_Z38__device_stub__Z13copyToNewNodeRP4nodePP4node, .-_Z38__device_stub__Z13copyToNewNodeRP4nodePP4node
	.globl	_Z13copyToNewNodeRP4node
	.type	_Z13copyToNewNodeRP4node, @function
_Z13copyToNewNodeRP4node:
.LFB2087:
	.cfi_startproc
	endbr64
	jmp	_Z38__device_stub__Z13copyToNewNodeRP4nodePP4node
	.cfi_endproc
.LFE2087:
	.size	_Z13copyToNewNodeRP4node, .-_Z13copyToNewNodeRP4node
	.globl	_Z33__device_stub__Z12searchBetteriPiiPi
	.type	_Z33__device_stub__Z12searchBetteriPiiPi, @function
_Z33__device_stub__Z12searchBetteriPiiPi:
.LFB2088:
	.cfi_startproc
	endbr64
	subq	$88, %rsp
	.cfi_def_cfa_offset 96
	movl	%edi, 12(%rsp)
	leaq	24(%rsp), %rcx
	leaq	16(%rsp), %rdx
	movq	%rsi, (%rsp)
	leaq	32(%rsp), %rdi
	leaq	44(%rsp), %rsi
	movq	%fs:40, %rax
	movq	%rax, 72(%rsp)
	xorl	%eax, %eax
	leaq	12(%rsp), %rax
	movl	$1, 40(%rsp)
	movq	%rax, 56(%rsp)
	movq	%rsp, %rax
	movq	%rax, 64(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 32(%rsp)
	movq	%rax, 44(%rsp)
	movl	$1, 52(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L92
	pushq	24(%rsp)
	.cfi_def_cfa_offset 104
	leaq	_Z12searchBetteriPi(%rip), %rdi
	pushq	24(%rsp)
	.cfi_def_cfa_offset 112
	movq	60(%rsp), %rcx
	movl	68(%rsp), %r8d
	movq	48(%rsp), %rsi
	movl	56(%rsp), %edx
	leaq	72(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 104
	popq	%rdx
	.cfi_def_cfa_offset 96
.L92:
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	je	.L94
	call	__stack_chk_fail@PLT
.L94:
	addq	$88, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2088:
	.size	_Z33__device_stub__Z12searchBetteriPiiPi, .-_Z33__device_stub__Z12searchBetteriPiiPi
	.globl	_Z12searchBetteriPi
	.type	_Z12searchBetteriPi, @function
_Z12searchBetteriPi:
.LFB2089:
	.cfi_startproc
	endbr64
	jmp	_Z33__device_stub__Z12searchBetteriPiiPi
	.cfi_endproc
.LFE2089:
	.size	_Z12searchBetteriPi, .-_Z12searchBetteriPi
	.globl	_Z35__device_stub__Z14containsBetteriPiiPi
	.type	_Z35__device_stub__Z14containsBetteriPiiPi, @function
_Z35__device_stub__Z14containsBetteriPiiPi:
.LFB2090:
	.cfi_startproc
	endbr64
	subq	$88, %rsp
	.cfi_def_cfa_offset 96
	movl	%edi, 12(%rsp)
	leaq	24(%rsp), %rcx
	leaq	16(%rsp), %rdx
	movq	%rsi, (%rsp)
	leaq	32(%rsp), %rdi
	leaq	44(%rsp), %rsi
	movq	%fs:40, %rax
	movq	%rax, 72(%rsp)
	xorl	%eax, %eax
	leaq	12(%rsp), %rax
	movl	$1, 40(%rsp)
	movq	%rax, 56(%rsp)
	movq	%rsp, %rax
	movq	%rax, 64(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 32(%rsp)
	movq	%rax, 44(%rsp)
	movl	$1, 52(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L97
	pushq	24(%rsp)
	.cfi_def_cfa_offset 104
	leaq	_Z14containsBetteriPi(%rip), %rdi
	pushq	24(%rsp)
	.cfi_def_cfa_offset 112
	movq	60(%rsp), %rcx
	movl	68(%rsp), %r8d
	movq	48(%rsp), %rsi
	movl	56(%rsp), %edx
	leaq	72(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 104
	popq	%rdx
	.cfi_def_cfa_offset 96
.L97:
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	je	.L99
	call	__stack_chk_fail@PLT
.L99:
	addq	$88, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2090:
	.size	_Z35__device_stub__Z14containsBetteriPiiPi, .-_Z35__device_stub__Z14containsBetteriPiiPi
	.globl	_Z14containsBetteriPi
	.type	_Z14containsBetteriPi, @function
_Z14containsBetteriPi:
.LFB2091:
	.cfi_startproc
	endbr64
	jmp	_Z35__device_stub__Z14containsBetteriPiiPi
	.cfi_endproc
.LFE2091:
	.size	_Z14containsBetteriPi, .-_Z14containsBetteriPi
	.globl	_Z26__device_stub__Z6searchiPiiPi
	.type	_Z26__device_stub__Z6searchiPiiPi, @function
_Z26__device_stub__Z6searchiPiiPi:
.LFB2092:
	.cfi_startproc
	endbr64
	subq	$88, %rsp
	.cfi_def_cfa_offset 96
	movl	%edi, 12(%rsp)
	leaq	24(%rsp), %rcx
	leaq	16(%rsp), %rdx
	movq	%rsi, (%rsp)
	leaq	32(%rsp), %rdi
	leaq	44(%rsp), %rsi
	movq	%fs:40, %rax
	movq	%rax, 72(%rsp)
	xorl	%eax, %eax
	leaq	12(%rsp), %rax
	movl	$1, 40(%rsp)
	movq	%rax, 56(%rsp)
	movq	%rsp, %rax
	movq	%rax, 64(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 32(%rsp)
	movq	%rax, 44(%rsp)
	movl	$1, 52(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L102
	pushq	24(%rsp)
	.cfi_def_cfa_offset 104
	leaq	_Z6searchiPi(%rip), %rdi
	pushq	24(%rsp)
	.cfi_def_cfa_offset 112
	movq	60(%rsp), %rcx
	movl	68(%rsp), %r8d
	movq	48(%rsp), %rsi
	movl	56(%rsp), %edx
	leaq	72(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 104
	popq	%rdx
	.cfi_def_cfa_offset 96
.L102:
	movq	72(%rsp), %rax
	subq	%fs:40, %rax
	je	.L104
	call	__stack_chk_fail@PLT
.L104:
	addq	$88, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2092:
	.size	_Z26__device_stub__Z6searchiPiiPi, .-_Z26__device_stub__Z6searchiPiiPi
	.globl	_Z6searchiPi
	.type	_Z6searchiPi, @function
_Z6searchiPi:
.LFB2093:
	.cfi_startproc
	endbr64
	jmp	_Z26__device_stub__Z6searchiPiiPi
	.cfi_endproc
.LFE2093:
	.size	_Z6searchiPi, .-_Z6searchiPi
	.globl	_Z28__device_stub__Z6searchPiS_iPiS_i
	.type	_Z28__device_stub__Z6searchPiS_iPiS_i, @function
_Z28__device_stub__Z6searchPiS_iPiS_i:
.LFB2094:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 24(%rsp)
	leaq	48(%rsp), %rcx
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movl	%edx, 12(%rsp)
	leaq	40(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 80(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 88(%rsp)
	leaq	12(%rsp), %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L107
	pushq	48(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z6searchPiS_i(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 144
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	96(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L107:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L109
	call	__stack_chk_fail@PLT
.L109:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2094:
	.size	_Z28__device_stub__Z6searchPiS_iPiS_i, .-_Z28__device_stub__Z6searchPiS_iPiS_i
	.globl	_Z6searchPiS_i
	.type	_Z6searchPiS_i, @function
_Z6searchPiS_i:
.LFB2095:
	.cfi_startproc
	endbr64
	jmp	_Z28__device_stub__Z6searchPiS_iPiS_i
	.cfi_endproc
.LFE2095:
	.size	_Z6searchPiS_i, .-_Z6searchPiS_i
	.globl	_Z33__device_stub__Z4testRP4nodePiS2_PP4nodePiS2_
	.type	_Z33__device_stub__Z4testRP4nodePiS2_PP4nodePiS2_, @function
_Z33__device_stub__Z4testRP4nodePiS2_PP4nodePiS2_:
.LFB2096:
	.cfi_startproc
	endbr64
	subq	$120, %rsp
	.cfi_def_cfa_offset 128
	movq	%rdi, 24(%rsp)
	leaq	48(%rsp), %rcx
	leaq	56(%rsp), %rdi
	movq	%rsi, 16(%rsp)
	leaq	68(%rsp), %rsi
	movq	%rdx, 8(%rsp)
	leaq	40(%rsp), %rdx
	movq	%fs:40, %rax
	movq	%rax, 104(%rsp)
	xorl	%eax, %eax
	leaq	24(%rsp), %rax
	movl	$1, 64(%rsp)
	movq	%rax, 80(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 88(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 96(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 56(%rsp)
	movq	%rax, 68(%rsp)
	movl	$1, 76(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L112
	pushq	48(%rsp)
	.cfi_def_cfa_offset 136
	leaq	_Z4testRP4nodePiS2_(%rip), %rdi
	pushq	48(%rsp)
	.cfi_def_cfa_offset 144
	movq	84(%rsp), %rcx
	movl	92(%rsp), %r8d
	movq	72(%rsp), %rsi
	movl	80(%rsp), %edx
	leaq	96(%rsp), %r9
	call	_Z16cudaLaunchKernelIcE9cudaErrorPT_4dim3S3_PPvmP11CUstream_st.isra.0
	popq	%rax
	.cfi_def_cfa_offset 136
	popq	%rdx
	.cfi_def_cfa_offset 128
.L112:
	movq	104(%rsp), %rax
	subq	%fs:40, %rax
	je	.L114
	call	__stack_chk_fail@PLT
.L114:
	addq	$120, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2096:
	.size	_Z33__device_stub__Z4testRP4nodePiS2_PP4nodePiS2_, .-_Z33__device_stub__Z4testRP4nodePiS2_PP4nodePiS2_
	.globl	_Z4testRP4nodePiS2_
	.type	_Z4testRP4nodePiS2_, @function
_Z4testRP4nodePiS2_:
.LFB2097:
	.cfi_startproc
	endbr64
	jmp	_Z33__device_stub__Z4testRP4nodePiS2_PP4nodePiS2_
	.cfi_endproc
.LFE2097:
	.size	_Z4testRP4nodePiS2_, .-_Z4testRP4nodePiS2_
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"_Z4testRP4nodePiS2_"
.LC1:
	.string	"_Z6searchPiS_i"
.LC2:
	.string	"_Z6searchiPi"
.LC3:
	.string	"_Z14containsBetteriPi"
.LC4:
	.string	"_Z12searchBetteriPi"
.LC5:
	.string	"_Z13copyToNewNodeRP4node"
.LC6:
	.string	"_Z8addValueP4node"
.LC7:
	.string	"_Z8addValuei"
.LC8:
	.string	"_Z8copyNodeiPi"
.LC9:
	.string	"_Z8copyNodeP4nodePi"
.LC10:
	.string	"_Z9insertVali"
.LC11:
	.string	"_Z13createNewNodeRP4node"
.LC12:
	.string	"_Z8bulkLoadRP4nodePiS2_"
.LC13:
	.string	"_Z10fillLeavesRP4nodePiS2_"
.LC14:
	.string	"_Z8fillRootRP4nodePiS2_ii"
.LC15:
	.string	"_Z9fillLevelRP4nodePiS2_ii"
.LC16:
	.string	"_Z10buildLevelRP4nodePiS2_ii"
.LC17:
	.string	"_Z9buildRootRP4nodePiS2_"
.LC18:
	.string	"_Z11buildLeavesRP4nodePiS2_"
.LC19:
	.string	"tempData"
.LC20:
	.string	"tempData1"
.LC21:
	.string	"root1"
.LC22:
	.string	"globalCurr"
.LC23:
	.string	"globalCurrs"
.LC24:
	.string	"newNode"
.LC25:
	.string	"globalIdx"
.LC26:
	.string	"tempKeys"
.LC27:
	.string	"tempPointers"
.LC28:
	.string	"globalPointerIdx"
.LC29:
	.string	"globalCurr1"
.LC30:
	.string	"globalCurr2"
.LC31:
	.string	"foundChild"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2099:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	leaq	_ZL15__fatDeviceText(%rip), %rdi
	call	__cudaRegisterFatBinary@PLT
	pushq	$0
	.cfi_def_cfa_offset 24
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC0(%rip), %rdx
	movq	%rax, %rdi
	movq	%rax, %rbx
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	leaq	_Z4testRP4nodePiS2_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC1(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z6searchPiS_i(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC2(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z6searchiPi(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC3(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z14containsBetteriPi(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC4(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z12searchBetteriPi(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC5(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z13copyToNewNodeRP4node(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC6(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z8addValueP4node(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC7(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z8addValuei(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	leaq	.LC8(%rip), %rdx
	pushq	$0
	.cfi_def_cfa_offset 24
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	_Z8copyNodeiPi(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC9(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z8copyNodeP4nodePi(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC10(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z9insertVali(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC11(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z13createNewNodeRP4node(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC12(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z8bulkLoadRP4nodePiS2_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC13(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z10fillLeavesRP4nodePiS2_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC14(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z8fillRootRP4nodePiS2_ii(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC15(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z9fillLevelRP4nodePiS2_ii(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	movq	%rbx, %rdi
	leaq	.LC16(%rip), %rdx
	pushq	$0
	.cfi_def_cfa_offset 24
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	_Z10buildLevelRP4nodePiS2_ii(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC17(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z9buildRootRP4nodePiS2_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	xorl	%r9d, %r9d
	orl	$-1, %r8d
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC18(%rip), %rdx
	movq	%rbx, %rdi
	leaq	_Z11buildLeavesRP4nodePiS2_(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 40
	pushq	$0
	.cfi_def_cfa_offset 48
	call	__cudaRegisterFunction@PLT
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	xorl	%r8d, %r8d
	movq	%rbx, %rdi
	pushq	$0
	.cfi_def_cfa_offset 24
	leaq	.LC19(%rip), %rdx
	movl	$8000000, %r9d
	leaq	_ZL8tempData(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 32
	movq	%rdx, %rcx
	call	__cudaRegisterVar@PLT
	popq	%rax
	.cfi_def_cfa_offset 24
	popq	%rdx
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC20(%rip), %rdx
	xorl	%r8d, %r8d
	movq	%rbx, %rdi
	movq	%rdx, %rcx
	movl	$8000000, %r9d
	leaq	_ZL9tempData1(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rcx
	.cfi_def_cfa_offset 24
	popq	%rsi
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC21(%rip), %rdx
	xorl	%r8d, %r8d
	movq	%rbx, %rdi
	movq	%rdx, %rcx
	movl	$8, %r9d
	leaq	_ZL5root1(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rdi
	.cfi_def_cfa_offset 24
	popq	%r8
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC22(%rip), %rdx
	xorl	%r8d, %r8d
	movq	%rbx, %rdi
	movq	%rdx, %rcx
	movl	$8, %r9d
	leaq	_ZL10globalCurr(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%r9
	.cfi_def_cfa_offset 24
	popq	%r10
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC23(%rip), %rdx
	xorl	%r8d, %r8d
	movq	%rbx, %rdi
	movq	%rdx, %rcx
	movl	$4096, %r9d
	leaq	_ZL11globalCurrs(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%r11
	.cfi_def_cfa_offset 24
	popq	%rax
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC24(%rip), %rdx
	xorl	%r8d, %r8d
	movq	%rbx, %rdi
	movq	%rdx, %rcx
	movl	$8, %r9d
	leaq	_ZL7newNode(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rax
	.cfi_def_cfa_offset 24
	popq	%rdx
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC25(%rip), %rdx
	xorl	%r8d, %r8d
	movq	%rbx, %rdi
	movq	%rdx, %rcx
	movl	$4, %r9d
	leaq	_ZL9globalIdx(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rcx
	.cfi_def_cfa_offset 24
	popq	%rsi
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC26(%rip), %rdx
	xorl	%r8d, %r8d
	movq	%rbx, %rdi
	movq	%rdx, %rcx
	movl	$2048, %r9d
	leaq	_ZL8tempKeys(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rdi
	.cfi_def_cfa_offset 24
	popq	%r8
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC27(%rip), %rdx
	xorl	%r8d, %r8d
	movq	%rbx, %rdi
	movq	%rdx, %rcx
	movl	$4096, %r9d
	leaq	_ZL12tempPointers(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%r9
	.cfi_def_cfa_offset 24
	popq	%r10
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC28(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$4, %r9d
	leaq	_ZL16globalPointerIdx(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%r11
	.cfi_def_cfa_offset 24
	popq	%rax
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC29(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$8, %r9d
	leaq	_ZL11globalCurr1(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rax
	.cfi_def_cfa_offset 24
	popq	%rdx
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC30(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movq	%rdx, %rcx
	movl	$8, %r9d
	leaq	_ZL11globalCurr2(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rcx
	.cfi_def_cfa_offset 24
	popq	%rsi
	.cfi_def_cfa_offset 16
	pushq	$0
	.cfi_def_cfa_offset 24
	pushq	$0
	.cfi_def_cfa_offset 32
	leaq	.LC31(%rip), %rdx
	movq	%rbx, %rdi
	xorl	%r8d, %r8d
	movl	$8, %r9d
	movq	%rdx, %rcx
	leaq	_ZL10foundChild(%rip), %rsi
	call	__cudaRegisterVar@PLT
	popq	%rdi
	.cfi_def_cfa_offset 24
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	popq	%r8
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rbx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE2099:
	.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
	.local	_ZL10foundChild
	.comm	_ZL10foundChild,8,8
	.local	_ZL11globalCurr2
	.comm	_ZL11globalCurr2,8,8
	.local	_ZL11globalCurr1
	.comm	_ZL11globalCurr1,8,8
	.local	_ZL16globalPointerIdx
	.comm	_ZL16globalPointerIdx,4,4
	.local	_ZL12tempPointers
	.comm	_ZL12tempPointers,4096,32
	.local	_ZL8tempKeys
	.comm	_ZL8tempKeys,2048,32
	.local	_ZL9globalIdx
	.comm	_ZL9globalIdx,4,4
	.local	_ZL7newNode
	.comm	_ZL7newNode,8,8
	.local	_ZL11globalCurrs
	.comm	_ZL11globalCurrs,4096,32
	.local	_ZL10globalCurr
	.comm	_ZL10globalCurr,8,8
	.local	_ZL5root1
	.comm	_ZL5root1,8,8
	.local	_ZL9tempData1
	.comm	_ZL9tempData1,8000000,32
	.local	_ZL8tempData
	.comm	_ZL8tempData,8000000,32
	.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:

/*
 * tree.cuh
 *
 *  Created on: Sep 12, 2016
 *      Author: piotr
 */

#ifndef TREE_CUH_
#define TREE_CUH_
#define ORDER 512
#define INPUTSIZE 1000000
#define BLOCKSIZE 1024

typedef struct node {
	node ** pointers;
	int * keys;
	struct node * parent;
	bool is_leaf;
	int num_keys;
	struct node * next;
} node;

__device__ void make_node(node*&);
__device__ void make_leaf(node*&);
__global__ void bulkLoad(node*&, int*, int*);
__global__ void buildLeaves(node*&, int*, int*);
__global__ void fillLeaves(node*&, int*, int*);
__global__ void test(node*&, int*, int*);
__global__ void buildLevel(node*&, int*, int*, int, int);
__global__ void fillLevel(node*&, int*, int*, int, int);
__global__ void fillRoot(node*&, int*, int*, int, int);
__global__ void buildRoot(node*&, int*, int*);
__global__ void search(int, int*);
__global__ void search(int*, int*, int);
__global__ void insertVal(int);
__device__ node* find(int);
__device__ void addKey(node*, int);
__device__ int contains(node*, int);
__global__ void searchBetter(int val, int* result);
__global__ void containsBetter(int val, int* result);
__global__ void copyNode(int val, int* full);
__global__ void copyNode(node* node, int* full);
__global__ void addValue(int val);
__global__ void addValue(node* nn);
__global__ void createNewNode(node*& nnode);
__global__ void copyToNewNode(node*& nnode);
#endif /* TREE_CUH_ */

	.text
	.file	"tree.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
