//#include <stdlib.h>
//#include <stdio.h>
//#include <curand_kernel.h>
//#include "cuda_runtime.h"
//#include "device_launch_parameters.h"
//
////#include "../../common/book.h"
//#include "../../common/cpu_anim.h"
//#include "../../common/common.h"
//#include "../../common/TexUtils.h"
//#include "../../common/PlyBlock.h"






//
//#define DIM 1024
//#define PI 3.1415926535897932f
//#define MAX_TEMP 1.0f
//#define MIN_TEMP 0.0001f
//#define SPEED   0.00005f
//
//
////// this kernel takes in a 2-d array of floats
////// it updates the value-of-interest by a scaled value based
////// on itself and its nearest neighbors
//__global__ void blend_kernel(float *dst, bool dstOut, cudaTextureObject_t texIn, cudaTextureObject_t texOut) {
//	// map from threadIdx/BlockIdx to pixel position
//	int x = threadIdx.x + blockIdx.x * blockDim.x;
//	int y = threadIdx.y + blockIdx.y * blockDim.y;
//	int offset = x + y * blockDim.x * gridDim.x;
//
//	int left = offset - 1;
//	int right = offset + 1;
//	if (x == 0)   left++;
//	if (x == DIM - 1) right--;
//
//	int top = offset - DIM;
//	int bottom = offset + DIM;
//	if (y == 0)   top += DIM;
//	if (y == DIM - 1) bottom -= DIM;
//
//	float   t, l, c, r, b;
//	if (dstOut) {
//		t = tex1Dfetch<float>(texIn, top);
//		l = tex1Dfetch<float>(texIn, left);
//		c = tex1Dfetch<float>(texIn, offset);
//		r = tex1Dfetch<float>(texIn, right);
//		b = tex1Dfetch<float>(texIn, bottom);
//
//	}
//	else {
//		t = tex1Dfetch<float>(texOut, top);
//		l = tex1Dfetch<float>(texOut, left);
//		c = tex1Dfetch<float>(texOut, offset);
//		r = tex1Dfetch<float>(texOut, right);
//		b = tex1Dfetch<float>(texOut, bottom);
//	}
//	dst[offset] = c + SPEED * (t + b + r + l - 4 * c);
//}
//
//
//
//void anim_gpu(AppBlock *appBlock, int ticks) {
//	//HANDLE_ERROR(cudaEventRecord(appBlock->start, 0));
//	//dim3    blocks(DIM / 16, DIM / 16);
//	//dim3    threads(16, 16);
//	//CPUAnimBitmap  *bitmap = appBlock->bitmap;
//
//	//// since tex is global and bound, we have to use a flag to
//	//// select which is in/out per iteration
//	//volatile bool dstOut = true;
//	//for (int i = 0; i<100; i++) {
//	//	float   *in, *out;
//	//	if (dstOut) {
//	//		in = appBlock->d->dev_inSrc;
//	//		out = appBlock->d->dev_outSrc;
//	//	}
//	//	else {
//	//		out = appBlock->d->dev_inSrc;
//	//		in = appBlock->d->dev_outSrc;
//	//	}
//	//	blend_kernel << <blocks, threads >> >(out, dstOut, *appBlock->d->texIn, *appBlock->d->texOut);
//	//	dstOut = !dstOut;
//	//}
//
//	//float_to_color << <blocks, threads >> >(appBlock->d->output_bitmap, appBlock->d->dev_outSrc);
//
//	//HANDLE_ERROR(cudaMemcpy(bitmap->get_ptr(),
//	//	d->output_bitmap,
//	//	bitmap->image_size(),
//	//	cudaMemcpyDeviceToHost));
//
//	//HANDLE_ERROR(cudaEventRecord(appBlock->stop, 0));
//	//HANDLE_ERROR(cudaEventSynchronize(d->stop));
//	//float   elapsedTime;
//	//HANDLE_ERROR(cudaEventElapsedTime(&elapsedTime,
//	//	d->start, d->stop));
//	//d->totalTime += elapsedTime;
//	//++d->frames;
//	//printf("Average Time per frame:  %3.1f ms\n", d->totalTime / d->frames);
//	//printf("tic:  %d\n\n", ticks);
//}
//
//// clean up memory allocated on the GPU
//
//void anim_exit(AppBlock *d) {
//
//	//cudaDestroyTextureObject(*(d->texIn));
//	//cudaDestroyTextureObject(*d->texOut);
//	//cudaDestroyTextureObject(*d->texConst);
//
//	//HANDLE_ERROR(cudaFree(d->dev_inSrc));
//	//HANDLE_ERROR(cudaFree(d->dev_outSrc));
//	//HANDLE_ERROR(cudaFree(d->dev_constSrc));
//
//	//HANDLE_ERROR(cudaEventDestroy(d->start));
//	//HANDLE_ERROR(cudaEventDestroy(d->stop));
//}
//
//
////int main(void)
////{
////	PlyBlock   data;
////	CPUAnimBitmap bitmap(DIM, DIM, &data);
////	data.bitmap = &bitmap;
////	data.totalTime = 0;
////	data.frames = 0;
////	HANDLE_ERROR(cudaEventCreate(&data.start));
////	HANDLE_ERROR(cudaEventCreate(&data.stop));
////
////	int imageSize = bitmap.image_size();
////
////	HANDLE_ERROR(cudaMalloc((void**)&data.output_bitmap,
////		imageSize));
////
////	// assume float == 4 chars in size (ie rgba)
////	HANDLE_ERROR(cudaMalloc((void**)&data.dev_inSrc, imageSize));
////	HANDLE_ERROR(cudaMalloc((void**)&data.dev_outSrc, imageSize));
////	HANDLE_ERROR(cudaMalloc((void**)&data.dev_constSrc, imageSize));
////
////
////	data.texIn = TexObjFloat1D(data.dev_inSrc, imageSize);
////	data.texOut = TexObjFloat1D(data.dev_outSrc, imageSize);
////	data.texConst = TexObjFloat1D(data.dev_constSrc, imageSize);
////
////
////
////	// intialize the constant data
////	float *temp = RndFloat0to1(DIM*DIM);
////
////	HANDLE_ERROR(cudaMemcpy(data.dev_constSrc, temp, imageSize,
////		cudaMemcpyHostToDevice));
////
////
////	HANDLE_ERROR(cudaMemcpy(data.dev_inSrc, temp, imageSize,
////		cudaMemcpyHostToDevice));
////	free(temp);
////
////	bitmap.anim_and_exit((void(*)(void*, int))anim_gpu,
////		(void(*)(void*))anim_exit);
////}

	.file	"tmpxft_003443b4_00000000-6_Isingo.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 <stdlib.h>
//#include <stdio.h>
//#include <curand_kernel.h>
//#include "cuda_runtime.h"
//#include "device_launch_parameters.h"
//
////#include "../../common/book.h"
//#include "../../common/cpu_anim.h"
//#include "../../common/common.h"
//#include "../../common/TexUtils.h"
//#include "../../common/PlyBlock.h"






//
//#define DIM 1024
//#define PI 3.1415926535897932f
//#define MAX_TEMP 1.0f
//#define MIN_TEMP 0.0001f
//#define SPEED   0.00005f
//
//
////// this kernel takes in a 2-d array of floats
////// it updates the value-of-interest by a scaled value based
////// on itself and its nearest neighbors
//__global__ void blend_kernel(float *dst, bool dstOut, cudaTextureObject_t texIn, cudaTextureObject_t texOut) {
//	// map from threadIdx/BlockIdx to pixel position
//	int x = threadIdx.x + blockIdx.x * blockDim.x;
//	int y = threadIdx.y + blockIdx.y * blockDim.y;
//	int offset = x + y * blockDim.x * gridDim.x;
//
//	int left = offset - 1;
//	int right = offset + 1;
//	if (x == 0)   left++;
//	if (x == DIM - 1) right--;
//
//	int top = offset - DIM;
//	int bottom = offset + DIM;
//	if (y == 0)   top += DIM;
//	if (y == DIM - 1) bottom -= DIM;
//
//	float   t, l, c, r, b;
//	if (dstOut) {
//		t = tex1Dfetch<float>(texIn, top);
//		l = tex1Dfetch<float>(texIn, left);
//		c = tex1Dfetch<float>(texIn, offset);
//		r = tex1Dfetch<float>(texIn, right);
//		b = tex1Dfetch<float>(texIn, bottom);
//
//	}
//	else {
//		t = tex1Dfetch<float>(texOut, top);
//		l = tex1Dfetch<float>(texOut, left);
//		c = tex1Dfetch<float>(texOut, offset);
//		r = tex1Dfetch<float>(texOut, right);
//		b = tex1Dfetch<float>(texOut, bottom);
//	}
//	dst[offset] = c + SPEED * (t + b + r + l - 4 * c);
//}
//
//
//
//void anim_gpu(AppBlock *appBlock, int ticks) {
//	//HANDLE_ERROR(cudaEventRecord(appBlock->start, 0));
//	//dim3    blocks(DIM / 16, DIM / 16);
//	//dim3    threads(16, 16);
//	//CPUAnimBitmap  *bitmap = appBlock->bitmap;
//
//	//// since tex is global and bound, we have to use a flag to
//	//// select which is in/out per iteration
//	//volatile bool dstOut = true;
//	//for (int i = 0; i<100; i++) {
//	//	float   *in, *out;
//	//	if (dstOut) {
//	//		in = appBlock->d->dev_inSrc;
//	//		out = appBlock->d->dev_outSrc;
//	//	}
//	//	else {
//	//		out = appBlock->d->dev_inSrc;
//	//		in = appBlock->d->dev_outSrc;
//	//	}
//	//	blend_kernel << <blocks, threads >> >(out, dstOut, *appBlock->d->texIn, *appBlock->d->texOut);
//	//	dstOut = !dstOut;
//	//}
//
//	//float_to_color << <blocks, threads >> >(appBlock->d->output_bitmap, appBlock->d->dev_outSrc);
//
//	//HANDLE_ERROR(cudaMemcpy(bitmap->get_ptr(),
//	//	d->output_bitmap,
//	//	bitmap->image_size(),
//	//	cudaMemcpyDeviceToHost));
//
//	//HANDLE_ERROR(cudaEventRecord(appBlock->stop, 0));
//	//HANDLE_ERROR(cudaEventSynchronize(d->stop));
//	//float   elapsedTime;
//	//HANDLE_ERROR(cudaEventElapsedTime(&elapsedTime,
//	//	d->start, d->stop));
//	//d->totalTime += elapsedTime;
//	//++d->frames;
//	//printf("Average Time per frame:  %3.1f ms\n", d->totalTime / d->frames);
//	//printf("tic:  %d\n\n", ticks);
//}
//
//// clean up memory allocated on the GPU
//
//void anim_exit(AppBlock *d) {
//
//	//cudaDestroyTextureObject(*(d->texIn));
//	//cudaDestroyTextureObject(*d->texOut);
//	//cudaDestroyTextureObject(*d->texConst);
//
//	//HANDLE_ERROR(cudaFree(d->dev_inSrc));
//	//HANDLE_ERROR(cudaFree(d->dev_outSrc));
//	//HANDLE_ERROR(cudaFree(d->dev_constSrc));
//
//	//HANDLE_ERROR(cudaEventDestroy(d->start));
//	//HANDLE_ERROR(cudaEventDestroy(d->stop));
//}
//
//
////int main(void)
////{
////	PlyBlock   data;
////	CPUAnimBitmap bitmap(DIM, DIM, &data);
////	data.bitmap = &bitmap;
////	data.totalTime = 0;
////	data.frames = 0;
////	HANDLE_ERROR(cudaEventCreate(&data.start));
////	HANDLE_ERROR(cudaEventCreate(&data.stop));
////
////	int imageSize = bitmap.image_size();
////
////	HANDLE_ERROR(cudaMalloc((void**)&data.output_bitmap,
////		imageSize));
////
////	// assume float == 4 chars in size (ie rgba)
////	HANDLE_ERROR(cudaMalloc((void**)&data.dev_inSrc, imageSize));
////	HANDLE_ERROR(cudaMalloc((void**)&data.dev_outSrc, imageSize));
////	HANDLE_ERROR(cudaMalloc((void**)&data.dev_constSrc, imageSize));
////
////
////	data.texIn = TexObjFloat1D(data.dev_inSrc, imageSize);
////	data.texOut = TexObjFloat1D(data.dev_outSrc, imageSize);
////	data.texConst = TexObjFloat1D(data.dev_constSrc, imageSize);
////
////
////
////	// intialize the constant data
////	float *temp = RndFloat0to1(DIM*DIM);
////
////	HANDLE_ERROR(cudaMemcpy(data.dev_constSrc, temp, imageSize,
////		cudaMemcpyHostToDevice));
////
////
////	HANDLE_ERROR(cudaMemcpy(data.dev_inSrc, temp, imageSize,
////		cudaMemcpyHostToDevice));
////	free(temp);
////
////	bitmap.anim_and_exit((void(*)(void*, int))anim_gpu,
////		(void(*)(void*))anim_exit);
////}

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