//
// Created by luchin on 29-07-21.
//
#include <cassert>
#include <iostream>

//static cudaError_t checkCuda(cudaError_t result) {
//    if (result != cudaSuccess) {
//        fprintf(stderr, "CUDA Runtime Error: %s\n", cudaGetErrorString(result));
//        assert(result == cudaSuccess);
//    }
//    return result;
//}
//
//Par_CUDA::Par_CUDA() : AbstractGoL() {
//    int devId = 0;
//    cudaDeviceProp prop;
//    checkCuda(cudaGetDeviceProperties(&prop, devId));
//    checkCuda(cudaSetDevice(devId));
//    int total_size = sizeof(char) * LARGO * LARGO;
//    d_grid = nullptr;
//    checkCuda(cudaMalloc(&d_grid, total_size));
//}
//
//__global__ void step(char *grid) {
//#ifdef CUDA_USE_2D
//    int x = blockIdx.x * blockDim.x + threadIdx.x;
//    int y = blockIdx.y * blockDim.y + threadIdx.y;
//#else
//    int tmp = blockIdx.x * blockDim.x + threadIdx.x;
//    int x = tmp / LARGO;
//    int y = tmp % LARGO;
//#endif
//
//    // contamos los vecinos
//    //    printf("x is %d and y is %d\n", x, y);
//    if (x > LARGO || y > LARGO) return;
//    x += LARGO; // nos aseguramos de que x-1 sea positivo
//    y += LARGO;
//    int x_m = (x - 1) % LARGO;
//    int x_p = (x + 1) % LARGO;
//    int y_m = (y - 1) % LARGO;
//    int y_p = (y + 1) % LARGO;
//    x = x % LARGO;
//    y = y % LARGO;
//    int num_neighbors =
//            grid[x_m * LARGO + y_m] + grid[x * LARGO + y_m] + grid[x_p * LARGO + y_m] +
//            grid[x_m * LARGO + y] + grid[x_p * LARGO + y] +
//            grid[x_m * LARGO + y_p] + grid[x * LARGO + y_p] + grid[x_p * LARGO + y_p];
//    char alive = grid[x * LARGO + y];
//
//    __syncthreads();
//    // reemplazamos los lugares donde corresponde
//    if ((alive && num_neighbors == 2) || num_neighbors == 3) {
//        grid[x * LARGO + y] = 1;
//    } else {
//        grid[x * LARGO + y] = 0;
//    }
//}
//
//
//void Par_CUDA::run_game(int num_steps) {
//#ifdef CUDA_USE_2D
//    dim3 dimGrid((LARGO + 7) / 8, (LARGO + 7) / 88, 1);
//    dim3 dimBlock(8, 8, 1);
//#else
//    dim3 dimGrid((LARGO * LARGO + 7) / 8, 1, 1);
//    dim3 dimBlock(8, 1, 1);
//#endif
//    cudaMemcpy(d_grid, h_grid, sizeof(char) * LARGO * LARGO, cudaMemcpyHostToDevice);
//    for (int i = 0; i < num_steps; i++) {
//        step<<<dimGrid, dimBlock>>>(d_grid);
//    }
//    cudaMemcpy(h_grid, d_grid, sizeof(char) * LARGO * LARGO, cudaMemcpyDeviceToHost);
//}
	.file	"tmpxft_0020a910_00000000-6_swe_cuda.cudafe1.cpp"
	.text
#APP
	.globl _ZSt21ios_base_library_initv

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

//
// Created by luchin on 29-07-21.
//

#include <hip/hip_runtime.h>
#include <cassert>
#include <iostream>

//static cudaError_t checkCuda(cudaError_t result) {
//    if (result != cudaSuccess) {
//        fprintf(stderr, "CUDA Runtime Error: %s\n", cudaGetErrorString(result));
//        assert(result == cudaSuccess);
//    }
//    return result;
//}
//
//Par_CUDA::Par_CUDA() : AbstractGoL() {
//    int devId = 0;
//    cudaDeviceProp prop;
//    checkCuda(cudaGetDeviceProperties(&prop, devId));
//    checkCuda(cudaSetDevice(devId));
//    int total_size = sizeof(char) * LARGO * LARGO;
//    d_grid = nullptr;
//    checkCuda(cudaMalloc(&d_grid, total_size));
//}
//
//__global__ void step(char *grid) {
//#ifdef CUDA_USE_2D
//    int x = blockIdx.x * blockDim.x + threadIdx.x;
//    int y = blockIdx.y * blockDim.y + threadIdx.y;
//#else
//    int tmp = blockIdx.x * blockDim.x + threadIdx.x;
//    int x = tmp / LARGO;
//    int y = tmp % LARGO;
//#endif
//
//    // contamos los vecinos
//    //    printf("x is %d and y is %d\n", x, y);
//    if (x > LARGO || y > LARGO) return;
//    x += LARGO; // nos aseguramos de que x-1 sea positivo
//    y += LARGO;
//    int x_m = (x - 1) % LARGO;
//    int x_p = (x + 1) % LARGO;
//    int y_m = (y - 1) % LARGO;
//    int y_p = (y + 1) % LARGO;
//    x = x % LARGO;
//    y = y % LARGO;
//    int num_neighbors =
//            grid[x_m * LARGO + y_m] + grid[x * LARGO + y_m] + grid[x_p * LARGO + y_m] +
//            grid[x_m * LARGO + y] + grid[x_p * LARGO + y] +
//            grid[x_m * LARGO + y_p] + grid[x * LARGO + y_p] + grid[x_p * LARGO + y_p];
//    char alive = grid[x * LARGO + y];
//
//    __syncthreads();
//    // reemplazamos los lugares donde corresponde
//    if ((alive && num_neighbors == 2) || num_neighbors == 3) {
//        grid[x * LARGO + y] = 1;
//    } else {
//        grid[x * LARGO + y] = 0;
//    }
//}
//
//
//void Par_CUDA::run_game(int num_steps) {
//#ifdef CUDA_USE_2D
//    dim3 dimGrid((LARGO + 7) / 8, (LARGO + 7) / 88, 1);
//    dim3 dimBlock(8, 8, 1);
//#else
//    dim3 dimGrid((LARGO * LARGO + 7) / 8, 1, 1);
//    dim3 dimBlock(8, 1, 1);
//#endif
//    cudaMemcpy(d_grid, h_grid, sizeof(char) * LARGO * LARGO, cudaMemcpyHostToDevice);
//    for (int i = 0; i < num_steps; i++) {
//        step<<<dimGrid, dimBlock>>>(d_grid);
//    }
//    cudaMemcpy(h_grid, d_grid, sizeof(char) * LARGO * LARGO, cudaMemcpyDeviceToHost);
//}
	.text
	.file	"swe_cuda.hip"
                                        # Start of file scope inline assembly
	.globl	_ZSt21ios_base_library_initv

                                        # End of file scope inline assembly
	.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
