//#include"cuda_helper.cuh"
//
//#include"npp.h"
//#include"nppcore.h"
//#include"nppdefs.h"
//#include"nppi.h"
//#include"npps.h"
//#include"nppversion.h"
//#define NPP_CALL(x){const NppStatus a=(x);if (a!=NPP_SUCCESS){printf("\nNPP Error(err_num=%d) \n", a);cudaDeviceReset();ASSERT(0);}}
//
//#define NUM_STREAMS  4
//typedef unsigned char u8;
//
//#include <npp.h>
//#pragma comment(lib, "cudart.lib")
//#pragma comment(lib, "nppi.lib")
//
//template<class T>
//__host__ void fillHostArray(T* data, const u32 num_elements){
//	for (u32 i = 0; i < num_elements; i++)
//		data[i] = rand() % (num_elements + 1);
//}
//void _nppTest(){
//
//	checkCudaErrors(cudaSetDevice(0));
//	const u32 num_bytes = (1024u * 255u*256)*sizeof(Npp8u);
//
//	Npp8u *host_src_ptr0;
//	Npp8u *host_src_ptr1;
//	Npp8u *host_dst_ptr0[NUM_STREAMS];
//
//
//	Npp8u *device_src_ptr0[NUM_STREAMS];
//	Npp8u *device_src_ptr1[NUM_STREAMS];
//	Npp8u *device_dst_ptr0[NUM_STREAMS];
//	
//
//	checkCudaErrors(cudaHostAlloc(&host_src_ptr0,num_bytes,cudaHostAllocDefault));
//	checkCudaErrors(cudaHostAlloc(&host_src_ptr1, num_bytes, cudaHostAllocDefault));
//	
//
//	cudaStream_t stream[NUM_STREAMS];
//	cudaEvent_t start, stop;
//	checkCudaErrors(cudaEventCreate(&start));
//	checkCudaErrors(cudaEventCreate(&stop));
//	for (u32 str = 0; str < NUM_STREAMS; str++){
//		checkCudaErrors(cudaHostAlloc(&host_dst_ptr0[str], num_bytes, cudaHostAllocDefault));
//		checkCudaErrors(cudaMalloc(&device_src_ptr0[str],num_bytes));
//		checkCudaErrors(cudaMalloc(&device_src_ptr1[str], num_bytes));
//		checkCudaErrors(cudaMalloc(&device_dst_ptr0[str], num_bytes));
//		checkCudaErrors(cudaStreamCreate(&stream[str]));
//	
//	}
//	fillHostArray(host_src_ptr0, num_bytes);
//	fillHostArray(host_src_ptr1, num_bytes);
//
//	checkCudaErrors(cudaEventRecord(start,0));
//
//	for (u32 str = 0; str < NUM_STREAMS; str++){
//	
//		nppSetStream(stream[str]);
//		checkCudaErrors(cudaMemcpyAsync(device_src_ptr0[str], host_src_ptr0,num_bytes,cudaMemcpyHostToDevice,stream[str]));
//		checkCudaErrors(cudaMemcpyAsync(device_src_ptr1[str], host_src_ptr1, num_bytes, cudaMemcpyHostToDevice, stream[str]));
//		nppsXor_8u(device_src_ptr0[str], device_src_ptr1[str], device_dst_ptr0[str],num_bytes);
//
//	}
//	for (u32 str = 0; str < NUM_STREAMS; str++){
//		nppSetStream(stream[str]);
//		checkCudaErrors(cudaMemcpyAsync(host_dst_ptr0[str], device_dst_ptr0[str], num_bytes,cudaMemcpyDeviceToHost,stream[str]));
//		checkCudaErrors(cudaStreamSynchronize(stream[str]));
//	}
//	checkCudaErrors(cudaEventRecord(stop, 0));
//
//	float elapasedTime = 0.f;
//	checkCudaErrors(cudaEventElapsedTime(&elapasedTime,start,stop));
//	printf("> Execute time:%3.1f\n", elapasedTime);
//	for (u32 str = 0; str < NUM_STREAMS; str++){
//		checkCudaErrors(cudaStreamDestroy(stream[str]));
//		checkCudaErrors(cudaFree(device_src_ptr0[str]));
//		checkCudaErrors(cudaFree(device_src_ptr1[str]));
//		checkCudaErrors(cudaFree(device_dst_ptr0[str]));
//		checkCudaErrors(cudaFreeHost(host_dst_ptr0[str]));
//	}
//	checkCudaErrors(cudaFreeHost(host_src_ptr0));
//	checkCudaErrors(cudaFreeHost(host_src_ptr1));
//
//}
//
//int main(int argc, char **argv){
//	printf("> %s Starting...\n\n", argv[0]);
//	srand((unsigned)time(NULL));
//	_nppTest();
//	checkCudaErrors(cudaDeviceReset());
//	EXIT_SUCCESS;
//};
	.file	"tmpxft_0036c563_00000000-6__nppTest.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2010:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2010:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2033:
	.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
.LFE2033:
	.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:


#include <hip/hip_runtime.h>
//#include"cuda_helper.cuh"
//
//#include"npp.h"
//#include"nppcore.h"
//#include"nppdefs.h"
//#include"nppi.h"
//#include"npps.h"
//#include"nppversion.h"
//#define NPP_CALL(x){const NppStatus a=(x);if (a!=NPP_SUCCESS){printf("\nNPP Error(err_num=%d) \n", a);cudaDeviceReset();ASSERT(0);}}
//
//#define NUM_STREAMS  4
//typedef unsigned char u8;
//
//#include <npp.h>
//#pragma comment(lib, "cudart.lib")
//#pragma comment(lib, "nppi.lib")
//
//template<class T>
//__host__ void fillHostArray(T* data, const u32 num_elements){
//	for (u32 i = 0; i < num_elements; i++)
//		data[i] = rand() % (num_elements + 1);
//}
//void _nppTest(){
//
//	checkCudaErrors(cudaSetDevice(0));
//	const u32 num_bytes = (1024u * 255u*256)*sizeof(Npp8u);
//
//	Npp8u *host_src_ptr0;
//	Npp8u *host_src_ptr1;
//	Npp8u *host_dst_ptr0[NUM_STREAMS];
//
//
//	Npp8u *device_src_ptr0[NUM_STREAMS];
//	Npp8u *device_src_ptr1[NUM_STREAMS];
//	Npp8u *device_dst_ptr0[NUM_STREAMS];
//	
//
//	checkCudaErrors(cudaHostAlloc(&host_src_ptr0,num_bytes,cudaHostAllocDefault));
//	checkCudaErrors(cudaHostAlloc(&host_src_ptr1, num_bytes, cudaHostAllocDefault));
//	
//
//	cudaStream_t stream[NUM_STREAMS];
//	cudaEvent_t start, stop;
//	checkCudaErrors(cudaEventCreate(&start));
//	checkCudaErrors(cudaEventCreate(&stop));
//	for (u32 str = 0; str < NUM_STREAMS; str++){
//		checkCudaErrors(cudaHostAlloc(&host_dst_ptr0[str], num_bytes, cudaHostAllocDefault));
//		checkCudaErrors(cudaMalloc(&device_src_ptr0[str],num_bytes));
//		checkCudaErrors(cudaMalloc(&device_src_ptr1[str], num_bytes));
//		checkCudaErrors(cudaMalloc(&device_dst_ptr0[str], num_bytes));
//		checkCudaErrors(cudaStreamCreate(&stream[str]));
//	
//	}
//	fillHostArray(host_src_ptr0, num_bytes);
//	fillHostArray(host_src_ptr1, num_bytes);
//
//	checkCudaErrors(cudaEventRecord(start,0));
//
//	for (u32 str = 0; str < NUM_STREAMS; str++){
//	
//		nppSetStream(stream[str]);
//		checkCudaErrors(cudaMemcpyAsync(device_src_ptr0[str], host_src_ptr0,num_bytes,cudaMemcpyHostToDevice,stream[str]));
//		checkCudaErrors(cudaMemcpyAsync(device_src_ptr1[str], host_src_ptr1, num_bytes, cudaMemcpyHostToDevice, stream[str]));
//		nppsXor_8u(device_src_ptr0[str], device_src_ptr1[str], device_dst_ptr0[str],num_bytes);
//
//	}
//	for (u32 str = 0; str < NUM_STREAMS; str++){
//		nppSetStream(stream[str]);
//		checkCudaErrors(cudaMemcpyAsync(host_dst_ptr0[str], device_dst_ptr0[str], num_bytes,cudaMemcpyDeviceToHost,stream[str]));
//		checkCudaErrors(cudaStreamSynchronize(stream[str]));
//	}
//	checkCudaErrors(cudaEventRecord(stop, 0));
//
//	float elapasedTime = 0.f;
//	checkCudaErrors(cudaEventElapsedTime(&elapasedTime,start,stop));
//	printf("> Execute time:%3.1f\n", elapasedTime);
//	for (u32 str = 0; str < NUM_STREAMS; str++){
//		checkCudaErrors(cudaStreamDestroy(stream[str]));
//		checkCudaErrors(cudaFree(device_src_ptr0[str]));
//		checkCudaErrors(cudaFree(device_src_ptr1[str]));
//		checkCudaErrors(cudaFree(device_dst_ptr0[str]));
//		checkCudaErrors(cudaFreeHost(host_dst_ptr0[str]));
//	}
//	checkCudaErrors(cudaFreeHost(host_src_ptr0));
//	checkCudaErrors(cudaFreeHost(host_src_ptr1));
//
//}
//
//int main(int argc, char **argv){
//	printf("> %s Starting...\n\n", argv[0]);
//	srand((unsigned)time(NULL));
//	_nppTest();
//	checkCudaErrors(cudaDeviceReset());
//	EXIT_SUCCESS;
//};
	.text
	.file	"_nppTest.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
