|
1 | 1 | import argparse |
| 2 | +import functools |
| 3 | +import random |
2 | 4 | import subprocess |
3 | 5 |
|
| 6 | +import pandas as pd |
| 7 | +import torch |
| 8 | +import torch.nn.functional |
| 9 | +import triton |
| 10 | + |
| 11 | +import ops.ninetoothed.torch |
| 12 | +import ops.triton.torch |
| 13 | +import rotary_position_embedding |
| 14 | + |
4 | 15 | PROMPTS = ( |
5 | 16 | "The emergence of deep learning domain-specific languages (DSLs) has substantially reduced the obstacles in developing high-performance, cross-platform compute kernels, but current DSLs", |
6 | 17 | "Driven by recent advancements in the AI industry, the AI accelerator sector has increasingly diversified, with vendors developing their own hardware architectures and programming models, such as NVIDIA", |
|
15 | 26 | ALL_MAX_NEW_TOKENS = (128, 512, 2048) |
16 | 27 |
|
17 | 28 |
|
| 29 | +def _run_task(op_name, dtype, device, *arg_shapes, **kwarg_shapes): |
| 30 | + ninetoothed_op = getattr(ops.ninetoothed.torch, op_name) |
| 31 | + triton_op = getattr(ops.triton.torch, op_name) |
| 32 | + |
| 33 | + if op_name == "rotary_position_embedding": |
| 34 | + torch_op = rotary_position_embedding.torch_rotary_position_embedding |
| 35 | + else: |
| 36 | + torch_op = ( |
| 37 | + getattr(torch, op_name) |
| 38 | + if hasattr(torch, op_name) |
| 39 | + else getattr(torch.nn.functional, op_name) |
| 40 | + ) |
| 41 | + |
| 42 | + if op_name == "rms_norm": |
| 43 | + torch_op = functools.partial(torch_op, normalized_shape=arg_shapes[0][-1:]) |
| 44 | + elif op_name == "softmax": |
| 45 | + torch_op = functools.partial(torch_op, dim=-1) |
| 46 | + |
| 47 | + args = tuple( |
| 48 | + torch.randn(shape, dtype=dtype, device=device) if shape else random.gauss(0, 1) |
| 49 | + for shape in arg_shapes |
| 50 | + ) |
| 51 | + kwargs = { |
| 52 | + key: torch.randn(shape, dtype=dtype, device=device) |
| 53 | + if shape |
| 54 | + else random.gauss(0, 1) |
| 55 | + for key, shape in kwarg_shapes.items() |
| 56 | + } |
| 57 | + |
| 58 | + arg_shape_string = ", ".join(str(shape) for shape in arg_shapes) |
| 59 | + kwarg_shape_string = ", ".join( |
| 60 | + f"{key}={shape}" for key, shape in kwarg_shapes.items() |
| 61 | + ) |
| 62 | + shape_string = ( |
| 63 | + f"{arg_shape_string}, {kwarg_shape_string}" |
| 64 | + if kwarg_shape_string |
| 65 | + else arg_shape_string |
| 66 | + ) |
| 67 | + |
| 68 | + task_description = f"{op_name}({shape_string})" |
| 69 | + |
| 70 | + return task_description, _benchmark_ops( |
| 71 | + (ninetoothed_op, triton_op, torch_op), *args, **kwargs |
| 72 | + ) |
| 73 | + |
| 74 | + |
| 75 | +def _benchmark_ops(ops, *args, **kwargs): |
| 76 | + assert all( |
| 77 | + torch.allclose( |
| 78 | + op(*args, **kwargs), ops[0](*args, **kwargs), rtol=0.01, atol=0.01 |
| 79 | + ) |
| 80 | + for op in ops[1:] |
| 81 | + ) |
| 82 | + |
| 83 | + return tuple(triton.testing.do_bench(lambda: op(*args, **kwargs)) for op in ops) |
| 84 | + |
| 85 | + |
18 | 86 | if __name__ == "__main__": |
19 | 87 | parser = argparse.ArgumentParser(description="Run experiments.") |
20 | 88 |
|
|
29 | 97 |
|
30 | 98 | model_name_or_path = args.model |
31 | 99 |
|
| 100 | + random.seed(0) |
| 101 | + torch.manual_seed(0) |
| 102 | + |
32 | 103 | radon_commands = ( |
33 | 104 | ( |
34 | 105 | "radon", |
|
50 | 121 | with open("code_metrics.tex", "w") as f: |
51 | 122 | subprocess.run(("python", "compare_code_metrics.py"), stdout=f, check=True) |
52 | 123 |
|
| 124 | + dtype = torch.float16 |
| 125 | + device = "cuda" |
| 126 | + |
| 127 | + tasks = ( |
| 128 | + ("add", ((4096 * 4096,), (4096 * 4096,)), {}), |
| 129 | + ( |
| 130 | + "addmm", |
| 131 | + ((4096, 4096), (4096, 4096), (4096, 4096)), |
| 132 | + {"beta": (), "alpha": ()}, |
| 133 | + ), |
| 134 | + ("bmm", ((4, 2048, 2048), (4, 2048, 2048)), {}), |
| 135 | + ("conv2d", ((4, 512, 14, 14), (512, 512, 3, 3)), {}), |
| 136 | + ("mm", ((4096, 4096), (4096, 4096)), {}), |
| 137 | + ("rms_norm", ((4096, 4096),), {}), |
| 138 | + ("rotary_position_embedding", ((4, 1024, 48, 64), (1024, 32), (1024, 32)), {}), |
| 139 | + ( |
| 140 | + "scaled_dot_product_attention", |
| 141 | + ((4, 48, 1024, 64), (4, 48, 1024, 64), (4, 48, 1024, 64)), |
| 142 | + {}, |
| 143 | + ), |
| 144 | + ("silu", ((4096 * 4096,),), {}), |
| 145 | + ("softmax", ((4096, 4096),), {}), |
| 146 | + ) |
| 147 | + |
| 148 | + data = {"Task": [], "NineToothed": [], "Triton": [], "PyTorch": []} |
| 149 | + |
| 150 | + for name, args, kwargs in tasks: |
| 151 | + description, results = _run_task(name, dtype, device, *args, **kwargs) |
| 152 | + |
| 153 | + data["Task"].append(description) |
| 154 | + |
| 155 | + for i, provider in enumerate(("NineToothed", "Triton", "PyTorch")): |
| 156 | + data[provider].append(results[i]) |
| 157 | + |
| 158 | + df = pd.DataFrame(data) |
| 159 | + |
| 160 | + df.set_index("Task").to_csv("performance-metrics.csv") |
| 161 | + |
53 | 162 | for max_new_tokens in ALL_MAX_NEW_TOKENS: |
54 | 163 | for backend in BACKENDS: |
55 | 164 | with open(f"infer_{max_new_tokens}_{backend}.json", "w") as f: |
|
0 commit comments