#include "cuda_RandomForest_Constants.cu"

namespace Bagging{
	__global__ void kernel_entry(paramPack_Kernel params);
	__host__ void cuda_RandomForest_UpdateConstants(void* src);
}

namespace ExtremeCreateNodes{
	__global__ void kernel_entry(paramPack_Kernel params);
	__host__ void cuda_RandomForest_UpdateConstants(void* src);
}

namespace ExtremeFindSplit{
	__global__ void kernel_entry(paramPack_Kernel params);
	__host__ void cuda_RandomForest_UpdateConstants(void* src);
}

namespace ExtremeMakeSplit{
	__global__ void kernel_entry(paramPack_Kernel params);
	__host__ void cuda_RandomForest_UpdateConstants(void* src);
}

namespace Sort{
	__global__ void kernel_entry(paramPack_Kernel params);
	__host__ void cuda_RandomForest_UpdateConstants(void* src);
}

namespace FindSplit{
	__global__ void kernel_entry(paramPack_Kernel params);
	__host__ void cuda_RandomForest_UpdateConstants(void* src);
}

namespace EvaluateSplit{
	__global__ void kernel_entry(paramPack_Kernel params);
	__host__ void cuda_RandomForest_UpdateConstants(void* src);
}

namespace RandomForest_SplitData_Kernel{
	__global__ void kernel_entry(paramPack_Kernel params);
	__host__ void cuda_RandomForest_UpdateConstants(void* src);
}

namespace RandomForest_Kernel_Classify{
	__global__ void kernel_entry(paramPack_Kernel params);
	__host__ void cuda_RandomForest_UpdateConstants(void* src);
}

enum KernelID {KID_Bagging = 0, KID_Build, KID_KeplerBuild, KID_Sort, KID_FindSplit, KID_Split, KID_EvaluateSplit, KID_Classify, KID_ExtremeFindSplit, KID_ExtremeCreateNodes, KID_ExtremeMakeSplit};
extern "C" void cuda_RandomForest_Wrapper_UpdateConstants(void* src, int id){
	switch(id){
	case KID_Bagging:
		Bagging::cuda_RandomForest_UpdateConstants(src);
		break;
	case KID_Sort:
		Sort::cuda_RandomForest_UpdateConstants(src);
		break;
	case KID_FindSplit:
		FindSplit::cuda_RandomForest_UpdateConstants(src);
		break;
	case KID_Split:
		RandomForest_SplitData_Kernel::cuda_RandomForest_UpdateConstants(src);
		break;
	case KID_EvaluateSplit:
		EvaluateSplit::cuda_RandomForest_UpdateConstants(src);
		break;
	case KID_Classify:
		RandomForest_Kernel_Classify::cuda_RandomForest_UpdateConstants(src);
		break;
	case KID_ExtremeFindSplit:
		ExtremeFindSplit::cuda_RandomForest_UpdateConstants(src);
		break;
	case KID_ExtremeCreateNodes:
		ExtremeCreateNodes::cuda_RandomForest_UpdateConstants(src);
		break;
	case KID_ExtremeMakeSplit:
		ExtremeMakeSplit::cuda_RandomForest_UpdateConstants(src);
		break;
	default:
		break;
	}	
}

dim3 getGrid(unsigned int num_threads){
	int gridX = ceil((c_precision(num_threads)/c_precision(thread_group_size)));

	return dim3(gridX,1,1);
}

extern "C" void cuda_RandomForest_Wrapper_Bagging(unsigned int num_threads, void* params){
	dim3 grid = getGrid(num_threads);
    dim3 threads(thread_group_size, 1, 1);
    Bagging::kernel_entry<<< grid, threads >>>(*((Bagging::paramPack_Kernel*)params));
}

extern "C" void cuda_RandomForest_Wrapper_Build(unsigned int num_threads, void* params){
}

extern "C" void cuda_RandomForest_Wrapper_KeplerBuild(unsigned int num_threads, void* params){
}

extern "C" void cuda_RandomForest_Wrapper_Sort(unsigned int num_threads, void* params){
	dim3 grid = getGrid(num_threads);
    dim3 threads(thread_group_size, 1, 1);
    Sort::kernel_entry<<< grid, threads >>>(*((Sort::paramPack_Kernel*)params));
}

extern "C" void cuda_RandomForest_Wrapper_FindSplit(unsigned int num_threads, void* params){
	dim3 grid = getGrid(num_threads);
    dim3 threads(thread_group_size, 1, 1);
    FindSplit::kernel_entry<<< grid, threads >>>(*((FindSplit::paramPack_Kernel*)params));
}

extern "C" void cuda_RandomForest_Wrapper_EvaluateSplit(unsigned int num_threads, void* params){
	dim3 grid = getGrid(num_threads);
    dim3 threads(thread_group_size, 1, 1);
	EvaluateSplit::kernel_entry<<< grid, threads >>>(*((EvaluateSplit::paramPack_Kernel*)params));
}

extern "C" void cuda_RandomForest_Wrapper_SplitData(unsigned int num_threads, void* params){
	dim3 grid = getGrid(num_threads);
    dim3 threads(thread_group_size, 1, 1);
	RandomForest_SplitData_Kernel::kernel_entry<<< grid, threads >>>(*((RandomForest_SplitData_Kernel::paramPack_Kernel*)params));
}

extern "C" void cuda_RandomForest_Wrapper_Classify(unsigned int num_threads, void* params){
	dim3 grid = getGrid(num_threads);
    dim3 threads(thread_group_size, 1, 1);
	RandomForest_Kernel_Classify::kernel_entry<<< grid, threads >>>(*((RandomForest_Kernel_Classify::paramPack_Kernel*)params));
}

extern "C" void cuda_RandomForest_Wrapper_ExtremeFindSplit(unsigned int num_threads, void* params){
	dim3 grid = getGrid(num_threads);
    dim3 threads(thread_group_size, 1, 1);
    ExtremeFindSplit::kernel_entry<<< grid, threads >>>(*((ExtremeFindSplit::paramPack_Kernel*)params));
}

extern "C" void cuda_RandomForest_Wrapper_ExtremeCreateNodes(unsigned int num_threads, void* params){
	dim3 grid = getGrid(num_threads);
    dim3 threads(thread_group_size, 1, 1);
    ExtremeCreateNodes::kernel_entry<<< grid, threads >>>(*((ExtremeCreateNodes::paramPack_Kernel*)params));
}

extern "C" void cuda_RandomForest_Wrapper_ExtremeMakeSplit(unsigned int num_threads, void* params){
	dim3 grid = getGrid(num_threads);
    dim3 threads(thread_group_size, 1, 1);
    ExtremeMakeSplit::kernel_entry<<< grid, threads >>>(*((ExtremeMakeSplit::paramPack_Kernel*)params));
}
	.file	"tmpxft_003349a9_00000000-6_cuda_RandomForest_Wrappers.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2023:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2023:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	cuda_RandomForest_Wrapper_UpdateConstants
	.type	cuda_RandomForest_Wrapper_UpdateConstants, @function
cuda_RandomForest_Wrapper_UpdateConstants:
.LFB2008:
	.cfi_startproc
	endbr64
	cmpl	$10, %esi
	ja	.L2
	leaq	.L5(%rip), %rdx
	movl	%esi, %esi
	movslq	(%rdx,%rsi,4), %rax
	addq	%rdx, %rax
	notrack jmp	*%rax
	.section	.rodata
	.align 4
	.align 4
.L5:
	.long	.L13-.L5
	.long	.L2-.L5
	.long	.L2-.L5
	.long	.L12-.L5
	.long	.L11-.L5
	.long	.L10-.L5
	.long	.L9-.L5
	.long	.L8-.L5
	.long	.L7-.L5
	.long	.L6-.L5
	.long	.L4-.L5
	.text
.L13:
	jmp	_ZN7Bagging33cuda_RandomForest_UpdateConstantsEPv@PLT
.L12:
	jmp	_ZN4Sort33cuda_RandomForest_UpdateConstantsEPv@PLT
.L11:
	jmp	_ZN9FindSplit33cuda_RandomForest_UpdateConstantsEPv@PLT
.L10:
	jmp	_ZN29RandomForest_SplitData_Kernel33cuda_RandomForest_UpdateConstantsEPv@PLT
.L9:
	jmp	_ZN13EvaluateSplit33cuda_RandomForest_UpdateConstantsEPv@PLT
.L8:
	jmp	_ZN28RandomForest_Kernel_Classify33cuda_RandomForest_UpdateConstantsEPv@PLT
.L7:
	jmp	_ZN16ExtremeFindSplit33cuda_RandomForest_UpdateConstantsEPv@PLT
.L6:
	jmp	_ZN18ExtremeCreateNodes33cuda_RandomForest_UpdateConstantsEPv@PLT
.L4:
	jmp	_ZN16ExtremeMakeSplit33cuda_RandomForest_UpdateConstantsEPv@PLT
.L2:
	ret
	.cfi_endproc
.LFE2008:
	.size	cuda_RandomForest_Wrapper_UpdateConstants, .-cuda_RandomForest_Wrapper_UpdateConstants
	.globl	_Z7getGridj
	.type	_Z7getGridj, @function
_Z7getGridj:
.LFB2009:
	.cfi_startproc
	endbr64
	movl	%edi, %edi
	subq	$40, %rsp
	.cfi_def_cfa_offset 48
	cvtsi2ssq	%rdi, %xmm0
	mulss	.LC0(%rip), %xmm0
	call	ceilf@PLT
	cvttss2sil	%xmm0, %eax
	movl	%eax, 20(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 24(%rsp)
	movl	28(%rsp), %edx
	movq	20(%rsp), %rax
	addq	$40, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2009:
	.size	_Z7getGridj, .-_Z7getGridj
	.globl	cuda_RandomForest_Wrapper_Bagging
	.type	cuda_RandomForest_Wrapper_Bagging, @function
cuda_RandomForest_Wrapper_Bagging:
.LFB2010:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movq	%rsi, %rbx
	subq	$32, %rsp
	.cfi_def_cfa_offset 48
	call	_Z7getGridj
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1, %ecx
	movl	%edx, %esi
	movl	$67108865, %edx
	movq	%rax, %rdi
	salq	$6, %rdx
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L19
	subq	$64, %rsp
	.cfi_def_cfa_offset 112
	movl	$16, %ecx
	movq	%rbx, %rsi
	movq	%rsp, %rdi
	rep movsl
	call	_ZN7Bagging12kernel_entryENS_16paramPack_KernelE@PLT
	addq	$64, %rsp
	.cfi_def_cfa_offset 48
.L19:
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2010:
	.size	cuda_RandomForest_Wrapper_Bagging, .-cuda_RandomForest_Wrapper_Bagging
	.globl	cuda_RandomForest_Wrapper_Build
	.type	cuda_RandomForest_Wrapper_Build, @function
cuda_RandomForest_Wrapper_Build:
.LFB2011:
	.cfi_startproc
	endbr64
	ret
	.cfi_endproc
.LFE2011:
	.size	cuda_RandomForest_Wrapper_Build, .-cuda_RandomForest_Wrapper_Build
	.globl	cuda_RandomForest_Wrapper_KeplerBuild
	.type	cuda_RandomForest_Wrapper_KeplerBuild, @function
cuda_RandomForest_Wrapper_KeplerBuild:
.LFB2162:
	.cfi_startproc
	endbr64
	ret
	.cfi_endproc
.LFE2162:
	.size	cuda_RandomForest_Wrapper_KeplerBuild, .-cuda_RandomForest_Wrapper_KeplerBuild
	.globl	cuda_RandomForest_Wrapper_Sort
	.type	cuda_RandomForest_Wrapper_Sort, @function
cuda_RandomForest_Wrapper_Sort:
.LFB2013:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movq	%rsi, %rbx
	subq	$32, %rsp
	.cfi_def_cfa_offset 48
	call	_Z7getGridj
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1, %ecx
	movl	%edx, %esi
	movl	$67108865, %edx
	movq	%rax, %rdi
	salq	$6, %rdx
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L24
	subq	$64, %rsp
	.cfi_def_cfa_offset 112
	movl	$14, %ecx
	movq	%rbx, %rsi
	movq	%rsp, %rdi
	rep movsl
	call	_ZN4Sort12kernel_entryENS_16paramPack_KernelE@PLT
	addq	$64, %rsp
	.cfi_def_cfa_offset 48
.L24:
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2013:
	.size	cuda_RandomForest_Wrapper_Sort, .-cuda_RandomForest_Wrapper_Sort
	.globl	cuda_RandomForest_Wrapper_FindSplit
	.type	cuda_RandomForest_Wrapper_FindSplit, @function
cuda_RandomForest_Wrapper_FindSplit:
.LFB2014:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movq	%rsi, %rbx
	subq	$32, %rsp
	.cfi_def_cfa_offset 48
	call	_Z7getGridj
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1, %ecx
	movl	%edx, %esi
	movl	$67108865, %edx
	movq	%rax, %rdi
	salq	$6, %rdx
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L27
	addq	$-128, %rsp
	.cfi_def_cfa_offset 176
	movl	$30, %ecx
	movq	%rbx, %rsi
	movq	%rsp, %rdi
	rep movsl
	call	_ZN9FindSplit12kernel_entryENS_16paramPack_KernelE@PLT
	subq	$-128, %rsp
	.cfi_def_cfa_offset 48
.L27:
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2014:
	.size	cuda_RandomForest_Wrapper_FindSplit, .-cuda_RandomForest_Wrapper_FindSplit
	.globl	cuda_RandomForest_Wrapper_EvaluateSplit
	.type	cuda_RandomForest_Wrapper_EvaluateSplit, @function
cuda_RandomForest_Wrapper_EvaluateSplit:
.LFB2015:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movq	%rsi, %rbx
	subq	$32, %rsp
	.cfi_def_cfa_offset 48
	call	_Z7getGridj
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1, %ecx
	movl	%edx, %esi
	movl	$67108865, %edx
	movq	%rax, %rdi
	salq	$6, %rdx
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L30
	subq	$144, %rsp
	.cfi_def_cfa_offset 192
	movl	$34, %ecx
	movq	%rbx, %rsi
	movq	%rsp, %rdi
	rep movsl
	call	_ZN13EvaluateSplit12kernel_entryENS_16paramPack_KernelE@PLT
	addq	$144, %rsp
	.cfi_def_cfa_offset 48
.L30:
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2015:
	.size	cuda_RandomForest_Wrapper_EvaluateSplit, .-cuda_RandomForest_Wrapper_EvaluateSplit
	.globl	cuda_RandomForest_Wrapper_SplitData
	.type	cuda_RandomForest_Wrapper_SplitData, @function
cuda_RandomForest_Wrapper_SplitData:
.LFB2016:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movq	%rsi, %rbx
	subq	$32, %rsp
	.cfi_def_cfa_offset 48
	call	_Z7getGridj
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1, %ecx
	movl	%edx, %esi
	movl	$67108865, %edx
	movq	%rax, %rdi
	salq	$6, %rdx
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L33
	subq	$80, %rsp
	.cfi_def_cfa_offset 128
	movl	$20, %ecx
	movq	%rbx, %rsi
	movq	%rsp, %rdi
	rep movsl
	call	_ZN29RandomForest_SplitData_Kernel12kernel_entryENS_16paramPack_KernelE@PLT
	addq	$80, %rsp
	.cfi_def_cfa_offset 48
.L33:
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2016:
	.size	cuda_RandomForest_Wrapper_SplitData, .-cuda_RandomForest_Wrapper_SplitData
	.globl	cuda_RandomForest_Wrapper_Classify
	.type	cuda_RandomForest_Wrapper_Classify, @function
cuda_RandomForest_Wrapper_Classify:
.LFB2017:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movq	%rsi, %rbx
	subq	$32, %rsp
	.cfi_def_cfa_offset 48
	call	_Z7getGridj
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1, %ecx
	movl	%edx, %esi
	movl	$67108865, %edx
	movq	%rax, %rdi
	salq	$6, %rdx
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L36
	subq	$64, %rsp
	.cfi_def_cfa_offset 112
	movl	$14, %ecx
	movq	%rbx, %rsi
	movq	%rsp, %rdi
	rep movsl
	call	_ZN28RandomForest_Kernel_Classify12kernel_entryENS_16paramPack_KernelE@PLT
	addq	$64, %rsp
	.cfi_def_cfa_offset 48
.L36:
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2017:
	.size	cuda_RandomForest_Wrapper_Classify, .-cuda_RandomForest_Wrapper_Classify
	.globl	cuda_RandomForest_Wrapper_ExtremeFindSplit
	.type	cuda_RandomForest_Wrapper_ExtremeFindSplit, @function
cuda_RandomForest_Wrapper_ExtremeFindSplit:
.LFB2018:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movq	%rsi, %rbx
	subq	$32, %rsp
	.cfi_def_cfa_offset 48
	call	_Z7getGridj
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1, %ecx
	movl	%edx, %esi
	movl	$67108865, %edx
	movq	%rax, %rdi
	salq	$6, %rdx
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L39
	subq	$112, %rsp
	.cfi_def_cfa_offset 160
	movl	$28, %ecx
	movq	%rbx, %rsi
	movq	%rsp, %rdi
	rep movsl
	call	_ZN16ExtremeFindSplit12kernel_entryENS_16paramPack_KernelE@PLT
	addq	$112, %rsp
	.cfi_def_cfa_offset 48
.L39:
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2018:
	.size	cuda_RandomForest_Wrapper_ExtremeFindSplit, .-cuda_RandomForest_Wrapper_ExtremeFindSplit
	.globl	cuda_RandomForest_Wrapper_ExtremeCreateNodes
	.type	cuda_RandomForest_Wrapper_ExtremeCreateNodes, @function
cuda_RandomForest_Wrapper_ExtremeCreateNodes:
.LFB2019:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movq	%rsi, %rbx
	subq	$32, %rsp
	.cfi_def_cfa_offset 48
	call	_Z7getGridj
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1, %ecx
	movl	%edx, %esi
	movl	$67108865, %edx
	movq	%rax, %rdi
	salq	$6, %rdx
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L42
	subq	$96, %rsp
	.cfi_def_cfa_offset 144
	movl	$24, %ecx
	movq	%rbx, %rsi
	movq	%rsp, %rdi
	rep movsl
	call	_ZN18ExtremeCreateNodes12kernel_entryENS_16paramPack_KernelE@PLT
	addq	$96, %rsp
	.cfi_def_cfa_offset 48
.L42:
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2019:
	.size	cuda_RandomForest_Wrapper_ExtremeCreateNodes, .-cuda_RandomForest_Wrapper_ExtremeCreateNodes
	.globl	cuda_RandomForest_Wrapper_ExtremeMakeSplit
	.type	cuda_RandomForest_Wrapper_ExtremeMakeSplit, @function
cuda_RandomForest_Wrapper_ExtremeMakeSplit:
.LFB2020:
	.cfi_startproc
	endbr64
	pushq	%rbx
	.cfi_def_cfa_offset 16
	.cfi_offset 3, -16
	movq	%rsi, %rbx
	subq	$32, %rsp
	.cfi_def_cfa_offset 48
	call	_Z7getGridj
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1, %ecx
	movl	%edx, %esi
	movl	$67108865, %edx
	movq	%rax, %rdi
	salq	$6, %rdx
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L45
	subq	$96, %rsp
	.cfi_def_cfa_offset 144
	movl	$24, %ecx
	movq	%rbx, %rsi
	movq	%rsp, %rdi
	rep movsl
	call	_ZN16ExtremeMakeSplit12kernel_entryENS_16paramPack_KernelE@PLT
	addq	$96, %rsp
	.cfi_def_cfa_offset 48
.L45:
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE2020:
	.size	cuda_RandomForest_Wrapper_ExtremeMakeSplit, .-cuda_RandomForest_Wrapper_ExtremeMakeSplit
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2046:
	.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
.LFE2046:
	.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
	.section	.rodata.cst4,"aM",@progbits,4
	.align 4
.LC0:
	.long	1015021568
	.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:


#include <hip/hip_runtime.h>
#include "cuda_RandomForest_Constants.cu"

namespace Bagging{
	__global__ void kernel_entry(paramPack_Kernel params);
	__host__ void cuda_RandomForest_UpdateConstants(void* src);
}

namespace ExtremeCreateNodes{
	__global__ void kernel_entry(paramPack_Kernel params);
	__host__ void cuda_RandomForest_UpdateConstants(void* src);
}

namespace ExtremeFindSplit{
	__global__ void kernel_entry(paramPack_Kernel params);
	__host__ void cuda_RandomForest_UpdateConstants(void* src);
}

namespace ExtremeMakeSplit{
	__global__ void kernel_entry(paramPack_Kernel params);
	__host__ void cuda_RandomForest_UpdateConstants(void* src);
}

namespace Sort{
	__global__ void kernel_entry(paramPack_Kernel params);
	__host__ void cuda_RandomForest_UpdateConstants(void* src);
}

namespace FindSplit{
	__global__ void kernel_entry(paramPack_Kernel params);
	__host__ void cuda_RandomForest_UpdateConstants(void* src);
}

namespace EvaluateSplit{
	__global__ void kernel_entry(paramPack_Kernel params);
	__host__ void cuda_RandomForest_UpdateConstants(void* src);
}

namespace RandomForest_SplitData_Kernel{
	__global__ void kernel_entry(paramPack_Kernel params);
	__host__ void cuda_RandomForest_UpdateConstants(void* src);
}

namespace RandomForest_Kernel_Classify{
	__global__ void kernel_entry(paramPack_Kernel params);
	__host__ void cuda_RandomForest_UpdateConstants(void* src);
}

enum KernelID {KID_Bagging = 0, KID_Build, KID_KeplerBuild, KID_Sort, KID_FindSplit, KID_Split, KID_EvaluateSplit, KID_Classify, KID_ExtremeFindSplit, KID_ExtremeCreateNodes, KID_ExtremeMakeSplit};
extern "C" void cuda_RandomForest_Wrapper_UpdateConstants(void* src, int id){
	switch(id){
	case KID_Bagging:
		Bagging::cuda_RandomForest_UpdateConstants(src);
		break;
	case KID_Sort:
		Sort::cuda_RandomForest_UpdateConstants(src);
		break;
	case KID_FindSplit:
		FindSplit::cuda_RandomForest_UpdateConstants(src);
		break;
	case KID_Split:
		RandomForest_SplitData_Kernel::cuda_RandomForest_UpdateConstants(src);
		break;
	case KID_EvaluateSplit:
		EvaluateSplit::cuda_RandomForest_UpdateConstants(src);
		break;
	case KID_Classify:
		RandomForest_Kernel_Classify::cuda_RandomForest_UpdateConstants(src);
		break;
	case KID_ExtremeFindSplit:
		ExtremeFindSplit::cuda_RandomForest_UpdateConstants(src);
		break;
	case KID_ExtremeCreateNodes:
		ExtremeCreateNodes::cuda_RandomForest_UpdateConstants(src);
		break;
	case KID_ExtremeMakeSplit:
		ExtremeMakeSplit::cuda_RandomForest_UpdateConstants(src);
		break;
	default:
		break;
	}	
}

dim3 getGrid(unsigned int num_threads){
	int gridX = ceil((c_precision(num_threads)/c_precision(thread_group_size)));

	return dim3(gridX,1,1);
}

extern "C" void cuda_RandomForest_Wrapper_Bagging(unsigned int num_threads, void* params){
	dim3 grid = getGrid(num_threads);
    dim3 threads(thread_group_size, 1, 1);
    Bagging::kernel_entry<<< grid, threads >>>(*((Bagging::paramPack_Kernel*)params));
}

extern "C" void cuda_RandomForest_Wrapper_Build(unsigned int num_threads, void* params){
}

extern "C" void cuda_RandomForest_Wrapper_KeplerBuild(unsigned int num_threads, void* params){
}

extern "C" void cuda_RandomForest_Wrapper_Sort(unsigned int num_threads, void* params){
	dim3 grid = getGrid(num_threads);
    dim3 threads(thread_group_size, 1, 1);
    Sort::kernel_entry<<< grid, threads >>>(*((Sort::paramPack_Kernel*)params));
}

extern "C" void cuda_RandomForest_Wrapper_FindSplit(unsigned int num_threads, void* params){
	dim3 grid = getGrid(num_threads);
    dim3 threads(thread_group_size, 1, 1);
    FindSplit::kernel_entry<<< grid, threads >>>(*((FindSplit::paramPack_Kernel*)params));
}

extern "C" void cuda_RandomForest_Wrapper_EvaluateSplit(unsigned int num_threads, void* params){
	dim3 grid = getGrid(num_threads);
    dim3 threads(thread_group_size, 1, 1);
	EvaluateSplit::kernel_entry<<< grid, threads >>>(*((EvaluateSplit::paramPack_Kernel*)params));
}

extern "C" void cuda_RandomForest_Wrapper_SplitData(unsigned int num_threads, void* params){
	dim3 grid = getGrid(num_threads);
    dim3 threads(thread_group_size, 1, 1);
	RandomForest_SplitData_Kernel::kernel_entry<<< grid, threads >>>(*((RandomForest_SplitData_Kernel::paramPack_Kernel*)params));
}

extern "C" void cuda_RandomForest_Wrapper_Classify(unsigned int num_threads, void* params){
	dim3 grid = getGrid(num_threads);
    dim3 threads(thread_group_size, 1, 1);
	RandomForest_Kernel_Classify::kernel_entry<<< grid, threads >>>(*((RandomForest_Kernel_Classify::paramPack_Kernel*)params));
}

extern "C" void cuda_RandomForest_Wrapper_ExtremeFindSplit(unsigned int num_threads, void* params){
	dim3 grid = getGrid(num_threads);
    dim3 threads(thread_group_size, 1, 1);
    ExtremeFindSplit::kernel_entry<<< grid, threads >>>(*((ExtremeFindSplit::paramPack_Kernel*)params));
}

extern "C" void cuda_RandomForest_Wrapper_ExtremeCreateNodes(unsigned int num_threads, void* params){
	dim3 grid = getGrid(num_threads);
    dim3 threads(thread_group_size, 1, 1);
    ExtremeCreateNodes::kernel_entry<<< grid, threads >>>(*((ExtremeCreateNodes::paramPack_Kernel*)params));
}

extern "C" void cuda_RandomForest_Wrapper_ExtremeMakeSplit(unsigned int num_threads, void* params){
	dim3 grid = getGrid(num_threads);
    dim3 threads(thread_group_size, 1, 1);
    ExtremeMakeSplit::kernel_entry<<< grid, threads >>>(*((ExtremeMakeSplit::paramPack_Kernel*)params));
}
	.text
	.file	"cuda_RandomForest_Wrappers.hip"
	.globl	cuda_RandomForest_Wrapper_UpdateConstants # -- Begin function cuda_RandomForest_Wrapper_UpdateConstants
	.type	cuda_RandomForest_Wrapper_UpdateConstants,@function
cuda_RandomForest_Wrapper_UpdateConstants: # @cuda_RandomForest_Wrapper_UpdateConstants
	.cfi_startproc
# %bb.0:
	cmpl	$10, %esi
	ja	.LBB0_11
# %bb.1:
	movl	%esi, %eax
	jmpq	*.LJTI0_0(,%rax,8)
.LBB0_2:
	jmp	_ZN7Bagging33cuda_RandomForest_UpdateConstantsEPv # TAILCALL
.LBB0_4:
	jmp	_ZN9FindSplit33cuda_RandomForest_UpdateConstantsEPv # TAILCALL
.LBB0_7:
	jmp	_ZN28RandomForest_Kernel_Classify33cuda_RandomForest_UpdateConstantsEPv # TAILCALL
.LBB0_3:
	jmp	_ZN4Sort33cuda_RandomForest_UpdateConstantsEPv # TAILCALL
.LBB0_8:
	jmp	_ZN16ExtremeFindSplit33cuda_RandomForest_UpdateConstantsEPv # TAILCALL
.LBB0_9:
	jmp	_ZN18ExtremeCreateNodes33cuda_RandomForest_UpdateConstantsEPv # TAILCALL
.LBB0_5:
	jmp	_ZN29RandomForest_SplitData_Kernel33cuda_RandomForest_UpdateConstantsEPv # TAILCALL
.LBB0_6:
	jmp	_ZN13EvaluateSplit33cuda_RandomForest_UpdateConstantsEPv # TAILCALL
.LBB0_10:
	jmp	_ZN16ExtremeMakeSplit33cuda_RandomForest_UpdateConstantsEPv # TAILCALL
.LBB0_11:
	retq
.Lfunc_end0:
	.size	cuda_RandomForest_Wrapper_UpdateConstants, .Lfunc_end0-cuda_RandomForest_Wrapper_UpdateConstants
	.cfi_endproc
	.section	.rodata,"a",@progbits
	.p2align	3, 0x0
.LJTI0_0:
	.quad	.LBB0_2
	.quad	.LBB0_11
	.quad	.LBB0_11
	.quad	.LBB0_3
	.quad	.LBB0_4
	.quad	.LBB0_5
	.quad	.LBB0_6
	.quad	.LBB0_7
	.quad	.LBB0_8
	.quad	.LBB0_9
	.quad	.LBB0_10
                                        # -- End function
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function _Z7getGridj
.LCPI1_0:
	.long	0x3c800000                      # float 0.015625
	.text
	.globl	_Z7getGridj
	.type	_Z7getGridj,@function
_Z7getGridj:                            # @_Z7getGridj
	.cfi_startproc
# %bb.0:
	pushq	%rax
	.cfi_def_cfa_offset 16
	movl	%edi, %eax
	cvtsi2ss	%rax, %xmm0
	mulss	.LCPI1_0(%rip), %xmm0
	callq	ceilf@PLT
	cvttss2si	%xmm0, %eax
	btsq	$32, %rax
	movl	$1, %edx
	popq	%rcx
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end1:
	.size	_Z7getGridj, .Lfunc_end1-_Z7getGridj
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function cuda_RandomForest_Wrapper_Bagging
.LCPI2_0:
	.long	0x3c800000                      # float 0.015625
	.text
	.globl	cuda_RandomForest_Wrapper_Bagging
	.type	cuda_RandomForest_Wrapper_Bagging,@function
cuda_RandomForest_Wrapper_Bagging:      # @cuda_RandomForest_Wrapper_Bagging
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$64, %rsp
	.cfi_def_cfa_offset 80
	.cfi_offset %rbx, -16
	movl	%edi, %eax
	cvtsi2ss	%rax, %xmm0
	movq	%rsi, %rbx
	mulss	.LCPI2_0(%rip), %xmm0
	callq	ceilf@PLT
	cvttss2si	%xmm0, %edi
	btsq	$32, %rdi
	movabsq	$4294967296, %rdx               # imm = 0x100000000
	orq	$64, %rdx
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB2_2
# %bb.1:
	movups	(%rbx), %xmm0
	movups	16(%rbx), %xmm1
	movups	32(%rbx), %xmm2
	movups	48(%rbx), %xmm3
	movups	%xmm3, 48(%rsp)
	movups	%xmm2, 32(%rsp)
	movups	%xmm1, 16(%rsp)
	movups	%xmm0, (%rsp)
	callq	_ZN7Bagging27__device_stub__kernel_entryENS_16paramPack_KernelE
.LBB2_2:
	addq	$64, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end2:
	.size	cuda_RandomForest_Wrapper_Bagging, .Lfunc_end2-cuda_RandomForest_Wrapper_Bagging
	.cfi_endproc
                                        # -- End function
	.globl	cuda_RandomForest_Wrapper_Build # -- Begin function cuda_RandomForest_Wrapper_Build
	.type	cuda_RandomForest_Wrapper_Build,@function
cuda_RandomForest_Wrapper_Build:        # @cuda_RandomForest_Wrapper_Build
	.cfi_startproc
# %bb.0:
	retq
.Lfunc_end3:
	.size	cuda_RandomForest_Wrapper_Build, .Lfunc_end3-cuda_RandomForest_Wrapper_Build
	.cfi_endproc
                                        # -- End function
	.globl	cuda_RandomForest_Wrapper_KeplerBuild # -- Begin function cuda_RandomForest_Wrapper_KeplerBuild
	.type	cuda_RandomForest_Wrapper_KeplerBuild,@function
cuda_RandomForest_Wrapper_KeplerBuild:  # @cuda_RandomForest_Wrapper_KeplerBuild
	.cfi_startproc
# %bb.0:
	retq
.Lfunc_end4:
	.size	cuda_RandomForest_Wrapper_KeplerBuild, .Lfunc_end4-cuda_RandomForest_Wrapper_KeplerBuild
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function cuda_RandomForest_Wrapper_Sort
.LCPI5_0:
	.long	0x3c800000                      # float 0.015625
	.text
	.globl	cuda_RandomForest_Wrapper_Sort
	.type	cuda_RandomForest_Wrapper_Sort,@function
cuda_RandomForest_Wrapper_Sort:         # @cuda_RandomForest_Wrapper_Sort
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$64, %rsp
	.cfi_def_cfa_offset 80
	.cfi_offset %rbx, -16
	movl	%edi, %eax
	cvtsi2ss	%rax, %xmm0
	movq	%rsi, %rbx
	mulss	.LCPI5_0(%rip), %xmm0
	callq	ceilf@PLT
	cvttss2si	%xmm0, %edi
	btsq	$32, %rdi
	movabsq	$4294967296, %rdx               # imm = 0x100000000
	orq	$64, %rdx
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB5_2
# %bb.1:
	movq	48(%rbx), %rax
	movq	%rax, 48(%rsp)
	movups	(%rbx), %xmm0
	movups	16(%rbx), %xmm1
	movups	32(%rbx), %xmm2
	movups	%xmm2, 32(%rsp)
	movups	%xmm1, 16(%rsp)
	movups	%xmm0, (%rsp)
	callq	_ZN4Sort27__device_stub__kernel_entryENS_16paramPack_KernelE
.LBB5_2:
	addq	$64, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end5:
	.size	cuda_RandomForest_Wrapper_Sort, .Lfunc_end5-cuda_RandomForest_Wrapper_Sort
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function cuda_RandomForest_Wrapper_FindSplit
.LCPI6_0:
	.long	0x3c800000                      # float 0.015625
	.text
	.globl	cuda_RandomForest_Wrapper_FindSplit
	.type	cuda_RandomForest_Wrapper_FindSplit,@function
cuda_RandomForest_Wrapper_FindSplit:    # @cuda_RandomForest_Wrapper_FindSplit
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$128, %rsp
	.cfi_def_cfa_offset 144
	.cfi_offset %rbx, -16
	movl	%edi, %eax
	cvtsi2ss	%rax, %xmm0
	movq	%rsi, %rbx
	mulss	.LCPI6_0(%rip), %xmm0
	callq	ceilf@PLT
	cvttss2si	%xmm0, %edi
	btsq	$32, %rdi
	movabsq	$4294967296, %rdx               # imm = 0x100000000
	orq	$64, %rdx
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB6_2
# %bb.1:
	movl	$15, %ecx
	movq	%rsp, %rdi
	movq	%rbx, %rsi
	rep;movsq (%rsi), %es:(%rdi)
	callq	_ZN9FindSplit27__device_stub__kernel_entryENS_16paramPack_KernelE
.LBB6_2:
	addq	$128, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end6:
	.size	cuda_RandomForest_Wrapper_FindSplit, .Lfunc_end6-cuda_RandomForest_Wrapper_FindSplit
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function cuda_RandomForest_Wrapper_EvaluateSplit
.LCPI7_0:
	.long	0x3c800000                      # float 0.015625
	.text
	.globl	cuda_RandomForest_Wrapper_EvaluateSplit
	.type	cuda_RandomForest_Wrapper_EvaluateSplit,@function
cuda_RandomForest_Wrapper_EvaluateSplit: # @cuda_RandomForest_Wrapper_EvaluateSplit
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$144, %rsp
	.cfi_def_cfa_offset 160
	.cfi_offset %rbx, -16
	movl	%edi, %eax
	cvtsi2ss	%rax, %xmm0
	movq	%rsi, %rbx
	mulss	.LCPI7_0(%rip), %xmm0
	callq	ceilf@PLT
	cvttss2si	%xmm0, %edi
	btsq	$32, %rdi
	movabsq	$4294967296, %rdx               # imm = 0x100000000
	orq	$64, %rdx
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB7_2
# %bb.1:
	movl	$17, %ecx
	movq	%rsp, %rdi
	movq	%rbx, %rsi
	rep;movsq (%rsi), %es:(%rdi)
	callq	_ZN13EvaluateSplit27__device_stub__kernel_entryENS_16paramPack_KernelE
.LBB7_2:
	addq	$144, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end7:
	.size	cuda_RandomForest_Wrapper_EvaluateSplit, .Lfunc_end7-cuda_RandomForest_Wrapper_EvaluateSplit
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function cuda_RandomForest_Wrapper_SplitData
.LCPI8_0:
	.long	0x3c800000                      # float 0.015625
	.text
	.globl	cuda_RandomForest_Wrapper_SplitData
	.type	cuda_RandomForest_Wrapper_SplitData,@function
cuda_RandomForest_Wrapper_SplitData:    # @cuda_RandomForest_Wrapper_SplitData
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$80, %rsp
	.cfi_def_cfa_offset 96
	.cfi_offset %rbx, -16
	movl	%edi, %eax
	cvtsi2ss	%rax, %xmm0
	movq	%rsi, %rbx
	mulss	.LCPI8_0(%rip), %xmm0
	callq	ceilf@PLT
	cvttss2si	%xmm0, %edi
	btsq	$32, %rdi
	movabsq	$4294967296, %rdx               # imm = 0x100000000
	orq	$64, %rdx
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB8_2
# %bb.1:
	movl	$10, %ecx
	movq	%rsp, %rdi
	movq	%rbx, %rsi
	rep;movsq (%rsi), %es:(%rdi)
	callq	_ZN29RandomForest_SplitData_Kernel27__device_stub__kernel_entryENS_16paramPack_KernelE
.LBB8_2:
	addq	$80, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end8:
	.size	cuda_RandomForest_Wrapper_SplitData, .Lfunc_end8-cuda_RandomForest_Wrapper_SplitData
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function cuda_RandomForest_Wrapper_Classify
.LCPI9_0:
	.long	0x3c800000                      # float 0.015625
	.text
	.globl	cuda_RandomForest_Wrapper_Classify
	.type	cuda_RandomForest_Wrapper_Classify,@function
cuda_RandomForest_Wrapper_Classify:     # @cuda_RandomForest_Wrapper_Classify
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$64, %rsp
	.cfi_def_cfa_offset 80
	.cfi_offset %rbx, -16
	movl	%edi, %eax
	cvtsi2ss	%rax, %xmm0
	movq	%rsi, %rbx
	mulss	.LCPI9_0(%rip), %xmm0
	callq	ceilf@PLT
	cvttss2si	%xmm0, %edi
	btsq	$32, %rdi
	movabsq	$4294967296, %rdx               # imm = 0x100000000
	orq	$64, %rdx
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB9_2
# %bb.1:
	movq	48(%rbx), %rax
	movq	%rax, 48(%rsp)
	movups	(%rbx), %xmm0
	movups	16(%rbx), %xmm1
	movups	32(%rbx), %xmm2
	movups	%xmm2, 32(%rsp)
	movups	%xmm1, 16(%rsp)
	movups	%xmm0, (%rsp)
	callq	_ZN28RandomForest_Kernel_Classify27__device_stub__kernel_entryENS_16paramPack_KernelE
.LBB9_2:
	addq	$64, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end9:
	.size	cuda_RandomForest_Wrapper_Classify, .Lfunc_end9-cuda_RandomForest_Wrapper_Classify
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function cuda_RandomForest_Wrapper_ExtremeFindSplit
.LCPI10_0:
	.long	0x3c800000                      # float 0.015625
	.text
	.globl	cuda_RandomForest_Wrapper_ExtremeFindSplit
	.type	cuda_RandomForest_Wrapper_ExtremeFindSplit,@function
cuda_RandomForest_Wrapper_ExtremeFindSplit: # @cuda_RandomForest_Wrapper_ExtremeFindSplit
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$112, %rsp
	.cfi_def_cfa_offset 128
	.cfi_offset %rbx, -16
	movl	%edi, %eax
	cvtsi2ss	%rax, %xmm0
	movq	%rsi, %rbx
	mulss	.LCPI10_0(%rip), %xmm0
	callq	ceilf@PLT
	cvttss2si	%xmm0, %edi
	btsq	$32, %rdi
	movabsq	$4294967296, %rdx               # imm = 0x100000000
	orq	$64, %rdx
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB10_2
# %bb.1:
	movl	$14, %ecx
	movq	%rsp, %rdi
	movq	%rbx, %rsi
	rep;movsq (%rsi), %es:(%rdi)
	callq	_ZN16ExtremeFindSplit27__device_stub__kernel_entryENS_16paramPack_KernelE
.LBB10_2:
	addq	$112, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end10:
	.size	cuda_RandomForest_Wrapper_ExtremeFindSplit, .Lfunc_end10-cuda_RandomForest_Wrapper_ExtremeFindSplit
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function cuda_RandomForest_Wrapper_ExtremeCreateNodes
.LCPI11_0:
	.long	0x3c800000                      # float 0.015625
	.text
	.globl	cuda_RandomForest_Wrapper_ExtremeCreateNodes
	.type	cuda_RandomForest_Wrapper_ExtremeCreateNodes,@function
cuda_RandomForest_Wrapper_ExtremeCreateNodes: # @cuda_RandomForest_Wrapper_ExtremeCreateNodes
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$96, %rsp
	.cfi_def_cfa_offset 112
	.cfi_offset %rbx, -16
	movl	%edi, %eax
	cvtsi2ss	%rax, %xmm0
	movq	%rsi, %rbx
	mulss	.LCPI11_0(%rip), %xmm0
	callq	ceilf@PLT
	cvttss2si	%xmm0, %edi
	btsq	$32, %rdi
	movabsq	$4294967296, %rdx               # imm = 0x100000000
	orq	$64, %rdx
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB11_2
# %bb.1:
	movl	$12, %ecx
	movq	%rsp, %rdi
	movq	%rbx, %rsi
	rep;movsq (%rsi), %es:(%rdi)
	callq	_ZN18ExtremeCreateNodes27__device_stub__kernel_entryENS_16paramPack_KernelE
.LBB11_2:
	addq	$96, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end11:
	.size	cuda_RandomForest_Wrapper_ExtremeCreateNodes, .Lfunc_end11-cuda_RandomForest_Wrapper_ExtremeCreateNodes
	.cfi_endproc
                                        # -- End function
	.section	.rodata.cst4,"aM",@progbits,4
	.p2align	2, 0x0                          # -- Begin function cuda_RandomForest_Wrapper_ExtremeMakeSplit
.LCPI12_0:
	.long	0x3c800000                      # float 0.015625
	.text
	.globl	cuda_RandomForest_Wrapper_ExtremeMakeSplit
	.type	cuda_RandomForest_Wrapper_ExtremeMakeSplit,@function
cuda_RandomForest_Wrapper_ExtremeMakeSplit: # @cuda_RandomForest_Wrapper_ExtremeMakeSplit
	.cfi_startproc
# %bb.0:
	pushq	%rbx
	.cfi_def_cfa_offset 16
	subq	$96, %rsp
	.cfi_def_cfa_offset 112
	.cfi_offset %rbx, -16
	movl	%edi, %eax
	cvtsi2ss	%rax, %xmm0
	movq	%rsi, %rbx
	mulss	.LCPI12_0(%rip), %xmm0
	callq	ceilf@PLT
	cvttss2si	%xmm0, %edi
	btsq	$32, %rdi
	movabsq	$4294967296, %rdx               # imm = 0x100000000
	orq	$64, %rdx
	movl	$1, %esi
	movl	$1, %ecx
	xorl	%r8d, %r8d
	xorl	%r9d, %r9d
	callq	__hipPushCallConfiguration
	testl	%eax, %eax
	jne	.LBB12_2
# %bb.1:
	movl	$12, %ecx
	movq	%rsp, %rdi
	movq	%rbx, %rsi
	rep;movsq (%rsi), %es:(%rdi)
	callq	_ZN16ExtremeMakeSplit27__device_stub__kernel_entryENS_16paramPack_KernelE
.LBB12_2:
	addq	$96, %rsp
	.cfi_def_cfa_offset 16
	popq	%rbx
	.cfi_def_cfa_offset 8
	retq
.Lfunc_end12:
	.size	cuda_RandomForest_Wrapper_ExtremeMakeSplit, .Lfunc_end12-cuda_RandomForest_Wrapper_ExtremeMakeSplit
	.cfi_endproc
                                        # -- End function
	.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
