Ո՞րն է ձեր PyTorch Code- ի արագությունը:

Ոչ ոք, ով աշխատում է PyTorch model code- ի հետ, սկսում է հարցնել նույն հարցերը: Ինչպե՞ս կարող եմ արագ լինել, թե ինչպե՞ս կարող եմ արագ լինել, թե ինչպե՞ս կարող եմ արագ լինել Ինչպե՞ս կարող եմ արագ լինել, թե ինչպե՞ս կարող եմ արագ լինել, թե ինչպե՞ս կարող եմ արագ լինել Եթե դուք ML ինժեներ, հետազոտող, կամ պարզապես որոշել եք խաղալ անմիջական ML ռեժեսորի հետ շաբաթվա ընթացքում, դուք վերջապես փորձեք հասկանալ, թե ինչպես պետք է արագացնել ձեր կոդը: Երբ մենք կարող ենք անել այն, մենք պետք է իմանալ, թե ինչպես պետք է ճշգրտությամբ չափել կատարումը. Այնուհետեւ վերցնել ճշգրտական ճշգրտությունները այդ չափանիշների. Այս հոդվածը այն մասին է, որ ճշգրտությամբ benchmarking CUDA կամ PyTorch կոդը. Հիմնական Flow Օգտագործենք H100- ում մատեգորիաների մատեգորիաները, ինչպես նաեւ ամբողջ ժամանակը: «Kernel»- ի համար մենք նշանակում ենք մի գործառույթը (լուսանկարներ), որը աշխատում է GPU- ում: Ձեր կոդը կարող է ունի շատ այդ գործառույթներ, օրինակ: torch.zeros((16, 16), սարքավորումներ=”cuda) a + b, որտեղ a եւ b են tensors GPU մանրամասում A եւ B flash_attention (կ, k, v) եւ այնպես Քանի որ GPU- ում կատարվում է kernel- ը, այն Եթե մենք կարծում ենք, որ Python ֆայլը դասընթացային կոդը է, ապա այն փոխանցում է «պրակտային» աշխատանքը, ինչպիսիք են for-loops կամ if-statements, եւ իրական kernel- ի սկսումները: Այսպիսով, CPU-ը առաջ է, եւ երբ այն կտրում է մի kernel- ը, այն կտրում է այն kernel- ը GPU- ում կատարելու համար, ավելացնել այն մի գծի մեջ: Երբ հաջորդ kernel- ը ավելացվում է գծի մեջ, ոչ մի ապահովություն չէ, որ նախորդը արդեն ավարտվել է: Բեռնել Այսպիսով GPU-ը թույլ է տալիս փրկել տեւողությունը: Այն միշտ կարող է գտնել աշխատանքը այդ կերելային գոտի մեջ (հայած այն ժամանակ, երբ synchronization- ը բացառապես կոչվում է, կամ երբ կատարումը CPU- ի հետ կապված է): Արդյոք, թե ինչու սխալ է եւ թե ինչու սխալ է Երբ մենք առաջին անգամ կարծում ենք, որ Python- ում ժամանակը չափում է, բնական ինտեգրը այն է, որ հասկանում ենք ժամանակի մոդուլը եւ գործել ենք ինչ-որ բան այսպես: import time import torch def matmul(a, b): return a @ b def get_data(m, n, k): a = torch.randn(m, n, device="cuda", dtype=torch.bfloat16) b = torch.randn(n, k, device="cuda", dtype=torch.bfloat16) return a, b def benchmark_naive(m, n, k, n_iters: int = 100): a, b = get_data(m, n, k) times = [] for _ in range(n_iters): start = time.perf_counter_ns() matmul(a, b) end = time.perf_counter_ns() times.append(end - start) # perf_counter_ns returns nanoseconds, convert to ms return sum(times) / n_iters / 1e6 small_shapes_time = benchmark_naive(16, 32, 16) large_shapes_time = benchmark_naive(4096, 8192, 4096) print(f"(16, 32) by (32, 16) matmul time: {small_shapes_time:.5f} ms") print(f"(4096, 8192) by (8192, 4096) matmul time: {large_shapes_time:.5f} ms") # (16, 32) by (32, 16) matmul time: 0.01415 ms # (4096, 8192) by (8192, 4096) matmul time: 0.01350 ms Եթե մենք սեղմում ենք այս կոդը երկու տարբեր սարքերի վրա: եւ «Մենք կարող ենք տեսնել, որ matmul ժամանակում հիմնականում ոչ մի տարբերություն չի ունենում. Իհարկե, իմ գործում, դա իրականում ավելի արագ է միացվել մետրերը, որոնք 256 անգամ մեծ են. Այս արդյունքը շատ վատ է, եւ պետք է կարծում ենք, թե ինչ է այս կոդը իրականում չափում: 16x32 4096x8192 «Մենք երբեք չենք սպասում, որ GPU matmul- ը ավարտի, քանի որ բոլոր «պարգմանական» կոդը կատարվում են CPU-ի վրա: Ի՞նչն է այն ժամանակը, ինչը պետք է մեզ համար, որպեսզի matmul kernel- ը կտրել GPU- ի գծի մեջ եւ անցնել մեր կյանքի հետ: time.perf_counter_ns() Մենք կարող ենք ավելացնել Երբ մետաղալարը կտա CPU-ը GPU-ի համար, բայց դա դեռ չի հարմարավետ, քանի որ մենք մետաղալարում ենք մետաղալարների ժամանակը CPU-ի վրա, որը ներառում է առեւտրային մետաղալար: Մենք այսպես չենք բացահայտել իրական GPU-ի կատարումը: torch.cuda.synchronize() Հիմնական հոդված՝ CUDA events GPU- ի ժամանակը, այլեւ CPU- ի ժամանակը, ճիշտ է CUDA- ի գործառույթների հետ: Նրանք markers, որոնք մենք կարող ենք տեղադրել ուղղակիորեն GPU- ի գործառույթը: GPU- ը գրանցում է ժամանակագրերը, երբ այն հասկանում է յուրաքանչյուր CUDA- ի գործառույթը, այնպես որ մենք ստանում ենք գործառույթը գործառույթի ժամանակը սարքի վրա: import torch def matmul(a, b): return a @ b def get_data(m, n, k): a = torch.randn(m, n, device="cuda", dtype=torch.bfloat16) b = torch.randn(n, k, device="cuda", dtype=torch.bfloat16) return a, b def benchmark_cuda_events(m, n, k, n_iters: int = 100): a, b = get_data(m, n, k) start_events = [torch.cuda.Event(enable_timing=True) for _ in range(n_iters)] end_events = [torch.cuda.Event(enable_timing=True) for _ in range(n_iters)] for it in range(n_iters): start_events[it].record() matmul(a, b) end_events[it].record() torch.cuda.synchronize() times = [start.elapsed_time(end) for start, end in zip(start_events, end_events)] return sum(times) / n_iters small_shapes_time = benchmark_cuda_events(16, 32, 16) large_shapes_time = benchmark_cuda_events(4096, 8192, 4096) print(f"(16, 32) by (32, 16) matmul time: {small_shapes_time:.5f} ms") print(f"(4096, 8192) by (8192, 4096) matmul time: {large_shapes_time:.5f} ms") # (16, 32) by (32, 16) matmul time: 0.02093 ms # (4096, 8192) by (8192, 4096) matmul time: 0.34025 ms Հիմնական հոդվածը, ինչպիսիք են մեր մետաղական սերտիֆիկները Տեղադրում է, որ ժամանակակից սերտիֆիկները ներառված են: The Տեսանյութը հասկանում է, թե ինչ տարբերություն է սկսելու եւ ավարտելու ժամանակը: enable_timing=True .elapsed_time Այժմ մենք ստանում ենք իրական մանրամասներ, ոչ անսահմանափակապես արագ արդյունքները անսահմանափակ approach. From this benchmark, the matmul with large shapes takes at least 15x longer. Ավելի լավ է ներառում warmup iterations- ը առաջին անգամ, քանի որ կան հաճախ մեկ անգամ ծախսերը, որոնք ներառում են առաջին սկսում, ինչպիսիք են JIT- ը: Մենք ներառում ենք warmup- ը հաջորդ չափանիշներում: L2 սարքավորումներ Երբ մենք աշխատում ենք նույն սերունդը մի քանի անգամ նույն տվյալների հետ ( եւ դա այն է, ինչ մենք անում ենք), այդ տվյալները մնում են GPU- ի L2 սերունդում: NVIDIA GPUs- ում, կա մանրամասների հիերֆիկիա: HBM- ը մեծագույնն է եւ լայնագույնը, հետո անցնում են ավելի փոքր եւ արագ մանրամասներ (L2 մանրամասը → L1 մանրամասը կամ Shared Memory → մանրամասի մանրամասը): L2 մանրամասը համեմատաբար մեծ մանրամաս է, որը նվազեցնում է հիմնական մանրամասների մանրամասները: GPU Generation L2 per GPU HBM per GPU V100 Volta 6MB 32GB A100 Ampere 40MB 80GB H100 Hopper 50MB 80GB B200 Blackwell 126MB 192GB V100-ի մասին Հաջորդ 6 մգ 32GB բովանդակություն A100-ի մասին Ampera 40 մգ 80GB բովանդակություն H100-ի մասին Hopperի մասին 50 մգ 80GB բովանդակություն B200-ի համար Blackwell-ի մասին 126 մգ 192GB բովանդակություն Այս սխալով, մեր չափանիշները կարող են ցույց տալ ավելի լավ արդյունքներ, քան այն, ինչ մենք տեսնում ենք արտադրության, որտեղ տվյալները փոխանցում են միասին սխալների միջեւ: Եթե պետք է ստանալ ավելի ճշգրիտ չափանիշներ, մենք պետք է մաքուրել սխալը: Այն մի ճանապարհ է անել, որ մի մեծ սխալը տեղադրել է եւ նորարարել է այն յուրաքանչյուր սխալի սկսում: Սա, իհարկե, է հիորիկ, այլեւ ապահովված սխալը: import torch def flush_l2_cache(): # On H100, the L2 cache is 50MB, so we allocate something a bit bigger cache_size = 60 * 1024 * 1024 buffer = torch.zeros(cache_size // 4, dtype=torch.float32, device="cuda") buffer += 1 del buffer def matmul(a, b): return a @ b def get_data(m, n, k): a = torch.randn(m, n, device="cuda", dtype=torch.bfloat16) b = torch.randn(n, k, device="cuda", dtype=torch.bfloat16) return a, b def benchmark_with_cache_flush(m, n, k, n_iters: int = 100): a, b = get_data(m, n, k) for _ in range(10): matmul(a, b) # warmup torch.cuda.synchronize() start_events = [torch.cuda.Event(enable_timing=True) for _ in range(n_iters)] end_events = [torch.cuda.Event(enable_timing=True) for _ in range(n_iters)] for it in range(n_iters): flush_l2_cache() # flush the L2 between iterations start_events[it].record() matmul(a, b) end_events[it].record() torch.cuda.synchronize() times = [start.elapsed_time(end) for start, end in zip(start_events, end_events)] return sum(times) / n_iters small_shapes_time = benchmark_with_cache_flush(16, 32, 16) large_shapes_time = benchmark_with_cache_flush(4096, 8192, 4096) print(f"(16, 32) by (32, 16) matmul time: {small_shapes_time:.5f} ms") print(f"(4096, 8192) by (8192, 4096) matmul time: {large_shapes_time:.5f} ms") # (16, 32) by (32, 16) matmul time: 0.00634 ms # (4096, 8192) by (8192, 4096) matmul time: 0.40575 ms Իհարկե, այս արագության եւ նախորդի միջեւ արդյունքում կարեւոր տարբերություն էր, բայց իրականում մի քանի անգամ վերականգնումը կոդը մի քանի անգամ փոխում է արդյունքները, եւ մեր տարբերությունը լավ է տպագրված տարբերության մեջ - այնպես որ այս հատուկ խնդիրի համար սխալը չի ունենում կարեւոր ազդեցություն: Բջջային լուծումներ Արդյոք մենք պետք է իմանալ այս ամենը միայն այն համար, որ ինչ-որ բան ստուգենք: Հիմնական հարցը չէ: Իհարկե, կա մի գրասենյակ, որը կոչվում է , որը առաջին անգամ մշակվել է OpenAI- ի կողմից մինչեւ 2021-ին GPU- ի սերտիֆիկի գրելու համար Pythonic- ի ձեւով: Այժմ այն ունի տեղադրված փորձարկման մոդել, որը ունի ճիշտ benchmarking գործիքներ, որոնք կարող են կատարել բոլոր վերեւումներին: Triton # Source: https://github.com/triton-lang/triton/blob/main/python/triton/testing.py#L127C1-L190C64 def do_bench(fn, warmup=25, rep=100, grad_to_none=None, quantiles=None, return_mode="mean"): """ Benchmark the runtime of the provided function. By default, return the median runtime of :code:`fn` along with the 20-th and 80-th performance percentile. :param fn: Function to benchmark :type fn: Callable :param warmup: Warmup time (in ms) :type warmup: int :param rep: Repetition time (in ms) :type rep: int :param grad_to_none: Reset the gradient of the provided tensor to None :type grad_to_none: torch.tensor, optional :param quantiles: Performance percentile to return in addition to the median. :type quantiles: list[float], optional :param return_mode: The statistical measure to return. Options are "min", "max", "mean", "median", or "all". Default is "mean". :type return_mode: str """ assert return_mode in ["min", "max", "mean", "median", "all"] di = runtime.driver.active.get_device_interface() fn() di.synchronize() cache = runtime.driver.active.get_empty_cache_for_benchmark() # Estimate the runtime of the function start_event = di.Event(enable_timing=True) end_event = di.Event(enable_timing=True) start_event.record() for _ in range(5): runtime.driver.active.clear_cache(cache) fn() end_event.record() di.synchronize() estimate_ms = start_event.elapsed_time(end_event) / 5 # compute number of warmup and repeat n_warmup = max(1, int(warmup / estimate_ms)) n_repeat = max(1, int(rep / estimate_ms)) start_event = [di.Event(enable_timing=True) for i in range(n_repeat)] end_event = [di.Event(enable_timing=True) for i in range(n_repeat)] # Warm-up for _ in range(n_warmup): fn() # Benchmark for i in range(n_repeat): # we don't want `fn` to accumulate gradient values # if it contains a backward pass. So we clear the # provided gradients if grad_to_none is not None: for x in grad_to_none: x.grad = None # we clear the L2 cache before each run runtime.driver.active.clear_cache(cache) # record time of `fn` start_event[i].record() fn() end_event[i].record() # Record clocks di.synchronize() times = [s.elapsed_time(e) for s, e in zip(start_event, end_event)] return _summarize_statistics(times, quantiles, return_mode) Բացի այն, ինչ մենք կպահել ենք, Արդյոք, դուք կարող եք ստանալ մի քանի մանրամասներ, ոչ միայն միջին. Օրինակ, դուք կարող եք պահանջել մանրամասը կամ 99-րդ սենյակում: Այն ունի այլ օգտակար մանրամասներ, որոնք ներառում են docstring. do_bench Այս գործառույթը օգտագործելու համար մենք կարող ենք պարզապես զանգահարել: import triton bench_time = triton.testing.do_bench(lambda: matmul(a, b)) CPU-bound execution եւ CUDA գրաֆիկներ Երկուների համար, որոնք տեւում են շատ փոքր ժամանակ, լիցքավորումը կարող է լինել այնքան մեծ, որ այն ավելի մեծ է, քան կերելային ժամանակը: Այնպես էլ, երբ կերելային լիցքավորումների միջեւ կարեւոր CPU աշխատանքը (հարկե անպատմատացված) է: Սա կախված է մեր չափանիշներին: Ավելացրենք looong for-loop- ը մեր matmul kernel- ում, քանի որ մենք գիտենք, որ for-loops- ը իրականում արագ է CPU- ում Python- ում: def cpu_heavy_matmul(a, b): cnt = 0 for _ in range(100_000): cnt += 1 return a @ b Մենք տեսնում ենք, որ (4096, 8192) by (8192, 4096) matmul ժամանակը ավելացվել է 0.4ms- ից 2.2ms- ում - ավելի քան 5x! Այժմ, այնպես որ, թե ինչ է ձեր իրական դասընթացը տեսնում, դուք կարող եք ցանկանում է փոխել CPU- ի վերահսկողությունը համեմատության արդյունքներիից կամ չեք: Օրինակ, եթե ձեր դասընթացը արդեն CPU- ի վերահսկողությունը է, ապա նույն խնդիրները կունենա արտադրության ընթացքում: Այնուհետեւ, ամենահսկողությամբ, այն չի նշանակում, որ CPU- ի վերահսկողության ժամանակը մեր դասընթացների համար կախվածվի: Սակայն, սովորական դասընթացների ընթացքում, CPU-ը չի պետք է վերահսկողել օպտիմալ դասընթացային կատարումը - սովորաբար CPU-ը աշխատում է GPU- ի վերահսկողության ընթացքում: Այս դեպքում, նույնիսկ եթե այդ մեկ կունել ունի կարեւոր վերահսկողություն, դա չի փոխել կունելի վերահսկողության Այսպիսով, մենք կարող ենք օգտագործել CUDA գրաֆներ: CUDA գրաֆները թույլ են տալիս մեզ գրել մի շարք գրաֆներ մեկ անգամ եւ վերցնել նրանց մի փոքր CPU- ի ներգրավմանով: Triton- ում, կա , որը ապահովում է ճիշտ նույն գործառույթը: Այս օգտագործումը մեր նոր matmul- ի համար, մենք ստանում ենք առաջին կերելային ժամանակը (բեռնում է CPU- ի վերեւում): triton.testing.do_bench_cudagraph a, b = get_data(4096, 8192, 4096) triton.testing.do_bench_cudagraph(lambda: cpu_heavy_matmul(a, b)) # back to ~0.4ms! Ահա, CUDA գրաֆիկները ապրում են իրենց սեփական սահմանափակումներով (static shapes, static control flow եւ այլն: Նրանք ավելի լավ են հարմարվում GPU- ի որոշ բաղադրիչների բաղադրման համար, քան e2e դասընթացների բաղադրիչների համար: Դա է Դա է Ոչ միայն համակարգերը սխալ են - տեսնում են տվյալները Երբեմն նույնիսկ ճշգրիտ benchmarking մեթոդը չի խոսում ամբողջ պատմությունը, քանի որ արտադրության պայմանները տարբեր են միասին միջավայրի, որտեղ մենք աշխատում ենք benchmark- ը: Տեսեք, որ մենք ցանկանում ենք համեմատել համատեղելի gemm կետը մի mix-of-experts կետում: Հիմնական MoE կետում, tokens- ը մանրաթել է տարբեր փորձագետներին, որը հիմնված է router- ի հավանականության վրա: Այսպիսով, յուրաքանչյուր փորձագետի ( եւ համատեղելի gemm կետում) տեսող իրական աշխատանքային տպագրությունը հիմնված է այն, թե ինչպես հարմարված է routing- ը: Հիմնականորեն, մենք ստեղծում ենք անմիջական ռեժիմերային հավասարություններներ եւ պտուտակային tokens հիմնված նրանց վրա, սակայն, ամբողջ դասընթացում, մեր ռեժիմեր կարող է լինել անսահմանափակ, օգտագործելով մի քանի մասնագետներ շատ ավելի մեծ է, քան մյուսները. Տեսենք, թե ինչպիսիք են, որ յուրաքանչյուր մասնագետի համար ուղարկված tokens- ի բաղադրատոմսերը ազդում են մեր չափանիշներին: import numpy as np import torch import triton def sample_expert_assignments( seq_len: int, num_experts: int, top_k: int, use_beta: bool = True, alpha: float = 1.0, beta: float = 5.0, ) -> tuple[torch.Tensor, torch.Tensor]: if use_beta: expert_weights = np.random.beta(alpha, beta, num_experts) expert_weights = expert_weights / expert_weights.sum() else: expert_weights = np.ones(num_experts) / num_experts gumbel_noise = -np.log(-np.log(np.random.uniform(0, 1, (seq_len, num_experts)))) log_weights = np.log(expert_weights + 1e-10) scores = log_weights[None, :] + gumbel_noise expert_indices = torch.from_numpy(np.argsort(scores, axis=1)[:, -top_k:]) tokens_per_expert = torch.bincount( expert_indices.flatten(), minlength=num_experts ).to(torch.int32) return expert_indices, tokens_per_expert def get_tokens( seq_len: int, hidden: int, num_experts: int, top_k: int, expert_indices: torch.Tensor, tokens_per_expert: torch.Tensor, device: str = "cuda", ) -> torch.Tensor: x_original = torch.randn((seq_len, hidden), dtype=torch.bfloat16, device=device) token_indices = torch.arange(seq_len, device=device)[:, None].expand(-1, top_k) expert_flat = expert_indices.to(device).flatten() token_flat = token_indices.flatten() token_sorted = token_flat[torch.argsort(expert_flat)] return x_original[token_sorted] def get_tensors( seq_len: int = 1024, hidden: int = 4096, intermediate: int = 1536, num_experts: int = 128, top_k: int = 8, use_beta: bool = True, ) -> tuple[torch.Tensor, torch.Tensor, torch.Tensor]: expert_indices, tokens_per_expert = sample_expert_assignments( seq_len, num_experts, top_k, use_beta=use_beta, ) x = get_tokens( seq_len, hidden, num_experts, top_k, expert_indices, tokens_per_expert, device="cuda", ) w = torch.randn( (num_experts, hidden, intermediate), dtype=torch.bfloat16, device="cuda" ) offsets = torch.cumsum(tokens_per_expert, dim=0).to( dtype=torch.int32, device="cuda" ) return x, w, offsets def benchmark( seq_len: int = 1024, hidden: int = 4096, intermediate: int = 1536, num_experts: int = 128, top_k: int = 8, use_beta: bool = True, num_iters: int = 250, ): x, w, offsets = get_tensors( seq_len, hidden, intermediate, num_experts, top_k, use_beta, ) return triton.testing.do_bench_cudagraph( lambda: torch._grouped_mm(x, w, offs=offsets), rep=num_iters ) params = { "seq_len": 4096, "hidden": 4096, "intermediate": 1536, "num_experts": 128, "top_k": 8, } uniform = benchmark(**params, use_beta=False) beta = benchmark(**params, use_beta=True) print(f"Uniform: {uniform:.5f} ms") print(f"Beta: {beta:.5f} ms") # Uniform: 0.96811 ms # Beta: 1.02422 ms Շնորհակալություն, որ GroupedGEMM kernel- ը 6% արագ է, քան այն ժամանակ, երբ մենք նմուշում ենք բաա-նից (հարկե, մենք օգտագործում ենք սխալային բեռների մոդելման համար)! Միացեք ամեն ինչ Հաջորդը միասին: When writing the benchmarking code by hand, do not forget to: Flush the L2 cache Use CUDA events Use cuda.synchronize to wait for the completion of all GPU work Alternatively, use or for a built-in solution. triton.testing.do_bench triton.testing.do_bench_cudagraph Regardless of approach, do not forget about the underlying data you’re using. Բոլորում, ձեր մանրամասը միայն այնքան օգտակար է, քանի որ այն կարեւոր է այն մասին, ինչ մենք կարող ենք տեսնել իրական արտադրանքի կատարում: Եթե մենք վախենում ենք CPU overhead- ի մասին, չեք վախենում այնից: Եթե փորձեք մանրամասել այն ժամանակ, երբ մանրամասում եք: