#ifndef GPUPRIMITIVE_DEF_CU
#define GPUPRIMITIVE_DEF_CU

#include "stdlib.h"
#include <stdio.h>
#include <cuda_runtime.h>

//unsigned int gpuMemSize = 0;

#  define CUDA_SAFE_CALL( call) {                                    \
    cudaError err = call;                                                    \
    if( cudaSuccess != err) {                                                \
        fprintf(stderr, "Cuda error in file '%s' in line %i : %s.\n",        \
                __FILE__, __LINE__, cudaGetErrorString( err) );              \
        exit(EXIT_FAILURE);                                                  \
    } }

#define GPUMALLOC(D_POINTER, SIZE) CUDA_SAFE_CALL( cudaMalloc( D_POINTER, SIZE) )
#define CPUMALLOC(H_POINTER, SIZE) CUDA_SAFE_CALL(cudaMallocHost (H_POINTER, SIZE))

#define CPUFREE(H_POINTER) if(H_POINTER!=NULL) CUDA_SAFE_CALL(cudaFreeHost ((void *)H_POINTER))
#define GPUFREE(D_POINTER) CUDA_SAFE_CALL( cudaFree( D_POINTER) )


#define TOGPU(D_POINTER,H_POINTER, SIZE)  CUDA_SAFE_CALL(cudaMemcpy(D_POINTER,H_POINTER, SIZE, cudaMemcpyHostToDevice))
#define FROMGPU(H_POINTER, D_POINTER, SIZE)  CUDA_SAFE_CALL(cudaMemcpy(H_POINTER, D_POINTER, SIZE, cudaMemcpyDeviceToHost))
#define GPUTOGPU(D_TO, D_FROM, SIZE)  CUDA_SAFE_CALL(cudaMemcpy(D_TO, D_FROM, SIZE, cudaMemcpyDeviceToDevice))

#define GPUTOGPU_CONSTANT(D_TO, D_FROM, SIZE, OFFSET) CUDA_SAFE_CALL(cudaMemcpyToSymbol(D_TO, D_FROM, SIZE, OFFSET,cudaMemcpyDeviceToDevice))


#define SHARED_MEMORY_PER_PROCESSOR (32*1024)

#define NLJ_NUM_PROCESSOR (16)//for GTX
#define NLJ_SHARED_MEM_PER_PROCESSOR (SHARED_MEMORY_PER_PROCESSOR)
#define NLJ_SHARED_MEM (NLJ_SHARED_MEM_PER_PROCESSOR*NLJ_NUM_PROCESSOR)
#define NLJ_MAX_NUM_BLOCK_PER_DIM (32*1024)
#define NLJ_NUM_THREADS_PER_BLOCK 512
#define NLJ_NUM_TUPLE_PER_THREAD 2
#define NLJ_S_BLOCK_SIZE (NLJ_NUM_THREADS_PER_BLOCK*NLJ_NUM_TUPLE_PER_THREAD)
#define NLJ_R_BLOCK_SIZE NLJ_NUM_THREADS_PER_BLOCK


#define PRED_EQUAL2(DATA) (DATA[0]==DATA[1])
#define PRED_EQUAL(V1,V2) (V1==V2)

//ke's definitions.
/////////////////////////////////////////////////////////////////////////defines
#ifdef max
#undef max
#define max(a,b) (((a) > (b)) ? (a) : (b))
#endif
#ifdef min
#undef min
#define min(a,b) (((a) < (b)) ? (a) : (b))
#endif

///////////////////////////////////////general define

#define _charHistOpt 0	//NOTE: if charOpt, then can't have too many (>255) duplicate pid!!



typedef int2 Rec;
struct/* __align__(16)*/ RecRS
{
	unsigned int val;
	int ridR;
	int ridS;
};
//end of Ke's definition.


#define TEST_MAX (1<<30)
#define TEST_MIN (0)

#define SHARED_MEM 1

#define COALESCED 1

//#define OUTPUT_INFO 1

#define BINARY_SEARCH 1

//#define BINARY_SEARCH_HASH 1

#define CONSTANT_BUFFER_SIZE (1024*64)



#endif

	.file	"tmpxft_00218bc9_00000000-6_GPUPrimitive_Def.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2029:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2029:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2052:
	.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
.LFE2052:
	.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 GPUPRIMITIVE_DEF_CU
#define GPUPRIMITIVE_DEF_CU

#include "stdlib.h"
#include <stdio.h>
#include <hip/hip_runtime.h>

//unsigned int gpuMemSize = 0;

#  define CUDA_SAFE_CALL( call) {                                    \
    hipError_t err = call;                                                    \
    if( hipSuccess != err) {                                                \
        fprintf(stderr, "Cuda error in file '%s' in line %i : %s.\n",        \
                __FILE__, __LINE__, hipGetErrorString( err) );              \
        exit(EXIT_FAILURE);                                                  \
    } }

#define GPUMALLOC(D_POINTER, SIZE) CUDA_SAFE_CALL( hipMalloc( D_POINTER, SIZE) )
#define CPUMALLOC(H_POINTER, SIZE) CUDA_SAFE_CALL(hipHostMalloc (H_POINTER, SIZE))

#define CPUFREE(H_POINTER) if(H_POINTER!=NULL) CUDA_SAFE_CALL(hipHostFree ((void *)H_POINTER))
#define GPUFREE(D_POINTER) CUDA_SAFE_CALL( hipFree( D_POINTER) )


#define TOGPU(D_POINTER,H_POINTER, SIZE)  CUDA_SAFE_CALL(hipMemcpy(D_POINTER,H_POINTER, SIZE, hipMemcpyHostToDevice))
#define FROMGPU(H_POINTER, D_POINTER, SIZE)  CUDA_SAFE_CALL(hipMemcpy(H_POINTER, D_POINTER, SIZE, hipMemcpyDeviceToHost))
#define GPUTOGPU(D_TO, D_FROM, SIZE)  CUDA_SAFE_CALL(hipMemcpy(D_TO, D_FROM, SIZE, hipMemcpyDeviceToDevice))

#define GPUTOGPU_CONSTANT(D_TO, D_FROM, SIZE, OFFSET) CUDA_SAFE_CALL(hipMemcpyToSymbol(D_TO, D_FROM, SIZE, OFFSET,hipMemcpyDeviceToDevice))


#define SHARED_MEMORY_PER_PROCESSOR (32*1024)

#define NLJ_NUM_PROCESSOR (16)//for GTX
#define NLJ_SHARED_MEM_PER_PROCESSOR (SHARED_MEMORY_PER_PROCESSOR)
#define NLJ_SHARED_MEM (NLJ_SHARED_MEM_PER_PROCESSOR*NLJ_NUM_PROCESSOR)
#define NLJ_MAX_NUM_BLOCK_PER_DIM (32*1024)
#define NLJ_NUM_THREADS_PER_BLOCK 512
#define NLJ_NUM_TUPLE_PER_THREAD 2
#define NLJ_S_BLOCK_SIZE (NLJ_NUM_THREADS_PER_BLOCK*NLJ_NUM_TUPLE_PER_THREAD)
#define NLJ_R_BLOCK_SIZE NLJ_NUM_THREADS_PER_BLOCK


#define PRED_EQUAL2(DATA) (DATA[0]==DATA[1])
#define PRED_EQUAL(V1,V2) (V1==V2)

//ke's definitions.
/////////////////////////////////////////////////////////////////////////defines
#ifdef max
#undef max
#define max(a,b) (((a) > (b)) ? (a) : (b))
#endif
#ifdef min
#undef min
#define min(a,b) (((a) < (b)) ? (a) : (b))
#endif

///////////////////////////////////////general define

#define _charHistOpt 0	//NOTE: if charOpt, then can't have too many (>255) duplicate pid!!



typedef int2 Rec;
struct/* __align__(16)*/ RecRS
{
	unsigned int val;
	int ridR;
	int ridS;
};
//end of Ke's definition.


#define TEST_MAX (1<<30)
#define TEST_MIN (0)

#define SHARED_MEM 1

#define COALESCED 1

//#define OUTPUT_INFO 1

#define BINARY_SEARCH 1

//#define BINARY_SEARCH_HASH 1

#define CONSTANT_BUFFER_SIZE (1024*64)



#endif

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