/*#include <GL/glut.h>
#include <stdio.h>
#include <cuda.h>

#define DIM 500

__global__ void kernel( unsigned char *ptr, int ticks )
{
	// 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;
	// now calculate the value at that position
	float fx = x - DIM/2;
	float fy = y - DIM/2;
	float d = sqrtf( fx * fx + fy * fy );
	float fact=16.0f;
	if(d<fact)
		fact=d;
	unsigned char color = (unsigned char)((255.0f * cos(d/2.0f - ticks/8.0f))/(d/fact));
	ptr[offset*4 + 0] = 0;
	ptr[offset*4 + 1] = color*0.88f;
	ptr[offset*4 + 2] = color*0.92f;
	ptr[offset*4 + 3] = 255;
}

void display_cb()
{
	glClear(GL_COLOR_BUFFER_BIT);
	glColor3f(1,1,0);

	unsigned char *gpu_bitmap;
	unsigned char *cpu_bitmap=(unsigned char*)malloc(sizeof(unsigned char)*DIM*DIM*4);
	cudaMalloc( (void**)&gpu_bitmap, DIM*DIM*4 );


	dim3 blocks(DIM/16,DIM/16);
	dim3 threads(16,16);

	int ticks =  0;
	while(ticks<50)
	{
		kernel<<<blocks,threads>>>( gpu_bitmap, ticks );
		ticks++;

		cudaMemcpy( cpu_bitmap, gpu_bitmap,DIM*DIM*4 , cudaMemcpyDeviceToHost );

		//visualize cpu_bitmap
		glBegin(GL_POINTS);
		for(int x=0; x<DIM; x++)
		{
			for(int y=0; y<DIM; y++)
			{
				int offset = x + y * DIM;

				glColor3f((cpu_bitmap [offset*4 + 0]/255.0f),
						(cpu_bitmap [offset*4 + 1]/255.0f),
						(cpu_bitmap [offset*4 + 1]/255.0f));
				glVertex2f(x,y);
			}
		}
		glEnd();
		glutSwapBuffers();

	}

	cudaFree( gpu_bitmap );
	glutPostRedisplay();

}

void reshape_cb (int w, int h)
{
	if (w==0||h==0) return;
	glViewport(0,0,w,h);
	glMatrixMode (GL_PROJECTION);
	glLoadIdentity ();
	gluOrtho2D(0,w,0,h);
	glMatrixMode (GL_MODELVIEW);
	glLoadIdentity ();
}



void initialize()
{
	glutInitDisplayMode (GLUT_RGBA|GLUT_DOUBLE);
	glutInitWindowSize (DIM,DIM);
	glutInitWindowPosition (100,100);
	glutCreateWindow ("Ventana OpenGL");
	glutDisplayFunc (display_cb);
	glutReshapeFunc (reshape_cb);
	glClearColor(0.f,0.f,0.f,1.f);
}

int main (int argc, char **argv)
{
	glutInit (&argc, argv);
	initialize();
	glutMainLoop();
	return 0;
}

*/

	.file	"tmpxft_00338325_00000000-6_E3_GeneracionImagenOpengl.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 <GL/glut.h>
#include <stdio.h>
#include <cuda.h>

#define DIM 500

__global__ void kernel( unsigned char *ptr, int ticks )
{
	// 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;
	// now calculate the value at that position
	float fx = x - DIM/2;
	float fy = y - DIM/2;
	float d = sqrtf( fx * fx + fy * fy );
	float fact=16.0f;
	if(d<fact)
		fact=d;
	unsigned char color = (unsigned char)((255.0f * cos(d/2.0f - ticks/8.0f))/(d/fact));
	ptr[offset*4 + 0] = 0;
	ptr[offset*4 + 1] = color*0.88f;
	ptr[offset*4 + 2] = color*0.92f;
	ptr[offset*4 + 3] = 255;
}

void display_cb()
{
	glClear(GL_COLOR_BUFFER_BIT);
	glColor3f(1,1,0);

	unsigned char *gpu_bitmap;
	unsigned char *cpu_bitmap=(unsigned char*)malloc(sizeof(unsigned char)*DIM*DIM*4);
	cudaMalloc( (void**)&gpu_bitmap, DIM*DIM*4 );


	dim3 blocks(DIM/16,DIM/16);
	dim3 threads(16,16);

	int ticks =  0;
	while(ticks<50)
	{
		kernel<<<blocks,threads>>>( gpu_bitmap, ticks );
		ticks++;

		cudaMemcpy( cpu_bitmap, gpu_bitmap,DIM*DIM*4 , cudaMemcpyDeviceToHost );

		//visualize cpu_bitmap
		glBegin(GL_POINTS);
		for(int x=0; x<DIM; x++)
		{
			for(int y=0; y<DIM; y++)
			{
				int offset = x + y * DIM;

				glColor3f((cpu_bitmap [offset*4 + 0]/255.0f),
						(cpu_bitmap [offset*4 + 1]/255.0f),
						(cpu_bitmap [offset*4 + 1]/255.0f));
				glVertex2f(x,y);
			}
		}
		glEnd();
		glutSwapBuffers();

	}

	cudaFree( gpu_bitmap );
	glutPostRedisplay();

}

void reshape_cb (int w, int h)
{
	if (w==0||h==0) return;
	glViewport(0,0,w,h);
	glMatrixMode (GL_PROJECTION);
	glLoadIdentity ();
	gluOrtho2D(0,w,0,h);
	glMatrixMode (GL_MODELVIEW);
	glLoadIdentity ();
}



void initialize()
{
	glutInitDisplayMode (GLUT_RGBA|GLUT_DOUBLE);
	glutInitWindowSize (DIM,DIM);
	glutInitWindowPosition (100,100);
	glutCreateWindow ("Ventana OpenGL");
	glutDisplayFunc (display_cb);
	glutReshapeFunc (reshape_cb);
	glClearColor(0.f,0.f,0.f,1.f);
}

int main (int argc, char **argv)
{
	glutInit (&argc, argv);
	initialize();
	glutMainLoop();
	return 0;
}

*/

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