/*=================================================================
 *
 *  nvtspmex.cu
 *  Author: Andrew Magis
 *  Calculate TSP scores on the GPU
 *  Inputs: Class 1 data, Class 2 data, N (size of cross-validation)
 *  Outputs: TSP primary scores, TSP secondary scores, TSP upper and lower bounds for CV
 *
 *
 *=================================================================*/

#include <math.h>
//#include "mex.h"
#include <vector>
#include <iostream>
#include "nvtsp.cuh"
using std::cout;
using std::endl;
#ifndef __NVTSP_CU_
#define __NVTSP_CU_
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <cuda.h>
#include <cuda_runtime.h>
#endif
void DisplayDeviceProperties(int device) {

    cudaDeviceProp deviceProp;
    memset(&deviceProp, 0, sizeof (deviceProp));
	
	printf("-----\n");
	
    if (cudaSuccess == cudaGetDeviceProperties(&deviceProp, device)) {
		printf("Device Name\t\t\t\t%s\n", deviceProp.name);
		printf("Total Global Memory\t\t\t%ld KB\n",deviceProp.totalGlobalMem / 1024);
		printf("Maximum threads per block\t\t%d\n", deviceProp.maxThreadsPerBlock);
		printf("Can map to host memory\t\t%d\n", deviceProp.canMapHostMemory);
	    printf("Number processors\t\t%d\n", deviceProp.multiProcessorCount);
	    printf("Compute Mode\t\t%d\n", deviceProp.computeMode);
        printf("Shared Memory per Block\t\t%d\n", deviceProp.sharedMemPerBlock/1024);
        printf("Max Grid Size[0]\t\t%d\n", deviceProp.maxGridSize[0]);	
        printf("Timeout \t\t%d\n", deviceProp.kernelExecTimeoutEnabled);  	
    } else {
        printf("\n%s", cudaGetErrorString(cudaGetLastError()));
    }
	
	printf("------\n");				
		
}

#define THREADS 16
#define ABSMACRO(X) ( ((X)<0) ? (-(X)) : (X) )
#define ABSBINARYMACRO(X) ( ((X)<0) ? (0) : (1) )
#define MINMACRO(X,Y) ( ((X)<(Y)) ? (X) : (Y) )
#define MAXMACRO(X,Y) ( ((X)>(Y)) ? (X) : (Y) )

__global__ void tspKernel(float *d_class1, float *d_class2, unsigned int n1, unsigned int n2, unsigned int m, unsigned int cvn, float *primary, float *secondary, float *lower, float *upper, int *vote) {
	
    float class1_score = 0.f;
	float class2_score = 0.f;
	float class1_rank = 0.f;
	float class2_rank = 0.f;	
	float temp_lower1, temp_lower2;
	float temp_upper1, temp_upper2;
	
	float n1_invert = __fdividef(1.f, (float)n1);
	float n2_invert = __fdividef(1.f, (float)n2);
	
	//We are only building a diagonal matrix here, so return if I am part of the diagonal
	//or below the diagonal
	if ((blockIdx.x*blockDim.x+threadIdx.x) > (blockIdx.y*blockDim.y+threadIdx.y)) {
	
		//Pointers to correct memory location for class1
		float *data1 = &d_class1[(blockIdx.x*blockDim.x + threadIdx.x)];
		float *data2 = &d_class1[(blockIdx.y*blockDim.y + threadIdx.y)];

		for (int i = 0; i < n1*m; i+=m) {
			//if (data1[i] <= data2[i]) {
			//	class1_score += 1.f;
			//}	
			class1_score += signbit(data1[i]-data2[i]);
			class1_rank += (float)(data1[i]-data2[i]);
		}
		temp_lower1 = __fdividef(class1_score-(float)cvn, (float)(n1-cvn));
		temp_upper1 = __fdividef(class1_score, (float)(n1-cvn));
		class1_score = class1_score * n1_invert;
		class1_rank = class1_rank * n1_invert;

		//Pointers to correct memory location for class2
		data1 =  &d_class2[(blockIdx.x*blockDim.x + threadIdx.x)];
		data2 =  &d_class2[(blockIdx.y*blockDim.y + threadIdx.y)];
	
		for (int i = 0; i < n2*m; i+=m) {
			//if (data1[i] <= data2[i]) {
			//	class2_score += 1.f;
			//}		
			class2_score += signbit(data1[i]-data2[i]);
			class2_rank += (float)(data1[i]-data2[i]);
		}
		temp_lower2 = __fdividef(class2_score, (float)(n2-cvn));
		temp_upper2 = __fdividef(class2_score-(float)cvn, (float)(n2-cvn));
		class2_score = class2_score * n2_invert;
		class2_rank = class2_rank * n2_invert;
		
		temp_lower1 = ABSMACRO(temp_lower1 - class2_score);
		temp_lower2 = ABSMACRO(class1_score - temp_lower2); 
		temp_upper1 = ABSMACRO(temp_upper1 - class2_score);
		temp_upper2 = ABSMACRO(class1_score - temp_upper2);
	
	}

	//Write the result to global memory
	primary[(blockIdx.x*blockDim.x + threadIdx.x)*m + (blockIdx.y*blockDim.y + threadIdx.y)] = ABSMACRO(class1_score-class2_score);
	secondary[(blockIdx.x*blockDim.x + threadIdx.x)*m + (blockIdx.y*blockDim.y + threadIdx.y)] = ABSMACRO(class1_rank-class2_rank);
	lower[(blockIdx.x*blockDim.x + threadIdx.x)*m + (blockIdx.y*blockDim.y + threadIdx.y)] = MINMACRO(temp_lower1, temp_lower2);
	upper[(blockIdx.x*blockDim.x + threadIdx.x)*m + (blockIdx.y*blockDim.y + threadIdx.y)] = MAXMACRO(temp_upper1, temp_upper2);
	vote[(blockIdx.x*blockDim.x + threadIdx.x)*m + (blockIdx.y*blockDim.y + threadIdx.y)] = ABSBINARYMACRO(class1_score-class2_score);

}
/**

If I am lucky this is the only bit I'm going to have to rewrite
**/
void nvwrapper( std::vector<double> & data, int dsSize, std::vector<int> & classSizes){

    size_t f_size = sizeof(float);
   // gonna have to pass this in as a parameter, but not interested in doing that  the moment.
	unsigned int cvn = 5;//(unsigned int)(cvn_temp[0]);

	//m is the number of rows (genes)
	//n is the number of chips (samples)
	unsigned int m1 = dsSize;
	unsigned int n1 = classSizes[0];
	unsigned int m2 = dsSize;
	unsigned int n2 = classSizes[1];
	
	//Create a padded m which is multiple of THREADS
	unsigned int m;
	if (m1 % THREADS == 0) {
		m = m1;
	} else {
		m = ((int)(m1 / THREADS) + 1) * THREADS;
	}

    /**Five outputs required (TSP primary scores, TSP secondary scores, lower bounds, upper bounds, vote)
    **/

    
    //build_output_containers
    float ** TSPPrimaryScores, ** TSPSecondaryScores,** lowerbounds,** upperbounds, ** vote;
    float ** help_allocation[5];
   for(int i=0;i<5;i++){
        help_allocation[i] = new float*[m1];
        for(int j=0;j<m1;j++){
            help_allocation[i][j] = new float[m1];
        }
    }
    TSPPrimaryScores = help_allocation[0];
    TSPSecondaryScores=help_allocation[1];
    lowerbounds=help_allocation[2];
    upperbounds=help_allocation[3];
    vote=help_allocation[4];
 
    int gpu_mem;	
    //the size our vectors 
    unsigned long int class_size[2] = {m*n1*f_size, m*n2*f_size};
	unsigned long int result_size = m*m * f_size;

    //sum of gpu memory we need
    unsigned long int necessary_gpu_memory = class_size[0] + class_size[1] + (5*result_size);
    //check  available memory
    availableMemory(0, necessary_gpu_memory);

    //Allocate space on the GPU to store the input data
    float * d_class[2];
    for(int i=0;i<2;i++){
        if( cudaMalloc( (void**)&d_class[i], class_size[i] ) != cudaSuccess){
            cout << "Memory allocating failure on GPU" << endl;
            exit(1);
        }
    }

    //Allocate space on the GPU to store the output data
	float * d_s[5];
    for(int i=0;i<5;i++){
        if( cudaMalloc( (void**)&d_s[i], result_size ) != cudaSuccess){
            cout << "Memory allocating failure on GPU" << endl;
            exit(1);
        }

    } 
   

    //Allocate space on the host to store the input data
    float * h_class[2];
    int data_i = 0;
    for(int i=0;i<2;i++){
        if(cudaMallocHost((void**)&h_class[i], class_size[i] ) != cudaSuccess){
            cout << "Memory allocating failure on HOST-input" << endl;
            exit(1);
        }
        //where data aligns assign data, where padding assign zero
        int h_class_elements = class_size[i]/f_size;
        for(int pad_i=0; pad_i < h_class_elements; pad_i++){
            if(pad_i % m < m1){
                if( data_i >= data.size()){
                    cout << data_i <<endl;
                    cout << pad_i << endl;
                    cout << m1 << endl;
                    cout << m << endl;
                    cout << class_size[i];
                }
                h_class[i][pad_i] = (float)data.at( data_i);
                
                data_i++;
            } else {
                h_class[i][pad_i] = 0.0;
            }
        }
    }

    //Allocate space on the host to store the output data
	float * h_s[5];
    for(int i=0;i<5;i++){
        if( cudaMallocHost( (void**)&h_s[i], result_size ) != cudaSuccess){
            cout << "Memory allocating failure on HOST-output" << endl;
            exit(1);
        }

    } 
						
	//Copy data to the GPU
    for(int i=0;i<2;i++){
	    if ( (cudaMemcpy(d_class[i], h_class[i], class_size[i], cudaMemcpyHostToDevice) != cudaSuccess)){
    		cout << "Error copying data to the GPU." << endl;
            exit(1);
        }
	}	
	//Set the dimension of the blocks and grids
	dim3 dimBlock(THREADS, THREADS);
	dim3 dimGrid(m/THREADS, m/THREADS);
    //here we go	
	tspKernel<<<dimGrid, dimBlock>>>(d_class[0], d_class[1], n1, n2, m, cvn, d_s[0], d_s[1], d_s[2], d_s[3], (int*)d_s[4]);
	//wait for it
    cudaThreadSynchronize();
		
	//Copy the memory back
    for(int i = 0; i<5;i++){
	    if (cudaMemcpy(h_s[i], d_s[i], result_size, cudaMemcpyDeviceToHost) != cudaSuccess){
		    cout << "Error copying data from the GPU.";	
            exit(0);
        }
    }
		
	//Finally, copy the padded array data into the output matrix
    for (int help_i = 0; help_i < 5; help_i++){ 
        float ** out = help_allocation[help_i]; 
        float * padded = h_s[help_i];
        for (int i = 0; i < m1; i++) {
            memcpy(out[i], padded, m1*sizeof(float));
            padded += m;        
        }
    }
	
    //Clear up memory on the device
    for(int i=0;i<5;i++){
	    if(i<2) cudaFree(d_class[i]);
	    cudaFree(d_s[i]); 
    }
	//Clear up memory on the host
    for(int i=0;i<5;i++){
	    if(i<2) cudaFree(h_class[i]);
	    cudaFree(h_s[i]); 
    }

    //TEMP*********************************************
    //just cleanup memory for now, worry about passing back later
    for(int i=0;i<5;i++){
        for(int j=0;j<m1;j++){
            delete [] help_allocation[i][j];
        }
        delete [] help_allocation[i];
    }
 
}

/**
returns the number of bytes of memory available on dev.
FYI, if you only have 1 gpu dev is just 0
**/
void availableMemory(int dev, unsigned long int necessary_gpu_memory){


    size_t available_gpu_memory, total_gpu_memory;
    CUcontext ctx;
    cuInit(dev);

    cuCtxCreate(&ctx, 0, dev);
    cuMemGetInfo(&available_gpu_memory, &total_gpu_memory);
    cuCtxDetach(ctx);
    if(necessary_gpu_memory >available_gpu_memory 	){
        cout << "not enough memory to run this calculation"<<endl;
        cout << "Necessary: ";
        cout << necessary_gpu_memory<<endl;
        cout << "Available: ";
        cout << available_gpu_memory<<endl;
        exit(1);
    }
}

	.file	"tmpxft_0032eff4_00000000-6_nvtsp.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB4003:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE4003:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.align 2
	.type	_ZNSolsEPFRSoS_E.isra.0, @function
_ZNSolsEPFRSoS_E.isra.0:
.LFB4664:
	.cfi_startproc
	jmp	*%rsi
	.cfi_endproc
.LFE4664:
	.size	_ZNSolsEPFRSoS_E.isra.0, .-_ZNSolsEPFRSoS_E.isra.0
	.section	.rodata.str1.1,"aMS",@progbits,1
.LC0:
	.string	"-----\n"
.LC1:
	.string	"Device Name\t\t\t\t%s\n"
.LC2:
	.string	"Total Global Memory\t\t\t%ld KB\n"
.LC3:
	.string	"Maximum threads per block\t\t%d\n"
.LC4:
	.string	"Can map to host memory\t\t%d\n"
.LC5:
	.string	"Number processors\t\t%d\n"
.LC6:
	.string	"Compute Mode\t\t%d\n"
.LC7:
	.string	"Shared Memory per Block\t\t%d\n"
.LC8:
	.string	"Max Grid Size[0]\t\t%d\n"
.LC9:
	.string	"Timeout \t\t%d\n"
.LC10:
	.string	"\n%s"
.LC11:
	.string	"------\n"
	.text
	.globl	_Z23DisplayDevicePropertiesi
	.type	_Z23DisplayDevicePropertiesi, @function
_Z23DisplayDevicePropertiesi:
.LFB3998:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movl	$258, %ecx
	movl	%edi, %ebp
	leaq	.LC0(%rip), %rsi
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	subq	$1048, %rsp
	.cfi_def_cfa_offset 1072
	movq	%fs:40, %rax
	movq	%rax, 1032(%rsp)
	xorl	%eax, %eax
	movq	%rsp, %rbx
	movq	%rbx, %rdi
	rep stosl
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	%ebp, %esi
	movq	%rbx, %rdi
	call	cudaGetDeviceProperties_v2@PLT
	testl	%eax, %eax
	jne	.L4
	movq	%rbx, %rdx
	leaq	.LC1(%rip), %rsi
	movl	$2, %edi
	call	__printf_chk@PLT
	movq	288(%rsp), %rdx
	leaq	.LC2(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	shrq	$10, %rdx
	call	__printf_chk@PLT
	movl	320(%rsp), %edx
	leaq	.LC3(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	400(%rsp), %edx
	leaq	.LC4(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	388(%rsp), %edx
	leaq	.LC5(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	404(%rsp), %edx
	leaq	.LC6(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movq	296(%rsp), %rdx
	leaq	.LC7(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	shrq	$10, %rdx
	call	__printf_chk@PLT
	movl	336(%rsp), %edx
	leaq	.LC8(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	movl	392(%rsp), %edx
	leaq	.LC9(%rip), %rsi
	xorl	%eax, %eax
	movl	$2, %edi
	call	__printf_chk@PLT
	jmp	.L5
.L4:
	call	cudaGetLastError@PLT
	movl	%eax, %edi
	call	cudaGetErrorString@PLT
	leaq	.LC10(%rip), %rsi
	movl	$2, %edi
	movq	%rax, %rdx
	xorl	%eax, %eax
	call	__printf_chk@PLT
.L5:
	xorl	%eax, %eax
	leaq	.LC11(%rip), %rsi
	movl	$2, %edi
	call	__printf_chk@PLT
	movq	1032(%rsp), %rax
	subq	%fs:40, %rax
	je	.L6
	call	__stack_chk_fail@PLT
.L6:
	addq	$1048, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3998:
	.size	_Z23DisplayDevicePropertiesi, .-_Z23DisplayDevicePropertiesi
	.section	.rodata.str1.1
.LC12:
	.string	"not enough memory to run this calculation"
.LC13:
	.string	"Necessary: "
.LC14:
	.string	"Available: "
	.text
	.globl	_Z15availableMemoryim
	.type	_Z15availableMemoryim, @function
_Z15availableMemoryim:
.LFB4000:
	.cfi_startproc
	endbr64
	pushq	%rbp
	.cfi_def_cfa_offset 16
	.cfi_offset 6, -16
	movq	%rsi, %rbp
	pushq	%rbx
	.cfi_def_cfa_offset 24
	.cfi_offset 3, -24
	movl	%edi, %ebx
	subq	$40, %rsp
	.cfi_def_cfa_offset 64
	movq	%fs:40, %rax
	movq	%rax, 24(%rsp)
	xorl	%eax, %eax
	call	cuInit@PLT
	xorl	%esi, %esi
	leaq	16(%rsp), %rdi
	movl	%ebx, %edx
	call	cuCtxCreate_v2@PLT
	movq	%rsp, %rdi
	leaq	8(%rsp), %rsi
	call	cuMemGetInfo_v2@PLT
	movq	16(%rsp), %rdi
	call	cuCtxDetach@PLT
	cmpq	%rbp, (%rsp)
	jnb	.L9
	leaq	_ZSt4cout(%rip), %rbx
	leaq	.LC12(%rip), %rsi
	movq	%rbx, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	leaq	.LC13(%rip), %rsi
	movq	%rbx, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	%rbp, %rsi
	movq	%rbx, %rdi
	call	_ZNSo9_M_insertImEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	leaq	.LC14(%rip), %rsi
	movq	%rbx, %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	(%rsp), %rsi
	movq	%rbx, %rdi
	call	_ZNSo9_M_insertImEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movl	$1, %edi
	call	exit@PLT
.L9:
	movq	24(%rsp), %rax
	subq	%fs:40, %rax
	je	.L11
	call	__stack_chk_fail@PLT
.L11:
	addq	$40, %rsp
	.cfi_def_cfa_offset 24
	popq	%rbx
	.cfi_def_cfa_offset 16
	popq	%rbp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4000:
	.size	_Z15availableMemoryim, .-_Z15availableMemoryim
	.globl	_Z44__device_stub__Z9tspKernelPfS_jjjjS_S_S_S_PiPfS_jjjjS_S_S_S_Pi
	.type	_Z44__device_stub__Z9tspKernelPfS_jjjjS_S_S_S_PiPfS_jjjjS_S_S_S_Pi, @function
_Z44__device_stub__Z9tspKernelPfS_jjjjS_S_S_S_PiPfS_jjjjS_S_S_S_Pi:
.LFB4025:
	.cfi_startproc
	endbr64
	subq	$248, %rsp
	.cfi_def_cfa_offset 256
	movq	256(%rsp), %rax
	movq	%rdi, 72(%rsp)
	leaq	96(%rsp), %rdi
	movq	%rsi, 64(%rsp)
	leaq	108(%rsp), %rsi
	movq	%rax, 40(%rsp)
	movq	264(%rsp), %rax
	movl	%edx, 60(%rsp)
	leaq	80(%rsp), %rdx
	movq	%rax, 32(%rsp)
	movq	272(%rsp), %rax
	movl	%ecx, 56(%rsp)
	leaq	88(%rsp), %rcx
	movq	%rax, 24(%rsp)
	movq	280(%rsp), %rax
	movl	%r8d, 52(%rsp)
	movq	%rax, 16(%rsp)
	movq	288(%rsp), %rax
	movl	%r9d, 48(%rsp)
	movq	%rax, 8(%rsp)
	movq	%fs:40, %rax
	movq	%rax, 232(%rsp)
	xorl	%eax, %eax
	leaq	72(%rsp), %rax
	movq	%rax, 144(%rsp)
	leaq	64(%rsp), %rax
	movq	%rax, 152(%rsp)
	leaq	60(%rsp), %rax
	movq	%rax, 160(%rsp)
	leaq	56(%rsp), %rax
	movq	%rax, 168(%rsp)
	leaq	52(%rsp), %rax
	movq	%rax, 176(%rsp)
	leaq	48(%rsp), %rax
	movq	%rax, 184(%rsp)
	leaq	40(%rsp), %rax
	movq	%rax, 192(%rsp)
	leaq	32(%rsp), %rax
	movq	%rax, 200(%rsp)
	leaq	24(%rsp), %rax
	movq	%rax, 208(%rsp)
	leaq	16(%rsp), %rax
	movq	%rax, 216(%rsp)
	leaq	8(%rsp), %rax
	movq	%rax, 224(%rsp)
	movabsq	$4294967297, %rax
	movq	%rax, 96(%rsp)
	movl	$1, 104(%rsp)
	movq	%rax, 108(%rsp)
	movl	$1, 116(%rsp)
	call	__cudaPopCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L13
	pushq	88(%rsp)
	.cfi_def_cfa_offset 264
	leaq	_Z9tspKernelPfS_jjjjS_S_S_S_Pi(%rip), %rdi
	pushq	88(%rsp)
	.cfi_def_cfa_offset 272
	movq	124(%rsp), %rcx
	movl	132(%rsp), %r8d
	movq	112(%rsp), %rsi
	movl	120(%rsp), %edx
	leaq	160(%rsp), %r9
	call	cudaLaunchKernel@PLT
	popq	%rax
	.cfi_def_cfa_offset 264
	popq	%rdx
	.cfi_def_cfa_offset 256
.L13:
	movq	232(%rsp), %rax
	subq	%fs:40, %rax
	je	.L15
	call	__stack_chk_fail@PLT
.L15:
	addq	$248, %rsp
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE4025:
	.size	_Z44__device_stub__Z9tspKernelPfS_jjjjS_S_S_S_PiPfS_jjjjS_S_S_S_Pi, .-_Z44__device_stub__Z9tspKernelPfS_jjjjS_S_S_S_PiPfS_jjjjS_S_S_S_Pi
	.globl	_Z9tspKernelPfS_jjjjS_S_S_S_Pi
	.type	_Z9tspKernelPfS_jjjjS_S_S_S_Pi, @function
_Z9tspKernelPfS_jjjjS_S_S_S_Pi:
.LFB4026:
	.cfi_startproc
	endbr64
	jmp	_Z44__device_stub__Z9tspKernelPfS_jjjjS_S_S_S_PiPfS_jjjjS_S_S_S_Pi
	.cfi_endproc
.LFE4026:
	.size	_Z9tspKernelPfS_jjjjS_S_S_S_Pi, .-_Z9tspKernelPfS_jjjjS_S_S_S_Pi
	.section	.rodata.str1.1
.LC15:
	.string	"Memory allocating failure on GPU"
.LC16:
	.string	"Memory allocating failure on HOST-input"
.LC17:
	.string	"vector::_M_range_check: __n (which is %zu) >= this->size() (which is %zu)"
.LC19:
	.string	"Memory allocating failure on HOST-output"
.LC20:
	.string	"Error copying data to the GPU."
.LC21:
	.string	"Error copying data from the GPU."
	.text
	.globl	_Z9nvwrapperRSt6vectorIdSaIdEEiRS_IiSaIiEE
	.type	_Z9nvwrapperRSt6vectorIdSaIdEEiRS_IiSaIiEE, @function
_Z9nvwrapperRSt6vectorIdSaIdEEiRS_IiSaIiEE:
.LFB3999:
	.cfi_startproc
	endbr64
	pushq	%r15
	.cfi_def_cfa_offset 16
	.cfi_offset 15, -16
	pushq	%r14
	.cfi_def_cfa_offset 24
	.cfi_offset 14, -24
	pushq	%r13
	.cfi_def_cfa_offset 32
	.cfi_offset 13, -32
	pushq	%r12
	.cfi_def_cfa_offset 40
	.cfi_offset 12, -40
	movq	%rdi, %r12
	pushq	%rbp
	.cfi_def_cfa_offset 48
	.cfi_offset 6, -48
	movl	%esi, %ebp
	pushq	%rbx
	.cfi_def_cfa_offset 56
	.cfi_offset 3, -56
	movl	%esi, %ebx
	subq	$296, %rsp
	.cfi_def_cfa_offset 352
	movq	%fs:40, %rax
	movq	%rax, 280(%rsp)
	movq	(%rdx), %rax
	movl	(%rax), %ecx
	movl	4(%rax), %eax
	movl	%ecx, 40(%rsp)
	movl	%eax, 44(%rsp)
	testl	$15, %esi
	je	.L19
	movl	%esi, %eax
	andl	$-16, %eax
	leal	16(%rax), %ebx
.L19:
	movl	%ebp, %r14d
	xorl	%r13d, %r13d
	leaq	0(,%r14,8), %rax
	movq	%rax, 8(%rsp)
	leaq	0(,%r14,4), %rax
	movq	%rax, 48(%rsp)
.L22:
	movq	8(%rsp), %rdi
	xorl	%r15d, %r15d
	call	_Znam@PLT
	movq	%rax, 160(%rsp,%r13,8)
	movq	%rax, %rcx
.L20:
	cmpl	%ebp, %r15d
	jnb	.L80
	movq	48(%rsp), %rdi
	movq	%rcx, 16(%rsp)
	call	_Znam@PLT
	movq	16(%rsp), %rcx
	movq	%rax, (%rcx,%r15,8)
	incq	%r15
	jmp	.L20
.L80:
	incq	%r13
	cmpq	$5, %r13
	jne	.L22
	movl	40(%rsp), %eax
	xorl	%edi, %edi
	imull	%ebx, %eax
	salq	$2, %rax
	movq	%rax, 8(%rsp)
	movq	%rax, 112(%rsp)
	movl	44(%rsp), %eax
	imull	%ebx, %eax
	salq	$2, %rax
	movq	%rax, 16(%rsp)
	movq	16(%rsp), %rdx
	movq	%rax, 120(%rsp)
	movl	%ebx, %eax
	imull	%ebx, %eax
	leaq	0(,%rax,4), %rcx
	movq	%rcx, 32(%rsp)
	imulq	$20, %rax, %rax
	movq	8(%rsp), %rcx
	leaq	(%rcx,%rdx), %rsi
	addq	%rax, %rsi
	call	_Z15availableMemoryim
	movq	8(%rsp), %rsi
	leaq	128(%rsp), %rdi
	call	cudaMalloc@PLT
	testl	%eax, %eax
	jne	.L76
	movq	16(%rsp), %rsi
	leaq	136(%rsp), %rdi
	call	cudaMalloc@PLT
	testl	%eax, %eax
	jne	.L76
	leaq	200(%rsp), %r13
	leaq	240(%rsp), %r15
.L26:
	movq	32(%rsp), %rsi
	movq	%r13, %rdi
	call	cudaMalloc@PLT
	testl	%eax, %eax
	je	.L25
.L76:
	leaq	.LC15(%rip), %rsi
	jmp	.L77
.L25:
	addq	$8, %r13
	cmpq	%r15, %r13
	jne	.L26
	xorl	%r15d, %r15d
	xorl	%r13d, %r13d
.L35:
	movq	112(%rsp,%r15), %rax
	leaq	144(%rsp,%r15), %rdi
	movq	%rax, %rsi
	movq	%rax, 24(%rsp)
	call	cudaMallocHost@PLT
	testl	%eax, %eax
	je	.L27
	leaq	.LC16(%rip), %rsi
.L77:
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movl	$1, %edi
.L78:
	call	exit@PLT
.L27:
	movq	24(%rsp), %rax
	xorl	%r8d, %r8d
	shrq	$2, %rax
	movl	%eax, 60(%rsp)
.L28:
	movl	%r8d, 56(%rsp)
	cmpl	%r8d, 60(%rsp)
	jle	.L81
	movl	56(%rsp), %eax
	xorl	%edx, %edx
	divl	%ebx
	cmpl	%ebp, %edx
	jnb	.L29
	movq	8(%r12), %rax
	subq	(%r12), %rax
	movslq	%r13d, %r9
	sarq	$3, %rax
	cmpq	%rax, %r9
	jb	.L30
	movl	%r13d, %esi
	leaq	_ZSt4cout(%rip), %rdi
	movq	%r8, 72(%rsp)
	movq	%r9, 64(%rsp)
	call	_ZNSolsEi@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movl	56(%rsp), %esi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZNSolsEi@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movq	%r14, %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZNSo9_M_insertImEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movl	%ebx, %esi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZNSo9_M_insertImEERSoT_@PLT
	movq	_ZSt4endlIcSt11char_traitsIcEERSt13basic_ostreamIT_T0_ES6_@GOTPCREL(%rip), %rsi
	movq	%rax, %rdi
	call	_ZNSolsEPFRSoS_E.isra.0
	movq	24(%rsp), %rsi
	leaq	_ZSt4cout(%rip), %rdi
	call	_ZNSo9_M_insertImEERSoT_@PLT
	movq	72(%rsp), %r8
	movq	64(%rsp), %r9
.L30:
	movq	(%r12), %rax
	movq	8(%r12), %rdx
	subq	%rax, %rdx
	sarq	$3, %rdx
	cmpq	%rdx, %r9
	jb	.L31
	movq	280(%rsp), %rax
	subq	%fs:40, %rax
	jne	.L79
	movq	%r9, %rsi
	leaq	.LC17(%rip), %rdi
	xorl	%eax, %eax
	call	_ZSt24__throw_out_of_range_fmtPKcz@PLT
.L31:
	movq	144(%rsp,%r15), %rdx
	cvtsd2ss	(%rax,%r9,8), %xmm0
	incl	%r13d
	movss	%xmm0, (%rdx,%r8,4)
	jmp	.L33
.L29:
	movq	144(%rsp,%r15), %rax
	movl	$0x00000000, (%rax,%r8,4)
.L33:
	incq	%r8
	jmp	.L28
.L81:
	addq	$8, %r15
	cmpq	$16, %r15
	jne	.L35
	leaq	240(%rsp), %r13
	leaq	280(%rsp), %r14
	movq	%r13, %r12
.L37:
	movq	32(%rsp), %rsi
	movq	%r13, %rdi
	call	cudaMallocHost@PLT
	testl	%eax, %eax
	je	.L36
	leaq	.LC19(%rip), %rsi
	jmp	.L77
.L36:
	addq	$8, %r13
	cmpq	%r14, %r13
	jne	.L37
	movq	8(%rsp), %rdx
	movq	144(%rsp), %rsi
	movl	$1, %ecx
	movq	128(%rsp), %rdi
	call	cudaMemcpy@PLT
	testl	%eax, %eax
	jne	.L38
	movq	16(%rsp), %rdx
	movq	152(%rsp), %rsi
	movl	$1, %ecx
	movq	136(%rsp), %rdi
	call	cudaMemcpy@PLT
	testl	%eax, %eax
	jne	.L38
	movl	%ebx, %eax
	xorl	%r9d, %r9d
	xorl	%r8d, %r8d
	movl	$1, %ecx
	shrl	$4, %eax
	movl	$1, %esi
	movabsq	$68719476752, %rdx
	movl	$1, 96(%rsp)
	movl	%eax, 100(%rsp)
	movl	%eax, 104(%rsp)
	movq	100(%rsp), %rdi
	movq	%rdx, 88(%rsp)
	movl	$1, 108(%rsp)
	call	__cudaPushCallConfiguration@PLT
	testl	%eax, %eax
	jne	.L39
	pushq	%rax
	.cfi_def_cfa_offset 360
	movl	$5, %r9d
	movl	%ebx, %r8d
	pushq	240(%rsp)
	.cfi_def_cfa_offset 368
	pushq	240(%rsp)
	.cfi_def_cfa_offset 376
	pushq	240(%rsp)
	.cfi_def_cfa_offset 384
	pushq	240(%rsp)
	.cfi_def_cfa_offset 392
	pushq	240(%rsp)
	.cfi_def_cfa_offset 400
	movl	92(%rsp), %ecx
	movl	88(%rsp), %edx
	movq	184(%rsp), %rsi
	movq	176(%rsp), %rdi
	call	_Z44__device_stub__Z9tspKernelPfS_jjjjS_S_S_S_PiPfS_jjjjS_S_S_S_Pi
	addq	$48, %rsp
	.cfi_def_cfa_offset 352
	jmp	.L39
.L38:
	leaq	.LC20(%rip), %rsi
	jmp	.L77
.L39:
	call	cudaThreadSynchronize@PLT
	xorl	%r14d, %r14d
	leaq	200(%rsp), %r13
.L42:
	movq	(%r14,%r13), %rsi
	movq	(%r12,%r14), %rdi
	movl	$2, %ecx
	movq	32(%rsp), %rdx
	call	cudaMemcpy@PLT
	testl	%eax, %eax
	je	.L41
	leaq	_ZSt4cout(%rip), %rdi
	leaq	.LC21(%rip), %rsi
	call	_ZStlsISt11char_traitsIcEERSt13basic_ostreamIcT_ES5_PKc@PLT
	xorl	%edi, %edi
	jmp	.L78
.L41:
	addq	$8, %r14
	cmpq	$40, %r14
	jne	.L42
	salq	$2, %rbx
	xorl	%eax, %eax
	leaq	160(%rsp), %r14
.L45:
	movq	(%rax,%r14), %r9
	movq	(%r12,%rax), %r8
	xorl	%edx, %edx
.L43:
	cmpl	%ebp, %edx
	jnb	.L82
	movq	(%r9,%rdx,8), %rdi
	movq	%r8, %rsi
	movq	48(%rsp), %rcx
	addq	%rbx, %r8
	rep movsb
	incq	%rdx
	jmp	.L43
.L82:
	addq	$8, %rax
	cmpq	$40, %rax
	jne	.L45
	xorl	%ebx, %ebx
.L47:
	cmpq	$1, %rbx
	ja	.L46
	movq	128(%rsp,%rbx,8), %rdi
	call	cudaFree@PLT
.L46:
	movq	0(%r13,%rbx,8), %rdi
	incq	%rbx
	call	cudaFree@PLT
	cmpq	$5, %rbx
	jne	.L47
	xorl	%ebx, %ebx
.L49:
	cmpq	$1, %rbx
	ja	.L48
	movq	144(%rsp,%rbx,8), %rdi
	call	cudaFree@PLT
.L48:
	movq	(%r12,%rbx,8), %rdi
	incq	%rbx
	call	cudaFree@PLT
	cmpq	$5, %rbx
	jne	.L49
	xorl	%r12d, %r12d
.L50:
	xorl	%ebx, %ebx
.L55:
	movq	(%r14,%r12,8), %rdi
	cmpl	%ebp, %ebx
	jnb	.L83
	movq	(%rdi,%rbx,8), %rdi
	testq	%rdi, %rdi
	je	.L51
	call	_ZdaPv@PLT
.L51:
	incq	%rbx
	jmp	.L55
.L83:
	testq	%rdi, %rdi
	je	.L53
	call	_ZdaPv@PLT
.L53:
	incq	%r12
	cmpq	$5, %r12
	jne	.L50
	movq	280(%rsp), %rax
	subq	%fs:40, %rax
	je	.L56
.L79:
	call	__stack_chk_fail@PLT
.L56:
	addq	$296, %rsp
	.cfi_def_cfa_offset 56
	popq	%rbx
	.cfi_def_cfa_offset 48
	popq	%rbp
	.cfi_def_cfa_offset 40
	popq	%r12
	.cfi_def_cfa_offset 32
	popq	%r13
	.cfi_def_cfa_offset 24
	popq	%r14
	.cfi_def_cfa_offset 16
	popq	%r15
	.cfi_def_cfa_offset 8
	ret
	.cfi_endproc
.LFE3999:
	.size	_Z9nvwrapperRSt6vectorIdSaIdEEiRS_IiSaIiEE, .-_Z9nvwrapperRSt6vectorIdSaIdEEiRS_IiSaIiEE
	.section	.rodata.str1.1
.LC22:
	.string	"_Z9tspKernelPfS_jjjjS_S_S_S_Pi"
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB4028:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 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	.LC22(%rip), %rdx
	movq	%rax, %rdi
	leaq	_Z9tspKernelPfS_jjjjS_S_S_S_Pi(%rip), %rsi
	pushq	$0
	.cfi_def_cfa_offset 40
	movq	%rdx, %rcx
	pushq	$0
	.cfi_def_cfa_offset 48
	movq	%rax, _ZL20__cudaFatCubinHandle(%rip)
	call	__cudaRegisterFunction@PLT
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	addq	$32, %rsp
	.cfi_def_cfa_offset 16
	call	__cudaRegisterFatBinaryEnd@PLT
	leaq	_ZL26__cudaUnregisterBinaryUtilv(%rip), %rdi
	popq	%rdx
	.cfi_def_cfa_offset 8
	jmp	atexit@PLT
	.cfi_endproc
.LFE4028:
	.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
	.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:

#ifndef NVTSP_CUH
#define NVTSP_CUH
#include <vector>
#include <hip/hip_runtime.h>

void nvwrapper(std::vector<double> & data, int dsSize, std::vector<int> & classSizes );
void availableMemory(int dev, unsigned long int needed_memory);
#endif


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