This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
using Accum = cutlass::epilogue::fusion::Sm90AccFetch; | |
using ElementD = cutlass::half_t; | |
using StrideD = cute::Stride<int64_t, cute::Int<1>, cute::Int<0>>; | |
using Bias = cutlass::epilogue::fusion::Sm90ColBroadcast< | |
0 /*Stages*/, typename EpilogueDescriptor::TileShape, cutlass::half_t, cutlass::half_t, | |
cute::Stride<cute::Int<1>, cute::Int<0>, cute::Int<0>> | |
>; |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
using Accum = cutlass::epilogue::fusion::Sm90AccFetch; | |
using ElementD = cutlass::half_t; | |
using StrideD = cute::Stride<int64_t, cute::Int<1>, cute::Int<0>>; | |
using Bias = cutlass::epilogue::fusion::Sm90ColBroadcast< | |
0 /*Stages*/, typename EpilogueDescriptor::TileShape, cutlass::half_t, cutlass::half_t, | |
cute::Stride<cute::Int<1>, cute::Int<0>, cute::Int<0>> | |
>; |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
// Place your key bindings in this file to override the defaults | |
[ | |
{ | |
"key": "ctrl+tab", | |
"command": "-workbench.action.quickOpenPreviousRecentlyUsedEditorInGroup" | |
}, | |
{ | |
"key": "ctrl+tab", | |
"command": "workbench.action.nextEditor" | |
}, |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
from ctypes import c_void_p, c_long, c_int | |
import torch | |
import math | |
import random | |
import os | |
import tempfile | |
from math import inf, nan | |
from cmath import nanj | |
from torch._inductor.hooks import run_intermediate_hooks | |
from torch._inductor.utils import maybe_profile |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
import torch | |
from torch._inductor.utils import fresh_inductor_cache | |
torch._logging.set_logs(fusion=True) | |
with fresh_inductor_cache(): | |
@torch.compile() | |
def foo(x, y): | |
return (x @ y).permute(1, 0).relu().permute(1, 0).sigmoid() |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
import torch | |
torch._logging.set_logs(fusion=True) | |
@torch.compile() | |
def foo(x, y): | |
return (x @ y).permute(1, 0).relu().permute(1, 0).contiguous() | |
foo(torch.ones(256, 256, device="cuda"), torch.ones(256, 256, device="cuda")) |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
import torch | |
torch._logging.set_logs(fusion=True) | |
@torch.compile() | |
def foo(x, y): | |
return (x @ y).permute(1, 0).relu().permute(1, 0).contiguous() | |
foo(torch.ones(256, 256, device="cuda"), torch.ones(256, 256, device="cuda")) |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
# mypy: allow-untyped-defs | |
from numpy import dtype | |
from torch._inductor.ir import ComputedBuffer, InputBuffer | |
from typing import Union | |
from ..cutlass_utils import try_import_cutlass | |
if try_import_cutlass(): |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
def render_evt(tile_desc, epilogue_schedule): | |
from .cutlass_lib_extensions import evt_extensions | |
LayoutType = cutlass_lib.LayoutType | |
DataType = cutlass_lib.DataType | |
CutlassTensor = evt_extensions.CutlassTensor | |
bias_code = """def example_epilogue(accum, alpha, C, beta, aux, bias): | |
F = alpha * accum + (beta * C + aux) | |
E = relu(F + 1) + bias |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
ERROR:common: | |
Traceback (most recent call last): | |
File "/data/users/mlazos/pytorch/benchmarks/dynamo/common.py", line 2216, in check_accuracy | |
new_result = self.run_n_iterations( | |
File "/data/users/mlazos/pytorch/benchmarks/dynamo/common.py", line 1930, in run_n_iterations | |
model_iter_fn(mod, inputs, collect_outputs=False) | |
File "/data/users/mlazos/pytorch/torch/_dynamo/eval_frame.py", line 658, in _fn | |
return fn(*args, **kwargs) | |
File "/data/users/mlazos/pytorch/benchmarks/dynamo/huggingface.py", line 531, in forward_and_backward_pass | |
cloned_inputs = clone_inputs(inputs) |