minseok
π Release BioPhys 6.0 Grand Master: 16GB (14.89GB) Gemma-4 100% Devour, Ecosystem Evolution, Solar MoE, SNN Autoregressive SDK, Dynamic PhaseVM
be99550 | __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 볡μ | |
| 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; | |
| } | |