import torch
import triton
import triton.language as tl
import time

@triton.jit
def copy_kernel(A, B, N, BLOCK_SIZE: tl.constexpr):
    pid = tl.program_id(0)
    offsets = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE)
    mask = offsets < N
    b = tl.load(B + offsets, mask=mask)
    tl.store(A + offsets, b, mask=mask)

def run_high_bw():
    N = 1024 * 1024 * 512  # 512M elements (1GB for BF16)
    dtype = torch.bfloat16
    A = torch.empty(N, device="cuda", dtype=dtype)
    B = torch.randn(N, device="cuda", dtype=dtype)
    grid = (triton.cdiv(N, 1024),)
    
    torch.cuda.synchronize()
    start = time.time()
    for _ in range(100): copy_kernel[grid](A, B, N, BLOCK_SIZE=1024)
    torch.cuda.synchronize()
    bw = (2 * N * 2) / ((time.time() - start) / 100) / 1e12
    print(f"H200 HBM3e COPY (BF16): {bw:.2f} TB/s")

if __name__ == "__main__":
    run_high_bw()