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

/* The purpose of these microkernels is to 
offer the user a sanity check. These microkernels 
take the exact same parameters as their "real" 
implementations and perform simple modifications 
so the user can be sure the kernel is unpacking 
and modifying the parameters the correct way. */ 

__device__ void ComputeTest(void* params){
  
  //Params| &table | offset | 
  //Bytes |   8    |   4    | 
  
 void *table = *((void**)params);
 int offset = *((int*)(((void**)params)+1));
  
  //Extract all the values. 
  int np = *((int*) table);
  int nd = *(((int*) table)+1);

  int size = np * nd;

  double *mass = (double*)(((int*)table)+2);
  double *pos = mass + 1;
  double *vel = pos + size; 
  double *acc = vel + size;
  double *f = acc + size;

  double *pe = f + size;
  double *ke = pe + size;

  int j;
  int tid = threadIdx.x % 32; 
  int k = offset + tid; 
  //Compute all the potential energy and forces.     
   for(j=0; j<np; j++){
      if(k == j){ continue; }

      int index = j + k *nd;
      f[index] += 1;
      pe[index] += 1;
      ke[index] += 1;
    }
}

	.file	"tmpxft_0035ae57_00000000-6_MDFake.cudafe1.cpp"
	.text
#APP

#NO_APP
	.type	_ZL26__cudaUnregisterBinaryUtilv, @function
_ZL26__cudaUnregisterBinaryUtilv:
.LFB2030:
	.cfi_startproc
	endbr64
	movq	_ZL20__cudaFatCubinHandle(%rip), %rdi
	jmp	__cudaUnregisterFatBinary@PLT
	.cfi_endproc
.LFE2030:
	.size	_ZL26__cudaUnregisterBinaryUtilv, .-_ZL26__cudaUnregisterBinaryUtilv
	.globl	_Z11ComputeTestPv
	.type	_Z11ComputeTestPv, @function
_Z11ComputeTestPv:
.LFB2027:
	.cfi_startproc
	endbr64
	pushq	%rax
	.cfi_def_cfa_offset 16
	popq	%rax
	.cfi_def_cfa_offset 8
	subq	$24, %rsp
	.cfi_def_cfa_offset 32
	movl	$1, 12(%rsp)
	movl	12(%rsp), %edi
	call	exit@PLT
	.cfi_endproc
.LFE2027:
	.size	_Z11ComputeTestPv, .-_Z11ComputeTestPv
	.section	.text.startup,"ax",@progbits
	.type	_ZL24__sti____cudaRegisterAllv, @function
_ZL24__sti____cudaRegisterAllv:
.LFB2053:
	.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
.LFE2053:
	.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>

/* The purpose of these microkernels is to 
offer the user a sanity check. These microkernels 
take the exact same parameters as their "real" 
implementations and perform simple modifications 
so the user can be sure the kernel is unpacking 
and modifying the parameters the correct way. */ 

__device__ void ComputeTest(void* params){
  
  //Params| &table | offset | 
  //Bytes |   8    |   4    | 
  
 void *table = *((void**)params);
 int offset = *((int*)(((void**)params)+1));
  
  //Extract all the values. 
  int np = *((int*) table);
  int nd = *(((int*) table)+1);

  int size = np * nd;

  double *mass = (double*)(((int*)table)+2);
  double *pos = mass + 1;
  double *vel = pos + size; 
  double *acc = vel + size;
  double *f = acc + size;

  double *pe = f + size;
  double *ke = pe + size;

  int j;
  int tid = threadIdx.x % 32; 
  int k = offset + tid; 
  //Compute all the potential energy and forces.     
   for(j=0; j<np; j++){
      if(k == j){ continue; }

      int index = j + k *nd;
      f[index] += 1;
      pe[index] += 1;
      ke[index] += 1;
    }
}

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