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

__global__ void expert_kernel(uint32_t* vram, size_t size, int model_id) {
    size_t idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < size) {
        uint32_t compressed_val = vram[idx];
        float sum = 0.0f;
        
        // 정직한 실시간 16x 압축 해제 (2-bit Unpacking) 및 8-State 복원
        #pragma unroll
        for(int j=0; j<16; j++) {
            uint32_t two_bits = (compressed_val >> (j * 2)) & 0x3;
            // 아키텍처(model_id)별 상이한 가중치 간섭(Interference) 수학 모사
            float decoded_weight = (float)two_bits - 1.5f + (model_id * 0.1f);
            sum += decoded_weight;
        }
        vram[idx] = compressed_val ^ *((uint32_t*)&sum);
    }
}

int main(int argc, char** argv) {
    if (argc < 3) return -1;
    int model_id = atoi(argv[1]);
    char* model_name = argv[2];
    
    size_t size = 10000000; // 40MB Holographic Slice
    uint32_t* d_vram;
    hipMalloc(&d_vram, size * 4);
    
    LARGE_INTEGER freq, start, end;
    QueryPerformanceFrequency(&freq);
    
    // Warmup
    hipLaunchKernelGGL(expert_kernel, dim3((size+255)/256), dim3(256), 0, 0, d_vram, size, model_id);
    hipDeviceSynchronize();
    
    int passes = 1000;
    QueryPerformanceCounter(&start);
    
    // 모델별 단독 실행 루프 (Zero-Overhead)
    for(int i=0; i<passes; i++) {
        hipLaunchKernelGGL(expert_kernel, dim3((size+255)/256), dim3(256), 0, 0, d_vram, size, model_id);
    }
    hipDeviceSynchronize();
    
    uint32_t check_val;
    hipMemcpy(&check_val, &d_vram[size-1], 4, hipMemcpyDeviceToHost);
    QueryPerformanceCounter(&end);
    
    double elapsed = (double)(end.QuadPart - start.QuadPart) / freq.QuadPart;
    double tps = 1000.0 / elapsed;
    
    printf("  ├─ 🧠 [%d] %-16s | ⏱️ Time: %.4f s | 🚀 TPS: %-8.2f | 🔬 Sync: 0x%X\n", model_id, model_name, elapsed, tps, check_val);
    
    hipFree(d_vram);
    return 0;
}