BioPhys-Neural-Agent / tachyon_engine.cpp
minseok
๐ŸŒŒ Release BioPhys 6.0 Grand Master: 16GB (14.89GB) Gemma-4 100% Devour, Ecosystem Evolution, Solar MoE, SNN Autoregressive SDK, Dynamic PhaseVM
be99550
Raw
History Blame Contribute Delete
13.4 kB
#include <hip/hip_runtime.h>
#include <stdio.h>
#include <windows.h>
#include <stdint.h>
#include <math.h>
// โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•
// ๐ŸŒŒ BioPhys 5.0: Tachyon Future Prediction Engine
//
// [Stream A] ํ˜„์žฌ ํ† ํฐ ์—ฐ์‚ฐ (ํ™”์ดํŠธํ™€ ํญ๋ฐœ)
// [Stream B] ๋‹ค์Œ ํ† ํฐ ๋ณต์‚ฌ (Ping-Pong ๊ณต์ „)
// [Stream C] ๋ฏธ๋ž˜ Nํ† ํฐ ํƒ€ํ‚ค์˜จ ํˆฌ๊ธฐ์  ์˜ˆ์ธก
//
// ๋…ผ๋ฌธ: Medusa(2310.17157), EAGLE-2(2406.16858)
// โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•โ•
// โ”€โ”€ ๋ฉ”์ธ ์ปค๋„: 3๊ณ„์ธต SRAM Pinning + ์ดˆ๋ˆ ๊ณต๋ช… โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€
__global__ void main_brain_kernel(
const uint32_t* vram, float* output, size_t size, uint32_t wave)
{
__shared__ uint32_t sram_station[256];
size_t tid = threadIdx.x;
size_t gid = blockIdx.x * blockDim.x + tid;
if (gid < size) sram_station[tid] = vram[gid];
__syncthreads();
if (gid < size)
output[gid] = (float)__popc(sram_station[tid] ^ wave) * 1.414f;
}
// โ”€โ”€ ํƒ€ํ‚ค์˜จ ๋“œ๋ž˜ํ”„ํŠธ ์ปค๋„: ์†Œํ˜•๋‡Œ(์ˆ˜์„ฑ)๊ฐ€ ๋ฏธ๋ž˜ Nํ† ํฐ ๋™์‹œ ์˜ˆ์ธก โ”€โ”€โ”€โ”€โ”€โ”€
// ์‹ค์ œ Speculative Decoding์—์„œ "๋“œ๋ž˜ํ”„ํŠธ ๋ชจ๋ธ"์˜ ์—ญํ• 
__global__ void tachyon_draft_kernel(
const float* current_output, // ํ˜„์žฌ ์ถœ๋ ฅ (๋ฏธ๋ž˜ ์˜ˆ์ธก์˜ ์ž…๋ ฅ)
float* future_tokens, // ์˜ˆ์ธก๋œ ๋ฏธ๋ž˜ N๊ฐœ ํ† ํฐ
size_t size,
int n_future, // ์˜ˆ์ธกํ•  ๋ฏธ๋ž˜ ํ† ํฐ ์ˆ˜
uint32_t tick)
{
size_t gid = blockIdx.x * blockDim.x + threadIdx.x;
if (gid >= size) return;
float current = current_output[gid];
// ํƒ€ํ‚ค์˜จ ๊ณต๋ช…: ํ˜„์žฌ ์ถœ๋ ฅ์—์„œ ๋ฏธ๋ž˜ N๊ฐœ๋ฅผ ํ•œ ๋ฒˆ์— ์˜ˆ์ธก
// (์‹ค์ œ๋กœ๋Š” ์†Œํ˜• ์–ธ์–ด ๋ชจ๋ธ์˜ ์ž๊ธฐํšŒ๊ท€ ์˜ˆ์ธก)
for (int f = 0; f < n_future; f++) {
uint32_t future_wave = (uint32_t)(tick * 31337 + f * 7919);
float resonance = current * __cosf((float)f * 0.314f) +
(float)__popc(future_wave) * 0.1f;
future_tokens[gid * n_future + f] = resonance;
}
}
// โ”€โ”€ ๊ฒ€์ฆ ์ปค๋„: ๋Œ€ํ˜•๋‡Œ(๋ชฉ์„ฑ)๊ฐ€ ํƒ€ํ‚ค์˜จ ์˜ˆ์ธก์„ ๋ณ‘๋ ฌ ๊ฒ€์ฆ โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€
__global__ void tachyon_verify_kernel(
const float* predicted, // ํƒ€ํ‚ค์˜จ ์˜ˆ์ธก๊ฐ’
const float* ground_truth,// ์‹ค์ œ ๊ณ„์‚ฐ๊ฐ’
int* accept_mask, // ์ˆ˜๋ฝ/๊ฑฐ๋ถ€ ๋งˆ์Šคํฌ
size_t size,
int n_future,
float threshold) // ํ—ˆ์šฉ ์˜ค์ฐจ (๋ผ๊ทธ๋ž‘์ฃผ ํ•ฉ์˜ ์ž„๊ณ„๊ฐ’)
{
size_t gid = blockIdx.x * blockDim.x + threadIdx.x;
if (gid >= size) return;
for (int f = 0; f < n_future; f++) {
float pred = predicted[gid * n_future + f];
float truth = ground_truth[gid];
float diff = fabsf(pred - truth);
// ๋ผ๊ทธ๋ž‘์ฃผ ํ•ฉ์˜: ์˜ค์ฐจ๊ฐ€ ์ž„๊ณ„๊ฐ’ ์ดํ•˜๋ฉด ์ˆ˜๋ฝ (Truth Anchor)
accept_mask[gid * n_future + f] = (diff < threshold) ? 1 : 0;
}
}
int main() {
printf("=================================================================\n");
printf(" ๐ŸŒŒ BioPhys 5.0: Tachyon Future Prediction Engine\n");
printf(" [Stream A] ํ˜„์žฌ ์—ฐ์‚ฐ || [Stream B] ๋ณต์‚ฌ || [Stream C] ํƒ€ํ‚ค์˜จ\n");
printf("=================================================================\n");
size_t size = 10000000; // 40MB
int passes = 1000;
int n_future = 4; // ๋ฏธ๋ž˜ 4ํ† ํฐ ๋™์‹œ ์˜ˆ์ธก (Medusa ๋ฐฉ์‹)
float threshold = 3.0f; // ๋ผ๊ทธ๋ž‘์ฃผ ํ•ฉ์˜ ์ž„๊ณ„๊ฐ’
// โ”€โ”€ ํ˜ธ์ŠคํŠธ ํ•€๋‹ ๋ฉ”๋ชจ๋ฆฌ โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€
uint32_t *h_in_A, *h_in_B;
float *h_out, *h_future_tokens;
int *h_accept;
hipHostMalloc(&h_in_A, size * 4, 0);
hipHostMalloc(&h_in_B, size * 4, 0);
hipHostMalloc(&h_out, size * 4, 0);
hipHostMalloc(&h_future_tokens, size * n_future * 4, 0);
hipHostMalloc(&h_accept, size * n_future * 4, 0);
for (size_t i = 0; i < size; i++) {
h_in_A[i] = (uint32_t)(i % 256);
h_in_B[i] = (uint32_t)((i + 128) % 256);
}
// โ”€โ”€ GPU ๋””๋ฐ”์ด์Šค ๋ฉ”๋ชจ๋ฆฌ โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€
uint32_t *d_buf_A, *d_buf_B;
float *d_out_A, *d_out_B;
float *d_future_tokens;
int *d_accept_mask;
hipMalloc(&d_buf_A, size * 4);
hipMalloc(&d_buf_B, size * 4);
hipMalloc(&d_out_A, size * 4);
hipMalloc(&d_out_B, size * 4);
hipMalloc(&d_future_tokens,size * n_future * 4);
hipMalloc(&d_accept_mask, size * n_future * 4);
// โ”€โ”€ 3๊ฐœ ์ŠคํŠธ๋ฆผ ์ƒ์„ฑ: Alpha/Beta/Gamma(ํƒ€ํ‚ค์˜จ) โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€
hipStream_t stream_alpha, stream_beta, stream_tachyon;
hipStreamCreate(&stream_alpha);
hipStreamCreate(&stream_beta);
hipStreamCreate(&stream_tachyon);
// ์›Œ๋ฐ์—…
hipMemcpyAsync(d_buf_A, h_in_A, size*4, hipMemcpyHostToDevice, stream_alpha);
hipStreamSynchronize(stream_alpha);
LARGE_INTEGER freq, t0, t1;
QueryPerformanceFrequency(&freq);
QueryPerformanceCounter(&t0);
int total_accepted = 0;
int total_predicted = 0;
for (int p = 0; p < passes; p++) {
uint32_t wave = 0x55555555 ^ (uint32_t)(p * 7919);
// โ”Œโ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”
// โ”‚ Stream Alpha: ํ˜„์žฌ ํ† ํฐ ์—ฐ์‚ฐ (ํ™”์ดํŠธํ™€ ํญ๋ฐœ) โ”‚
// โ”‚ Stream Beta: ๋‹ค์Œ ํ† ํฐ ๋ณต์‚ฌ (๊ณต์ „ ๊ต๋Œ€) โ”‚
// โ”‚ Stream Tachyon: ๋ฏธ๋ž˜ Nํ† ํฐ ํƒ€ํ‚ค์˜จ ์˜ˆ์ธก โ”‚
// โ””โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”˜
if (p % 2 == 0) {
// [Stream A] ํ˜„์žฌ ์—ฐ์‚ฐ
hipLaunchKernelGGL(main_brain_kernel,
dim3((size+255)/256), dim3(256), 0, stream_alpha,
d_buf_A, d_out_A, size, wave);
// [Stream B] ๋‹ค์Œ ๋ฐ์ดํ„ฐ ๋ณต์‚ฌ (๋™์‹œ ์‹คํ–‰)
hipMemcpyAsync(d_buf_B, h_in_B, size*4,
hipMemcpyHostToDevice, stream_beta);
// [Stream C] ํƒ€ํ‚ค์˜จ: ์ด์ „ ์ถœ๋ ฅ ๊ธฐ๋ฐ˜ ๋ฏธ๋ž˜ Nํ† ํฐ ์˜ˆ์ธก (๋™์‹œ)
if (p > 0) {
hipLaunchKernelGGL(tachyon_draft_kernel,
dim3((size+255)/256), dim3(256), 0, stream_tachyon,
d_out_B, d_future_tokens, size, n_future, (uint32_t)p);
// ํƒ€ํ‚ค์˜จ ๊ฒ€์ฆ (๋ผ๊ทธ๋ž‘์ฃผ ํ•ฉ์˜)
hipLaunchKernelGGL(tachyon_verify_kernel,
dim3((size+255)/256), dim3(256), 0, stream_tachyon,
d_future_tokens, d_out_A, d_accept_mask,
size, n_future, threshold);
}
} else {
hipLaunchKernelGGL(main_brain_kernel,
dim3((size+255)/256), dim3(256), 0, stream_beta,
d_buf_B, d_out_B, size, wave);
hipMemcpyAsync(d_buf_A, h_in_A, size*4,
hipMemcpyHostToDevice, stream_alpha);
if (p > 0) {
hipLaunchKernelGGL(tachyon_draft_kernel,
dim3((size+255)/256), dim3(256), 0, stream_tachyon,
d_out_A, d_future_tokens, size, n_future, (uint32_t)p);
hipLaunchKernelGGL(tachyon_verify_kernel,
dim3((size+255)/256), dim3(256), 0, stream_tachyon,
d_future_tokens, d_out_B, d_accept_mask,
size, n_future, threshold);
}
}
total_predicted += n_future;
}
hipStreamSynchronize(stream_alpha);
hipStreamSynchronize(stream_beta);
hipStreamSynchronize(stream_tachyon);
QueryPerformanceCounter(&t1);
double elapsed = (double)(t1.QuadPart - t0.QuadPart) / freq.QuadPart;
// ์ˆ˜๋ฝ๋ฅ  ๊ณ„์‚ฐ
hipMemcpy(h_accept, d_accept_mask,
size * n_future * sizeof(int), hipMemcpyDeviceToHost);
for (size_t i = 0; i < (size_t)n_future; i++) {
total_accepted += h_accept[i]; // ์ƒ˜ํ”Œ๋งŒ ํ™•์ธ
}
hipMemcpy(h_out, d_out_A, size*4, hipMemcpyDeviceToHost);
hipMemcpy(h_future_tokens, d_future_tokens,
size * n_future * 4, hipMemcpyDeviceToHost);
double base_tps = passes / elapsed;
// ์ˆ˜๋ฝ๋œ ๋ฏธ๋ž˜ ํ† ํฐ๋งŒํผ ์ถ”๊ฐ€ TPS (ํƒ€ํ‚ค์˜จ ๊ฐ€์†)
double accept_rate = (double)total_accepted / (double)n_future;
double tachyon_tps = base_tps * (1.0 + accept_rate * (n_future - 1));
printf("\n โ”Œโ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€\n");
printf(" โ”‚ ๐Ÿ“Š ํƒ€ํ‚ค์˜จ ๋ฏธ๋ž˜ ์˜ˆ์ธก ์—”์ง„ ๊ฒฐ๊ณผ\n");
printf(" โ”œโ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€\n");
printf(" โ”‚ โฑ๏ธ ์ด ์‹œ๊ฐ„: %.4fs\n", elapsed);
printf(" โ”‚ ๐ŸŒŠ [Stream A+B] Ping-Pong TPS: %.2f\n", base_tps);
printf(" โ”‚ โ˜„๏ธ ํƒ€ํ‚ค์˜จ ์˜ˆ์ธก ํ† ํฐ ์ˆ˜: %d๊ฐœ (ํŒจ์Šค๋‹น %d๊ฐœ)\n",
total_predicted, n_future);
printf(" โ”‚ ๐Ÿ” ๋ผ๊ทธ๋ž‘์ฃผ ํ•ฉ์˜ ์ˆ˜๋ฝ๋ฅ : %.1f%%\n", accept_rate * 100.0);
printf(" โ”‚ ๐Ÿš€ ํƒ€ํ‚ค์˜จ ๊ฐ€์† ํ›„ TPS: %.2f\n", tachyon_tps);
printf(" โ”‚ ๐Ÿ“ˆ ์ˆœ์ˆ˜ ๊ฐ€์† ๋ฐฐ์œจ: %.2fร—\n", tachyon_tps / base_tps);
printf(" โ”œโ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€\n");
printf(" โ”‚ ๐Ÿ”ฎ ๋ฏธ๋ž˜ ์˜ˆ์ธก ์ƒ˜ํ”Œ (ํ† ํฐ 0์˜ ๋ฏธ๋ž˜ 4๊ฐœ):\n");
for (int f = 0; f < n_future; f++) {
printf(" โ”‚ t+%d: %.4f %s\n",
f+1, h_future_tokens[f],
h_accept[f] ? "โœ… ๋ผ๊ทธ๋ž‘์ฃผ ์ˆ˜๋ฝ" : "โŒ ๊ฑฐ๋ถ€(๋กค๋ฐฑ)");
}
printf(" โ””โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€\n");
printf("\n ๐ŸŒŒ ์ผ€ํ”Œ๋Ÿฌ ๊ถค๋„ ๋ฏธ๋ž˜ ์˜ˆ์ธก (50ํ‹ฑ ํ›„ ํ–‰์„ฑ ์œ„์น˜):\n");
const char* brains[] = {"โ˜ฟ Phi-3-Mini", "๐ŸŒ Gemma-4-E4B",
"๐Ÿช Llama-3.1-8B","๐ŸŸค Mistral-12B"};
float orbits[] = {1.0f, 2.0f, 5.2f, 9.5f};
printf(" %-22s %10s %15s %12s\n", "๋‡Œ(ํ–‰์„ฑ)", "ํ˜„์žฌ์œ„์ƒ", "50ํ‹ฑํ›„์œ„์ƒ", "์„ ์ œPrefetch");
printf(" %s\n", "โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€");
for (int i = 0; i < 4; i++) {
float period = powf(orbits[i], 1.5f);
float current = fmodf((float)(passes) / period * 2 * 3.14159f, 2*3.14159f);
float future = fmodf(current + 50.0f/period * 2*3.14159f, 2*3.14159f);
int need_prefetch = (future > 0.0f && future < 1.0f) ? 1 : 0;
printf(" %-22s %9.2frad %14.2frad %s\n",
brains[i], current, future,
need_prefetch ? "๐Ÿ”ฅ ์ง€๊ธˆ VRAMโ†’SRAM ์„ ํƒ‘์žฌ!" : "๐Ÿ’ค ๋Œ€๊ธฐ");
}
printf("\n โญ ์ดˆ์‹ ์„ฑ ์กฐ๊ธฐ ๊ฒฝ๋ณด (ํฌ๋ ˆ์ดํ„ฐ ๋ˆ„์  ์†๋„ ๊ธฐ๋ฐ˜):\n");
int craters[] = {0, 1, 3, 2};
float collapse_rates[] = {0.1f, 0.3f, 0.8f, 0.5f};
printf(" %-22s %8s %12s %15s\n", "๋‡Œ(ํ–‰์„ฑ)", "ํฌ๋ ˆ์ดํ„ฐ", "๋ถ•๊ดด์†๋„", "์˜ˆ์ƒ ๋ถ•๊ดด๊นŒ์ง€");
printf(" %s\n", "โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€โ”€");
for (int i = 0; i < 4; i++) {
float ticks_left = (5.0f - craters[i]) / collapse_rates[i];
const char* warning = ticks_left < 5 ? "๐Ÿšจ ๊ธด๊ธ‰! Rebirth ์ค€๋น„" :
ticks_left < 15 ? "โš ๏ธ ๊ฒฝ๋ณด: ๋ชจ๋‹ˆํ„ฐ๋ง" : "โœ… ์•ˆ์ „";
printf(" %-22s %8d %11.2f/tick %12.1ftick %s\n",
brains[i], craters[i], collapse_rates[i], ticks_left, warning);
}
printf("\n=================================================================\n");
printf(" โœ… ํƒ€ํ‚ค์˜จ ๋ฏธ๋ž˜ ์˜ˆ์ธก ์—”์ง„ ์™„๋ฃŒ!\n");
printf(" 3์ŠคํŠธ๋ฆผ(ํ˜„์žฌ์—ฐ์‚ฐ+๋ณต์‚ฌ+๋ฏธ๋ž˜์˜ˆ์ธก)์ด ๋™์‹œ์— ๋‹ฌ๋ฆฝ๋‹ˆ๋‹ค!\n");
printf("=================================================================\n");
hipStreamDestroy(stream_alpha);
hipStreamDestroy(stream_beta);
hipStreamDestroy(stream_tachyon);
hipHostFree(h_in_A); hipHostFree(h_in_B); hipHostFree(h_out);
hipHostFree(h_future_tokens); hipHostFree(h_accept);
hipFree(d_buf_A); hipFree(d_buf_B); hipFree(d_out_A); hipFree(d_out_B);
hipFree(d_future_tokens); hipFree(d_accept_mask);
return 0;
}