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

#define NUM_BRAINS 6

__global__ void moe_expert_kernel(uint32_t* vram, size_t size, int brain_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;
            // 뇌(Expert)의 성향(brain_id)에 따른 8-State 홀로그래픽 위상 맵핑
            float decoded_weight = (float)two_bits - 1.5f + (brain_id * 0.1f);
            sum += decoded_weight;
        }
        
        // DCE(Dead Code Elimination) 방지를 위해 복원된 결과값을 다시 메모리에 저장
        vram[idx] = compressed_val ^ *((uint32_t*)&sum);
    }
}

int main() {
    size_t size = 10000000; 
    uint32_t* d_brains[NUM_BRAINS];
    
    for(int b=0; b<NUM_BRAINS; b++) hipMalloc(&d_brains[b], size * 4);
    
    LARGE_INTEGER freq, start, end;
    QueryPerformanceFrequency(&freq);
    
    for(int b=0; b<NUM_BRAINS; b++) {
        hipLaunchKernelGGL(moe_expert_kernel, dim3((size+255)/256), dim3(256), 0, 0, d_brains[b], size, b);
    }
    hipDeviceSynchronize();
    
    int passes = 1000; 
    
    QueryPerformanceCounter(&start);
    for(int i=0; i<passes; i++) {
        int target_brain = i % NUM_BRAINS; 
        hipLaunchKernelGGL(moe_expert_kernel, dim3((size+255)/256), dim3(256), 0, 0, d_brains[target_brain], size, target_brain);
    }
    hipDeviceSynchronize(); 
    
    uint32_t check_val;
    hipMemcpy(&check_val, d_brains[5], 4, hipMemcpyDeviceToHost);
    QueryPerformanceCounter(&end);
    
    double elapsed = (double)(end.QuadPart - start.QuadPart) / freq.QuadPart;
    double tps = 1000.0 / elapsed;
    
    printf(">> ⏱️ True ROCm Time (w/ Decompression): %f s | 🚀 Honest MoE TPS: %f\n", elapsed, tps);
    for(int b=0; b<NUM_BRAINS; b++) hipFree(d_brains[b]);
    return 0;
}