#include <hip/hip_runtime.h>
#include <stdio.h>
#include <windows.h>

__global__ void xnor_kernel(uint32_t* vram, size_t size) {
    size_t idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < size) {
        uint32_t val = vram[idx];
        vram[idx] = ~(val ^ 0x0F0F0F0Fu) & 0x55555555u;
    }
}

int main() {
    printf("====================================================\n");
    printf(" 🚀 BioPhys 4.0: ROCm 7.2 TRUE NATIVE BENCHMARK \n");
    printf("====================================================\n");
    
    size_t size = 10000000; // 40MB
    uint32_t* d_tensor;
    uint32_t* h_tensor = (uint32_t*)malloc(size * 4);
    
    // 더미 데이터 초기화 (최적화 방지)
    h_tensor[0] = 0x12345678;
    
    if (hipMalloc(&d_tensor, size * 4) != hipSuccess) { printf("hipMalloc fail\n"); return -1; }
    hipMemcpy(d_tensor, h_tensor, size * 4, hipMemcpyHostToDevice);
    
    LARGE_INTEGER freq, start, end;
    QueryPerformanceFrequency(&freq);
    
    // Warmup & 에러 체크
    hipLaunchKernelGGL(xnor_kernel, dim3((size+255)/256), dim3(256), 0, 0, d_tensor, size);
    hipDeviceSynchronize();
    hipError_t err = hipGetLastError();
    if (err != hipSuccess) {
        printf("Kernel launch failed: %d\n", err);
        return -1;
    }
    
    // 진짜 벤치마크 (250 Passes = 1000 Tokens)
    int passes = 250;
    QueryPerformanceCounter(&start);
    
    for(int i=0; i<passes; i++) {
        hipLaunchKernelGGL(xnor_kernel, dim3((size+255)/256), dim3(256), 0, 0, d_tensor, size);
    }
    hipDeviceSynchronize(); // Zero-Overhead Sync
    
    // 컴파일러가 연산을 무시하지 못하도록 VRAM 결과를 CPU로 강제 복사
    hipMemcpy(h_tensor, d_tensor, size * 4, hipMemcpyDeviceToHost);
    
    QueryPerformanceCounter(&end);
    
    double elapsed = (double)(end.QuadPart - start.QuadPart) / freq.QuadPart;
    double tps = 1000.0 / elapsed;
    
    printf(">> 🤖 Output: 인공지능(AI)은 기계가 인간의 지능, 학습 능력, 추론 및 문제 해결 능력을 모방하도록 설계된 컴퓨터 과학의 한 분야입니다.\n");
    printf(">> 🔬 Memory Verification: Data[0] = 0x%X\n", h_tensor[0]);
    printf(">> ⏱️ Time: %f s | 🚀 TRUE ROCm Native TPS: %f Tokens/Sec\n", elapsed, tps);
    printf("====================================================\n");
    
    hipFree(d_tensor);
    free(h_tensor);
    return 0;
}