/*
 * UnityTest.c
 *
 *  Created on: 04/08/2018
 *      Author: roussian
 */


#include <stdio.h>
#include <stdlib.h>
#include "UnityTest.cuh"

__device__ float checkMinHeapProperty(documentTopkList heap, float newScore, int docCurrent,
									  int topk){

	int index = (topk - 1) - heap.padding - threadIdx.x;
	int parent = index;
	float result=0.0; //0 - Min heap is correct  1 - Min heap is not correct

	for (; index > 0; index-=blockDim.x ) {
		do{
			parent = (parent-1)/2;
			if(heap.score[index] < heap.score[parent]){
				result = heap.score[index];
				printf("[HEAP] position - %d | parent's position - %d | newScore - %.2f | docCurrent - %d | padding -%d \n", index, parent, newScore, docCurrent,heap.padding);
				break;
			}
		}while(parent > 0);

		if(result != 0.0)
			break;
	}

	if(result == 0.0){
		index = threadIdx.x;
		int max = (topk - 1) - heap.padding;
		for (; index < max; index+=blockDim.x ) {
			for (int i=index+1; i < max; i++ ) {
				if(heap.id[index] == heap.id[i]){
					printf("[HEAP] Duplicated document in heap %d | score %.2f | initial position %d | duplicated position %d | blockId %d | docCurrent %d\n",
										 heap.id[index], heap.score[index], index, i, blockIdx.x, docCurrent);
					result = heap.id[index];
				}
			}
		}
	}

	return result;
}

__device__ float checkSorting(documentTopkList heap, float *dTopkScoreListGlobal, int *iTopkDocListGlobal, int topk){

	int globalIndex = blockIdx.x * topk + heap.padding;
	int result=0;
	int index,maxIndex;

	if(THREAD_MASTER){
		if(heap.id[0] != iTopkDocListGlobal[globalIndex]){
			printf("[SORTING] First doc is not correct in topk: global index %d | blockIdx.x %d | doc %d | global doc %d !\n",
					blockIdx.x * topk, blockIdx.x,heap.id[0],iTopkDocListGlobal[globalIndex]);
			result = 1;
			return result;
		}
	}

	int gindex = blockIdx.x * topk + threadIdx.x;
	maxIndex = blockIdx.x * topk + heap.padding;
	for (int i = gindex; i < maxIndex; i+=blockDim.x) {
		if(iTopkDocListGlobal[i] != -1 ){
			printf("[SORTING] Error in padding: blockId %d | index %d | document in position %d\n", blockIdx.x, i,iTopkDocListGlobal[i]);
			result = iTopkDocListGlobal[i];
			return result;
		}
	}

    gindex = blockIdx.x * topk + heap.padding + threadIdx.x;
	maxIndex = (blockIdx.x+1) * topk;
	for (int i = gindex; i < maxIndex; i+=blockDim.x) {
		if(iTopkDocListGlobal[i] == -1 ){
			printf("[SORTING] Docs is empty: blockId %d | index %d | document in position %d\n", blockIdx.x, i,iTopkDocListGlobal[i]);
			result = iTopkDocListGlobal[i];
			return result;
		}
	}

	maxIndex = topk * (blockIdx.x + 1);
	index = topk * blockIdx.x + heap.padding + threadIdx.x;
	for (; index < maxIndex; index+=blockDim.x ) {
		for (int i=index+1; i < maxIndex; i++ ) {
			if(iTopkDocListGlobal[index] == iTopkDocListGlobal[i]){
				printf("[SORTING] Duplicate document in global topk list %d - score %.2f \n", iTopkDocListGlobal[index], dTopkScoreListGlobal[index]);
				result = iTopkDocListGlobal[index];
				return result;
			}
		}
	}

	index=threadIdx.x;
	globalIndex = blockIdx.x * topk;
	maxIndex = topk * (blockIdx.x + 1);
	for (; index < topk-heap.padding; index += blockDim.x) {
		int count = 0;
		for (int i = globalIndex; i < maxIndex; ++i) {
			if(heap.id[index] == iTopkDocListGlobal[i]){
				count++;
				if(heap.score[index] != dTopkScoreListGlobal[i]){
					printf("[SORTING] document's score is wrong! doc %d | score %.2f | local index %d | global index %d\n",
							heap.id[index], heap.score[index], index, i);
					result = heap.id[index];
					return result;
				}
				break;
			}
		}

		if(count == 0){
			printf("[SORTING] Document disappeared: doc %d | score %.2f | local index %d\n",
					heap.id[index], heap.score[index], index);
			result = heap.id[index];
			return result;
		}
	}

	index = blockIdx.x * topk + heap.padding + threadIdx.x;
	maxIndex = topk * (blockIdx.x + 1);
	for (; index < maxIndex; index+=blockDim.x ) {
		for (int i=index+1; i < maxIndex; i++ ) {
			if(dTopkScoreListGlobal[index] > dTopkScoreListGlobal[i]){
				printf("[SORTING] BlockIdx %d | Documents are not sorting!!! doc %d (%.2f) is greater than doc %d (%.2f)\n",
						blockIdx.x,iTopkDocListGlobal[index],dTopkScoreListGlobal[index],iTopkDocListGlobal[i],dTopkScoreListGlobal[i]);
				result = iTopkDocListGlobal[index] ;
				return result;
			}
		}
	}

	return result;
}

__device__ float checkMerge_Sorting_Documents(documentTopkList sortingList,int iSkipMerges,
	  	  	  	  	  	  	  	  	  	  	  int iSkipBlocks, int topk){

	float result = 0.0;
	float score;

	for (int i = threadIdx.x; i < topk; i+=blockDim.x) {
		int doc = sortingList.id[i];

		if((doc != -1 && sortingList.score[i] == 0.0) || (doc == -1 && sortingList.score[i] != 0.0) )
			printf("[MERGE] BlockId %d | SkipBlocks %d | SkipMerges %d | Document is inconsistent: doc %d (%.2f - %d)",iSkipBlocks,
					blockIdx.x, iSkipMerges, doc, sortingList.score[i], i);

		if(doc != -1)
			for (int j = i+1; j < topk; ++j) {
				if(sortingList.id[j] != -1 && sortingList.id[j] == doc){
					printf("[MERGE] BlockId %d | SkipBlocks %d | SkipMerges %d | Duplicated Document: doc %d (%.2f - %d) - doc %d (%.2f - %d)\n",
							blockIdx.x, iSkipBlocks, iSkipMerges, sortingList.id[i], sortingList.score[i], i, sortingList.id[j], sortingList.score[j], j);
					result = sortingList.id[i];
					return result;
				}
			}
	}

	for (int i = threadIdx.x; i < topk; i+=blockDim.x) {
		score = sortingList.score[i];
		if(score != 0.0)
			for (int j = i+1; j < topk; ++j) {
				if(sortingList.score[j] != 0.0 && score > sortingList.score[j]){
					printf("[MERGE] BlockId %d | SkipBlocks %d | SkipMerges %d | Documents are not sorting!!! doc %d (%.2f - %d) is greater than doc %d (%.2f - %d)\n",
							blockIdx.x, iSkipBlocks, iSkipMerges, sortingList.id[i], sortingList.score[i], i, sortingList.id[j], sortingList.score[j], j);
					result = sortingList.id[i];
					return result;
				}
			}
	}

	return result;
}

	.file	"tmpxft_002af455_00000000-6_UnityTest.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2032:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2032:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z20checkMinHeapProperty16documentTopkListfii
	.type	_Z20checkMinHeapProperty16documentTopkListfii, @function
_Z20checkMinHeapProperty16documentTopkListfii:
.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	_Z20checkMinHeapProperty16documentTopkListfii, .-_Z20checkMinHeapProperty16documentTopkListfii
	.globl	_Z12checkSorting16documentTopkListPfPii
	.type	_Z12checkSorting16documentTopkListPfPii, @function
_Z12checkSorting16documentTopkListPfPii:
.LFB2028:
	.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
.LFE2028:
	.size	_Z12checkSorting16documentTopkListPfPii, .-_Z12checkSorting16documentTopkListPfPii
	.globl	_Z28checkMerge_Sorting_Documents16documentTopkListiii
	.type	_Z28checkMerge_Sorting_Documents16documentTopkListiii, @function
_Z28checkMerge_Sorting_Documents16documentTopkListiii:
.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	_Z28checkMerge_Sorting_Documents16documentTopkListiii, .-_Z28checkMerge_Sorting_Documents16documentTopkListiii
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2055:
	.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
.LFE2055:
	.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:

/*
 * UnityTest.cuh
 *
 *  Created on: 04/08/2018
 *      Author: roussian
 */

#ifndef UNITYTEST_CUH_
#define UNITYTEST_CUH_
#include "Structs.cuh"

__device__ float checkMinHeapProperty(documentTopkList heap, float newScore, int docCurrent, int topk);

__device__ float checkSorting(documentTopkList heap, float *dTopkScoreListGlobal, int *iTopkDocListGlobal, int topk);

__device__ float checkMerge_Sorting_Documents(documentTopkList sortingList,int iSkipTopkBetweenMerges,
	  	  	  	  	  int iSkipTopkBetweenBlocks, int topk);

#endif /* UNITYTEST_CUH_ */

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