Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

[int8 woq] make the scale type the same as input for bf16 autocast #534

Merged
merged 1 commit into from
Jul 29, 2024

Conversation

Valentine233
Copy link
Contributor

Under bf16 autocast, input's type would convert from fp32 to bf16 because of torch.mm. However, scale's type is still fp32, so as the final output. To fix the issue, we make scale's type the same as the output of torch.mm, to get a bf16 scale.

FX graph before:

permute_2: "i8[4096, 4096]" = torch.ops.aten.permute.default(arg68_1, [1, 0]);  arg68_1 = None
view_3: "f32[128, 4096]" = torch.ops.aten.view.default(mul_1, [-1, 4096])
convert_element_type_default_190: "bf16[4096, 4096]" = torch.ops.prims.convert_element_type.default(permute_2, torch.bfloat16);  permute_2 = None
convert_element_type_7: "bf16[128, 4096]" = torch.ops.prims.convert_element_type.default(view_3, torch.bfloat16);  view_3 = None
mm_1: "bf16[128, 4096]" = torch.ops.aten.mm.default(convert_element_type_7, convert_element_type_default_190);  convert_element_type_7 = convert_element_type_default_190 = None
mul_3: "f32[128, 4096]" = torch.ops.aten.mul.Tensor(mm_1, arg69_1);  mm_1 = arg69_1 = None
view_4: "f32[4, 32, 4096]" = torch.ops.aten.view.default(mul_3, [4, 32, 4096]);  mul_3 = None

FX graph after:

permute_2: "i8[4096, 4096]" = torch.ops.aten.permute.default(arg68_1, [1, 0]);  arg68_1 = None
view_3: "f32[128, 4096]" = torch.ops.aten.view.default(mul_1, [-1, 4096])
convert_element_type_default_158: "bf16[4096, 4096]" = torch.ops.prims.convert_element_type.default(permute_2, torch.bfloat16);  permute_2 = None
convert_element_type_8: "bf16[128, 4096]" = torch.ops.prims.convert_element_type.default(view_3, torch.bfloat16);  view_3 = None
mm_1: "bf16[128, 4096]" = torch.ops.aten.mm.default(convert_element_type_8, convert_element_type_default_158);  convert_element_type_8 = convert_element_type_default_158 = None
convert_element_type_11: "bf16[4096]" = torch.ops.prims.convert_element_type.default(arg69_1, torch.bfloat16);  arg69_1 = None
mul_3: "bf16[128, 4096]" = torch.ops.aten.mul.Tensor(mm_1, convert_element_type_11);  mm_1 = convert_element_type_11 = None
view_4: "bf16[4, 32, 4096]" = torch.ops.aten.view.default(mul_3, [4, 32, 4096]);  mul_3 = None

Copy link

pytorch-bot bot commented Jul 23, 2024

🔗 Helpful Links

🧪 See artifacts and rendered test results at hud.pytorch.org/pr/pytorch/ao/534

Note: Links to docs will display an error until the docs builds have been completed.

✅ No Failures

As of commit ea32965 with merge base 5787e9e (image):
💚 Looks good so far! There are no failures yet. 💚

This comment was automatically generated by Dr. CI and updates every 15 minutes.

@facebook-github-bot facebook-github-bot added the CLA Signed This label is managed by the Facebook bot. Authors need to sign the CLA before a PR can be reviewed. label Jul 23, 2024
@Valentine233
Copy link
Contributor Author

@jerryzh168 @jgong5 @leslie-fang-intel Please help review the PR, thanks!

@jerryzh168 jerryzh168 requested a review from HDCharles July 23, 2024 21:26
@jerryzh168
Copy link
Contributor

can you check the output_code diff for TORCH_LOGS='output_code' python tutorials/quantize_vit/run_vit_b_quant.py using int8 weight only quant? (change

quantize_(model, int8_dynamic_activation_int8_weight())
to use int8_weight_only())

@Valentine233
Copy link
Contributor Author

can you check the output_code diff for TORCH_LOGS='output_code' python tutorials/quantize_vit/run_vit_b_quant.py using int8 weight only quant? (change

quantize_(model, int8_dynamic_activation_int8_weight())

to use int8_weight_only())

Hi Jerry, I found no output_code diff and here is the output_code:

V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code] Output code:
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code] # AOT ID: ['0_inference']
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code] from ctypes import c_void_p, c_long
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code] import torch
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code] import math
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code] import random
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code] import os
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code] import tempfile
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code] from math import inf, nan
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code] from torch._inductor.hooks import run_intermediate_hooks
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code] from torch._inductor.utils import maybe_profile
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code] from torch._inductor.codegen.memory_planning import _align as align
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code] from torch import device, empty_strided
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code] from torch._inductor.async_compile import AsyncCompile
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code] from torch._inductor.select_algorithm import extern_kernels
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code] from torch._inductor.codegen.multi_kernel import MultiKernelCall
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code] aten = torch.ops.aten
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code] inductor_ops = torch.ops.inductor
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code] _quantized = torch.ops._quantized
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code] assert_size_stride = torch._C._dynamo.guards.assert_size_stride
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code] empty_strided_cpu = torch._C._dynamo.guards._empty_strided_cpu
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code] empty_strided_cuda = torch._C._dynamo.guards._empty_strided_cuda
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code] reinterpret_tensor = torch._C._dynamo.guards._reinterpret_tensor
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code] alloc_from_pool = torch.ops.inductor._alloc_from_pool
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code] async_compile = AsyncCompile()
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code] cpp_fused_convolution_0 = async_compile.cpp_pybinding(['const bfloat16*', 'const bfloat16*', 'bfloat16*', 'bfloat16*'], '''
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code] #include "/tmp/torchinductor_liaoxuan/ky/cky2bufythacofebk7ujv36e4pxyqcqbpsy5r4vojoprjiwcwfxf.h"
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code] extern "C"  void kernel(const bfloat16* in_ptr0,
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]                        const bfloat16* in_ptr1,
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]                        bfloat16* out_ptr0,
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]                        bfloat16* out_ptr1)
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code] {
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]     {
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]         #pragma GCC ivdep
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]         for(long x0=static_cast<long>(0L); x0<static_cast<long>(3L); x0+=static_cast<long>(1L))
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]         {
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]             #pragma GCC ivdep
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]             for(long x1=static_cast<long>(0L); x1<static_cast<long>(50176L); x1+=static_cast<long>(1L))
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]             {
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]                 auto tmp0 = in_ptr0[static_cast<long>(x1 + (50176L*x0))];
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]                 out_ptr0[static_cast<long>(x0 + (3L*x1))] = tmp0;
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]             }
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]         }
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]     }
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]     #pragma omp parallel num_threads(112)
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]     {
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]         int tid = omp_get_thread_num();
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]         {
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]             #pragma omp for
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]             for(long x0=static_cast<long>(0L); x0<static_cast<long>(768L); x0+=static_cast<long>(1L))
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]             {
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]                 #pragma GCC ivdep
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]                 for(long x1=static_cast<long>(0L); x1<static_cast<long>(3L); x1+=static_cast<long>(1L))
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]                 {
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]                     #pragma GCC ivdep
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]                     for(long x2=static_cast<long>(0L); x2<static_cast<long>(256L); x2+=static_cast<long>(1L))
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]                     {
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]                         auto tmp0 = in_ptr1[static_cast<long>(x2 + (256L*x1) + (768L*x0))];
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]                         out_ptr1[static_cast<long>(x1 + (3L*x2) + (768L*x0))] = tmp0;
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]                     }
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]                 }
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]             }
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]         }
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]     }
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code] }
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code] ''')
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code] cpp_fused_cat_1 = async_compile.cpp_pybinding(['const bfloat16*', 'const bfloat16*', 'bfloat16*', 'bfloat16*'], '''
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code] #include "/tmp/torchinductor_liaoxuan/ky/cky2bufythacofebk7ujv36e4pxyqcqbpsy5r4vojoprjiwcwfxf.h"
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code] extern "C"  void kernel(const bfloat16* in_ptr0,
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]                        const bfloat16* in_ptr1,
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]                        bfloat16* out_ptr0,
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]                        bfloat16* out_ptr1)
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code] {
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]     {
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]         for(long x0=static_cast<long>(0L); x0<static_cast<long>(768L); x0+=static_cast<long>(32L))
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]         {
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]             auto tmp0 = at::vec::Vectorized<bfloat16>::loadu(in_ptr0 + static_cast<long>(x0), 32);
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]             tmp0.store(out_ptr0 + static_cast<long>(x0), 32);
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]         }
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]     }
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]     {
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]         for(long x0=static_cast<long>(0L); x0<static_cast<long>(150528L); x0+=static_cast<long>(32L))
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]         {
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]             auto tmp0 = at::vec::Vectorized<bfloat16>::loadu(in_ptr1 + static_cast<long>(x0), 32);
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]             tmp0.store(out_ptr1 + static_cast<long>(x0), 32);
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]         }
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]     }
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code] }
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code] ''')
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code] async_compile.wait(globals())
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code] del async_compile
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code] def call(args):
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]     arg0_1, arg1_1, arg2_1, arg3_1 = args
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]     args.clear()
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]     assert_size_stride(arg0_1, (1, 1, 768), (768, 768, 1))
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]     assert_size_stride(arg1_1, (768, 3, 16, 16), (768, 256, 16, 1))
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]     assert_size_stride(arg2_1, (768, ), (1, ))
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]     assert_size_stride(arg3_1, (1, 3, 224, 224), (150528, 50176, 224, 1))
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]     buf0 = empty_strided_cpu((1, 3, 224, 224), (150528, 1, 672, 3), torch.bfloat16)
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]     buf1 = empty_strided_cpu((768, 3, 16, 16), (768, 1, 48, 3), torch.bfloat16)
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]     cpp_fused_convolution_0(arg3_1, arg1_1, buf0, buf1)
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]     del arg1_1
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]     del arg3_1
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]     # Source Nodes: [x], Original ATen: [aten.convolution]
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]     buf2 = extern_kernels.convolution(buf0, buf1, arg2_1, stride=(16, 16), padding=(0, 0), dilation=(1, 1), transposed=False, output_padding=(0, 0), groups=1)
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]     assert_size_stride(buf2, (1, 768, 14, 14), (150528, 1, 10752, 768))
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]     del arg2_1
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]     del buf0
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]     del buf1
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]     buf5 = empty_strided_cpu((1, 197, 768), (151296, 768, 1), torch.bfloat16)
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]     buf3 = reinterpret_tensor(buf5, (1, 1, 768), (151296, 768, 1), 0)  # alias
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]     buf4 = reinterpret_tensor(buf5, (1, 196, 768), (151296, 768, 1), 768)  # alias
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]     cpp_fused_cat_1(arg0_1, buf2, buf3, buf4)
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]     del arg0_1
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]     return (buf5, )
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code] def benchmark_compiled_module(times=10, repeat=10):
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]     from torch._dynamo.testing import rand_strided
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]     from torch._inductor.utils import print_performance
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]     arg0_1 = rand_strided((1, 1, 768), (768, 768, 1), device='cpu', dtype=torch.bfloat16)
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]     arg1_1 = rand_strided((768, 3, 16, 16), (768, 256, 16, 1), device='cpu', dtype=torch.bfloat16)
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]     arg2_1 = rand_strided((768, ), (1, ), device='cpu', dtype=torch.bfloat16)
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]     arg3_1 = rand_strided((1, 3, 224, 224), (150528, 50176, 224, 1), device='cpu', dtype=torch.bfloat16)
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]     fn = lambda: call([arg0_1, arg1_1, arg2_1, arg3_1])
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]     return print_performance(fn, times=times, repeat=repeat)
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code] if __name__ == "__main__":
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]     from torch._inductor.wrapper_benchmark import compiled_module_main
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]     compiled_module_main('None', benchmark_compiled_module)
V0723 18:40:34.857000 2313084 torch/_inductor/codecache.py:948] [0/0_1] [__output_code]
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code] Output code:
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code]
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code] # AOT ID: ['1_inference']
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code] from ctypes import c_void_p, c_long
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code] import torch
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code] import math
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code] import random
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code] import os
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code] import tempfile
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code] from math import inf, nan
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code] from torch._inductor.hooks import run_intermediate_hooks
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code] from torch._inductor.utils import maybe_profile
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code] from torch._inductor.codegen.memory_planning import _align as align
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code]
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code] from torch import device, empty_strided
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code] from torch._inductor.async_compile import AsyncCompile
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code] from torch._inductor.select_algorithm import extern_kernels
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code] from torch._inductor.codegen.multi_kernel import MultiKernelCall
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code]
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code] aten = torch.ops.aten
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code] inductor_ops = torch.ops.inductor
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code] _quantized = torch.ops._quantized
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code] assert_size_stride = torch._C._dynamo.guards.assert_size_stride
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code] empty_strided_cpu = torch._C._dynamo.guards._empty_strided_cpu
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code] empty_strided_cuda = torch._C._dynamo.guards._empty_strided_cuda
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code] reinterpret_tensor = torch._C._dynamo.guards._reinterpret_tensor
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code] alloc_from_pool = torch.ops.inductor._alloc_from_pool
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code] async_compile = AsyncCompile()
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code]
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code]
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code] cpp_fused_add_0 = async_compile.cpp_pybinding(['const bfloat16*', 'const bfloat16*', 'bfloat16*'], '''
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code] #include "/tmp/torchinductor_liaoxuan/ky/cky2bufythacofebk7ujv36e4pxyqcqbpsy5r4vojoprjiwcwfxf.h"
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code] extern "C"  void kernel(const bfloat16* in_ptr0,
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code]                        const bfloat16* in_ptr1,
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code]                        bfloat16* out_ptr0)
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code] {
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code]     {
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code]         for(long x0=static_cast<long>(0L); x0<static_cast<long>(151296L); x0+=static_cast<long>(16L))
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code]         {
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code]             auto tmp0 = at::vec::Vectorized<bfloat16>::loadu(in_ptr0 + static_cast<long>(x0), 16);
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code]             auto tmp2 = at::vec::Vectorized<bfloat16>::loadu(in_ptr1 + static_cast<long>(x0), 16);
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code]             auto tmp1 = at::vec::convert<float>(tmp0);
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code]             auto tmp3 = at::vec::convert<float>(tmp2);
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code]             auto tmp4 = tmp1 + tmp3;
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code]             auto tmp5 = at::vec::convert<bfloat16>(tmp4);
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code]             tmp5.store(out_ptr0 + static_cast<long>(x0), 16);
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code]         }
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code]     }
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code] }
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code] ''')
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code]
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code]
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code] async_compile.wait(globals())
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code] del async_compile
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code]
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code] def call(args):
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code]     arg0_1, arg1_1 = args
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code]     args.clear()
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code]     assert_size_stride(arg0_1, (1, 197, 768), (151296, 768, 1))
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code]     assert_size_stride(arg1_1, (1, 197, 768), (151296, 768, 1))
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code]     buf0 = empty_strided_cpu((1, 197, 768), (151296, 768, 1), torch.bfloat16)
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code]     cpp_fused_add_0(arg1_1, arg0_1, buf0)
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code]     del arg0_1
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code]     del arg1_1
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code]     return (buf0, )
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code]
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code]
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code] def benchmark_compiled_module(times=10, repeat=10):
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code]     from torch._dynamo.testing import rand_strided
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code]     from torch._inductor.utils import print_performance
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code]     arg0_1 = rand_strided((1, 197, 768), (151296, 768, 1), device='cpu', dtype=torch.bfloat16)
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code]     arg1_1 = rand_strided((1, 197, 768), (151296, 768, 1), device='cpu', dtype=torch.bfloat16)
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code]     fn = lambda: call([arg0_1, arg1_1])
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code]     return print_performance(fn, times=times, repeat=repeat)
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code]
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code]
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code] if __name__ == "__main__":
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code]     from torch._inductor.wrapper_benchmark import compiled_module_main
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code]     compiled_module_main('None', benchmark_compiled_module)
V0723 18:40:34.978000 2313084 torch/_inductor/codecache.py:948] [1/0_1] [__output_code]
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code] Output code:
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code] # AOT ID: ['2_inference']
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code] from ctypes import c_void_p, c_long
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code] import torch
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code] import math
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code] import random
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code] import os
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code] import tempfile
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code] from math import inf, nan
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code] from torch._inductor.hooks import run_intermediate_hooks
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code] from torch._inductor.utils import maybe_profile
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code] from torch._inductor.codegen.memory_planning import _align as align
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code] from torch import device, empty_strided
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code] from torch._inductor.async_compile import AsyncCompile
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code] from torch._inductor.select_algorithm import extern_kernels
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code] from torch._inductor.codegen.multi_kernel import MultiKernelCall
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code] aten = torch.ops.aten
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code] inductor_ops = torch.ops.inductor
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code] _quantized = torch.ops._quantized
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code] assert_size_stride = torch._C._dynamo.guards.assert_size_stride
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code] empty_strided_cpu = torch._C._dynamo.guards._empty_strided_cpu
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code] empty_strided_cuda = torch._C._dynamo.guards._empty_strided_cuda
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code] reinterpret_tensor = torch._C._dynamo.guards._reinterpret_tensor
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code] alloc_from_pool = torch.ops.inductor._alloc_from_pool
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code] async_compile = AsyncCompile()
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code] cpp_fused_native_layer_norm_0 = async_compile.cpp_pybinding(['const bfloat16*', 'const bfloat16*', 'const bfloat16*', 'float*', 'float*', 'bfloat16*'], '''
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code] #include "/tmp/torchinductor_liaoxuan/ky/cky2bufythacofebk7ujv36e4pxyqcqbpsy5r4vojoprjiwcwfxf.h"
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code] extern "C"  void kernel(const bfloat16* in_ptr0,
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]                        const bfloat16* in_ptr1,
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]                        const bfloat16* in_ptr2,
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]                        float* out_ptr0,
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]                        float* out_ptr1,
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]                        bfloat16* out_ptr2)
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code] {
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]     {
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]         #pragma GCC ivdep
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]         for(long x0=static_cast<long>(0L); x0<static_cast<long>(197L); x0+=static_cast<long>(1L))
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]         {
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]             {
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]                 Welford<float> tmp_acc0 = Welford<float>();
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]                 Welford<at::vec::Vectorized<float>> tmp_acc0_vec = Welford<at::vec::Vectorized<float>>();
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]                 static WeightRecp<at::vec::Vectorized<float>> weight_recps(static_cast<long>(48L));
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]                 for(long x1=static_cast<long>(0L); x1<static_cast<long>(768L); x1+=static_cast<long>(16L))
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]                 {
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]                     auto tmp0 = at::vec::Vectorized<bfloat16>::loadu(in_ptr0 + static_cast<long>(x1 + (768L*x0)), 16);
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]                     auto tmp1 = at::vec::convert<float>(tmp0);
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]                     tmp_acc0_vec = welford_combine(tmp_acc0_vec, tmp1, &weight_recps);
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]                 }
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]                 tmp_acc0 = welford_combine(tmp_acc0, welford_vec_reduce_all(tmp_acc0_vec));
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]                 out_ptr0[static_cast<long>(x0)] = static_cast<float>(tmp_acc0.mean);
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]                 out_ptr1[static_cast<long>(x0)] = static_cast<float>(tmp_acc0.m2);
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]             }
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]             for(long x1=static_cast<long>(0L); x1<static_cast<long>(768L); x1+=static_cast<long>(16L))
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]             {
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]                 auto tmp0 = at::vec::Vectorized<bfloat16>::loadu(in_ptr0 + static_cast<long>(x1 + (768L*x0)), 16);
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]                 auto tmp2 = out_ptr0[static_cast<long>(x0)];
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]                 auto tmp5 = out_ptr1[static_cast<long>(x0)];
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]                 auto tmp13 = at::vec::Vectorized<bfloat16>::loadu(in_ptr1 + static_cast<long>(x1), 16);
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]                 auto tmp16 = at::vec::Vectorized<bfloat16>::loadu(in_ptr2 + static_cast<long>(x1), 16);
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]                 auto tmp1 = at::vec::convert<float>(tmp0);
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]                 auto tmp3 = at::vec::Vectorized<float>(tmp2);
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]                 auto tmp4 = tmp1 - tmp3;
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]                 auto tmp6 = static_cast<float>(768.0);
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]                 auto tmp7 = tmp5 / tmp6;
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]                 auto tmp8 = static_cast<float>(1e-06);
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]                 auto tmp9 = decltype(tmp7)(tmp7 + tmp8);
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]                 auto tmp10 = 1 / std::sqrt(tmp9);
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]                 auto tmp11 = at::vec::Vectorized<float>(tmp10);
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]                 auto tmp12 = tmp4 * tmp11;
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]                 auto tmp14 = at::vec::convert<float>(tmp13);
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]                 auto tmp15 = tmp12 * tmp14;
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]                 auto tmp17 = at::vec::convert<float>(tmp16);
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]                 auto tmp18 = tmp15 + tmp17;
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]                 auto tmp19 = at::vec::convert<bfloat16>(tmp18);
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]                 tmp19.store(out_ptr2 + static_cast<long>(x1 + (768L*x0)), 16);
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]             }
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]         }
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]     }
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code] }
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code] ''')
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code] async_compile.wait(globals())
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code] del async_compile
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code] def call(args):
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]     arg0_1, arg1_1, arg2_1 = args
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]     args.clear()
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]     assert_size_stride(arg0_1, (768, ), (1, ))
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]     assert_size_stride(arg1_1, (768, ), (1, ))
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]     assert_size_stride(arg2_1, (1, 197, 768), (151296, 768, 1))
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]     buf0 = empty_strided_cpu((1, 197, 1), (197, 1, 197), torch.float32)
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]     buf1 = empty_strided_cpu((1, 197, 1), (197, 1, 197), torch.float32)
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]     buf3 = empty_strided_cpu((1, 197, 768), (151296, 768, 1), torch.bfloat16)
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]     cpp_fused_native_layer_norm_0(arg2_1, arg0_1, arg1_1, buf0, buf1, buf3)
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]     del arg0_1
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]     del arg1_1
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]     del arg2_1
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]     return (buf3, )
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code] def benchmark_compiled_module(times=10, repeat=10):
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]     from torch._dynamo.testing import rand_strided
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]     from torch._inductor.utils import print_performance
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]     arg0_1 = rand_strided((768, ), (1, ), device='cpu', dtype=torch.bfloat16)
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]     arg1_1 = rand_strided((768, ), (1, ), device='cpu', dtype=torch.bfloat16)
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]     arg2_1 = rand_strided((1, 197, 768), (151296, 768, 1), device='cpu', dtype=torch.bfloat16)
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]     fn = lambda: call([arg0_1, arg1_1, arg2_1])
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]     return print_performance(fn, times=times, repeat=repeat)
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code] if __name__ == "__main__":
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]     from torch._inductor.wrapper_benchmark import compiled_module_main
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]     compiled_module_main('None', benchmark_compiled_module)
V0723 18:40:35.032000 2313084 torch/_inductor/codecache.py:948] [2/0_1] [__output_code]

In addition, although the output_code can be printed, I encountered an error no matter this PR was applied or not.

Traceback (most recent call last):
  File "/home/liaoxuan/ao/tutorials/quantize_vit/run_vit_b_quant.py", line 42, in <module>
    benchmark_model(model, 20, input_tensor)
  File "/home/liaoxuan/anaconda3/envs/ao/lib/python3.10/site-packages/torchao-0.3.1-py3.10-linux-x86_64.egg/torchao/utils.py", line 84, in benchmark_model
    model(input_tensor)
  File "/home/liaoxuan/pytorch_ao/torch/nn/modules/module.py", line 1716, in _wrapped_call_impl
    return self._call_impl(*args, **kwargs)
  File "/home/liaoxuan/pytorch_ao/torch/nn/modules/module.py", line 1727, in _call_impl
    return forward_call(*args, **kwargs)
  File "/home/liaoxuan/pytorch_ao/torch/_dynamo/eval_frame.py", line 448, in _fn
    return fn(*args, **kwargs)
  File "/home/liaoxuan/pytorch_ao/torch/nn/modules/module.py", line 1716, in _wrapped_call_impl
    return self._call_impl(*args, **kwargs)
  File "/home/liaoxuan/pytorch_ao/torch/nn/modules/module.py", line 1727, in _call_impl
    return forward_call(*args, **kwargs)
  File "/home/liaoxuan/anaconda3/envs/ao/lib/python3.10/site-packages/torchvision-0.20.0a0+c2619ba-py3.10-linux-x86_64.egg/torchvision/models/vision_transformer.py", line 298, in forward
    x = self.encoder(x)
  File "/home/liaoxuan/pytorch_ao/torch/nn/modules/module.py", line 1716, in _wrapped_call_impl
    return self._call_impl(*args, **kwargs)
  File "/home/liaoxuan/pytorch_ao/torch/nn/modules/module.py", line 1727, in _call_impl
    return forward_call(*args, **kwargs)
  File "/home/liaoxuan/anaconda3/envs/ao/lib/python3.10/site-packages/torchvision-0.20.0a0+c2619ba-py3.10-linux-x86_64.egg/torchvision/models/vision_transformer.py", line 157, in forward
    return self.ln(self.layers(self.dropout(input)))
  File "/home/liaoxuan/pytorch_ao/torch/nn/modules/module.py", line 1716, in _wrapped_call_impl
    return self._call_impl(*args, **kwargs)
  File "/home/liaoxuan/pytorch_ao/torch/nn/modules/module.py", line 1727, in _call_impl
    return forward_call(*args, **kwargs)
  File "/home/liaoxuan/pytorch_ao/torch/nn/modules/container.py", line 249, in forward
    input = module(input)
  File "/home/liaoxuan/pytorch_ao/torch/nn/modules/module.py", line 1716, in _wrapped_call_impl
    return self._call_impl(*args, **kwargs)
  File "/home/liaoxuan/pytorch_ao/torch/nn/modules/module.py", line 1727, in _call_impl
    return forward_call(*args, **kwargs)
  File "/home/liaoxuan/anaconda3/envs/ao/lib/python3.10/site-packages/torchvision-0.20.0a0+c2619ba-py3.10-linux-x86_64.egg/torchvision/models/vision_transformer.py", line 113, in forward
    x, _ = self.self_attention(x, x, x, need_weights=False)
  File "/home/liaoxuan/pytorch_ao/torch/nn/modules/module.py", line 1716, in _wrapped_call_impl
    return self._call_impl(*args, **kwargs)
  File "/home/liaoxuan/pytorch_ao/torch/nn/modules/module.py", line 1727, in _call_impl
    return forward_call(*args, **kwargs)
  File "/home/liaoxuan/pytorch_ao/torch/nn/modules/activation.py", line 1281, in forward
    self.out_proj.weight,
  File "/home/liaoxuan/pytorch_ao/torch/nn/utils/parametrize.py", line 406, in get_parametrized
    return parametrization()
  File "/home/liaoxuan/pytorch_ao/torch/nn/modules/module.py", line 1716, in _wrapped_call_impl
    return self._call_impl(*args, **kwargs)
  File "/home/liaoxuan/pytorch_ao/torch/nn/modules/module.py", line 1727, in _call_impl
    return forward_call(*args, **kwargs)
  File "/home/liaoxuan/pytorch_ao/torch/nn/utils/parametrize.py", line 302, in forward
    x = self[0](*originals)
  File "/home/liaoxuan/pytorch_ao/torch/nn/modules/module.py", line 1716, in _wrapped_call_impl
    return self._call_impl(*args, **kwargs)
  File "/home/liaoxuan/pytorch_ao/torch/nn/modules/module.py", line 1727, in _call_impl
    return forward_call(*args, **kwargs)
  File "/home/liaoxuan/anaconda3/envs/ao/lib/python3.10/site-packages/torchao-0.3.1-py3.10-linux-x86_64.egg/torchao/utils.py", line 221, in forward
    rebuilt = tp.__tensor_unflatten__(inner_tensors, meta, None, None)
  File "/home/liaoxuan/anaconda3/envs/ao/lib/python3.10/site-packages/torchao-0.3.1-py3.10-linux-x86_64.egg/torchao/dtypes/affine_quantized_tensor.py", line 417, in __tensor_unflatten__
    return cls(int_data, scale, zero_point, layout_type)
  File "/home/liaoxuan/pytorch_ao/torch/_dynamo/convert_frame.py", line 1170, in __call__
    return self._torchdynamo_orig_callable(
  File "/home/liaoxuan/pytorch_ao/torch/_dynamo/convert_frame.py", line 979, in __call__
    result = self._inner_convert(
  File "/home/liaoxuan/pytorch_ao/torch/_dynamo/convert_frame.py", line 499, in __call__
    return _compile(
  File "/home/liaoxuan/anaconda3/envs/ao/lib/python3.10/contextlib.py", line 79, in inner
    return func(*args, **kwds)
  File "/home/liaoxuan/pytorch_ao/torch/_dynamo/convert_frame.py", line 877, in _compile
    raise InternalTorchDynamoError(str(e)).with_traceback(
  File "/home/liaoxuan/pytorch_ao/torch/_dynamo/convert_frame.py", line 850, in _compile
    guarded_code = compile_inner(code, one_graph, hooks, transform)
  File "/home/liaoxuan/pytorch_ao/torch/_dynamo/utils.py", line 246, in time_wrapper
    r = func(*args, **kwargs)
  File "/home/liaoxuan/pytorch_ao/torch/_utils_internal.py", line 85, in wrapper_function
    return StrobelightCompileTimeProfiler.profile_compile_time(
  File "/home/liaoxuan/pytorch_ao/torch/_strobelight/compile_time_profiler.py", line 129, in profile_compile_time
    return func(*args, **kwargs)
  File "/home/liaoxuan/pytorch_ao/torch/_dynamo/convert_frame.py", line 668, in compile_inner
    out_code = transform_code_object(code, transform)
  File "/home/liaoxuan/pytorch_ao/torch/_dynamo/bytecode_transformation.py", line 1284, in transform_code_object
    transformations(instructions, code_options)
  File "/home/liaoxuan/pytorch_ao/torch/_dynamo/convert_frame.py", line 194, in _fn
    return fn(*args, **kwargs)
  File "/home/liaoxuan/pytorch_ao/torch/_dynamo/convert_frame.py", line 610, in transform
    tracer.run()
  File "/home/liaoxuan/pytorch_ao/torch/_dynamo/symbolic_convert.py", line 2546, in run
    super().run()
  File "/home/liaoxuan/pytorch_ao/torch/_dynamo/symbolic_convert.py", line 910, in run
    while self.step():
  File "/home/liaoxuan/pytorch_ao/torch/_dynamo/symbolic_convert.py", line 822, in step
    self.dispatch_table[inst.opcode](self, inst)
  File "/home/liaoxuan/pytorch_ao/torch/_dynamo/symbolic_convert.py", line 1622, in STORE_ATTR
    if isinstance(obj, NNModuleVariable) and not isinstance(val, ConstantVariable):
  File "/home/liaoxuan/pytorch_ao/torch/_dynamo/variables/base.py", line 106, in __instancecheck__
    instance = instance.realize()
  File "/home/liaoxuan/pytorch_ao/torch/_dynamo/variables/lazy.py", line 58, in realize
    self._cache.realize()
  File "/home/liaoxuan/pytorch_ao/torch/_dynamo/variables/lazy.py", line 24, in realize
    self.vt = VariableBuilder(tx, self.source)(self.value)
  File "/home/liaoxuan/pytorch_ao/torch/_dynamo/variables/builder.py", line 329, in __call__
    vt = self._wrap(value)
  File "/home/liaoxuan/pytorch_ao/torch/_dynamo/variables/builder.py", line 508, in _wrap
    return self.wrap_tensor(value)
  File "/home/liaoxuan/pytorch_ao/torch/_dynamo/variables/builder.py", line 1402, in wrap_tensor
    self.assert_not_wrapped_by_this_graph(value)
  File "/home/liaoxuan/pytorch_ao/torch/_dynamo/variables/builder.py", line 1318, in assert_not_wrapped_by_this_graph
    if is_fake(value) and maybe_get_fake_mode(value) is self.tx.fake_mode:
  File "/home/liaoxuan/pytorch_ao/torch/_subclasses/fake_tensor.py", line 172, in is_fake
    attrs, _ = type(x).__tensor_flatten__(x)
  File "/home/liaoxuan/anaconda3/envs/ao/lib/python3.10/site-packages/torchao-0.3.1-py3.10-linux-x86_64.egg/torchao/dtypes/affine_quantized_tensor.py", line 409, in __tensor_flatten__
    return ["int_data", "scale", "zero_point"], [self.layout_type]
torch._dynamo.exc.InternalTorchDynamoError: 'PlainAQTLayout' object has no attribute 'layout_type'

from user code:
   File "/home/liaoxuan/anaconda3/envs/ao/lib/python3.10/site-packages/torchao-0.3.1-py3.10-linux-x86_64.egg/torchao/dtypes/affine_quantized_tensor.py", line 403, in __init__
    self.int_data = int_data

@Valentine233
Copy link
Contributor Author

@HDCharles @jerryzh168 Could you help review? Thanks.

@jerryzh168
Copy link
Contributor

please make sure the test passes before landing

@Valentine233 Valentine233 force-pushed the int8_woq_scale_type branch from 7b92973 to ea32965 Compare July 27, 2024 01:15
@jerryzh168
Copy link
Contributor

the error might be the same one as #515, I'll see if I can repro as well next week

@Valentine233
Copy link
Contributor Author

Hi @jerryzh168, how could I re-trigger the CIs?

@jerryzh168
Copy link
Contributor

Hi @jerryzh168, how could I re-trigger the CIs?

looks like it requires approval, I just approved the CI to run

@jerryzh168 jerryzh168 merged commit ec317fc into pytorch:main Jul 29, 2024
20 of 22 checks passed
yanbing-j pushed a commit to yanbing-j/ao that referenced this pull request Dec 9, 2024
yanbing-j pushed a commit to yanbing-j/ao that referenced this pull request Dec 9, 2024
* make --device fast the default

* Update iOS.md (pytorch#517)

* Update iOS.md

* Update iOS.md

* Pip to pip3 (pytorch#504)

* remove macos-12 test

* pip to pip3

* break aoti CI jobs separately (pytorch#500)

* init

* fixes

* more fixes

* fixes

* fix

* fix

* bug fix

* add objcopy update

* suppress int8

* undefined variable

---------

Co-authored-by: Michael Gschwind <[email protected]>

* Support llama3 in chat in run.cpp  (pytorch#486)

* refactor chat runner in preparation for llama3

* add sketch for llama3 prompt template and move to returning tokens

* fix tiktoken

* fixes to chat

* add default llama_ver

* Add tests for quantize json, add cuda device specification and precision to cuda.json (pytorch#519)

* remove code for no KV Cache path (pytorch#527)

* Update ADVANCED-USERS.md (pytorch#529)

Update Advanced Users description to reflect changes in the repo since the description was initially created.

* runner-aoti on cuda (pytorch#531)

* runner-aoti on cuda

* transfer results back to CPU

* transfer results back to CPU

* runner-aoti on cuda

* Update runner_build.md (pytorch#530)

Update description of runner and build process in runner_build.md

* clean up runner code a little (pytorch#532)

* clean up runner code a little

* update

* update

* pull out generate loop in chat

* updates

* edit docs

* typo

* move int8 linear class and function into qops.py (pytorch#534)

* add dtype tests for runner-aoti + runner-et (pytorch#539)

* add dtype tests for runner-aoti + runner-et

* typo

* Quantized embedding (pytorch#536)

* move int8 linear class and function into qops.py

* move Quantized Embedding to qops.py

* Move Linear int4 to qops (pytorch#537)

* move int8 linear class and function into qops.py

* move Quantized Embedding to qops.py

* move int4 linear to qops

* Revert "add dtype tests for runner-aoti + runner-et (pytorch#539)" (pytorch#548)

This reverts commit a7a24577a65be67ac9ae4dc05452f35d9c49e5d1.

* fix generate for llama3 (pytorch#538)

* fix generate for llama3

* switch more things to C

* remove C++ header

* add delegation visualization instructions (pytorch#551)

* Add dtype runner aoti (pytorch#552)

* add dtype tests for runner-aoti + runner-et

* typo

* add dtype test runner-aoti

* test sdpa with fp16 (pytorch#553)

* test sdpa with fp16

* kv cache fp32

* typo

* update (pytorch#560)

* Only support newest versions of lm-eval (pytorch#556)

Summary:
remove support for lm-eval 0.3 to reduce the options we have

Test Plan:
CI

Reviewers:

Subscribers:

Tasks:

Tags:

* split cpu eval CI by dtype (pytorch#554)

* split cpu eval CI by dtype

* fix

* differentiate names with checks

* keep one name the same as old

* fix

* Removing duplicate HF issue message from README (pytorch#559)

Co-authored-by: Michael Gschwind <[email protected]>

* doc updates (pytorch#567)

* Add VM-safe MPS check

---------

Co-authored-by: Anthony Shoumikhin <[email protected]>
Co-authored-by: metascroy <[email protected]>
Co-authored-by: Nikita Shulga <[email protected]>
Co-authored-by: lucylq <[email protected]>
Co-authored-by: Jerry Zhang <[email protected]>
Co-authored-by: Jack-Khuu <[email protected]>
yanbing-j pushed a commit to yanbing-j/ao that referenced this pull request Dec 9, 2024
* code beautification

* code beautification, move functions together

* make --device fast the default (pytorch#515)

* make --device fast the default

* Update iOS.md (pytorch#517)

* Update iOS.md

* Update iOS.md

* Pip to pip3 (pytorch#504)

* remove macos-12 test

* pip to pip3

* break aoti CI jobs separately (pytorch#500)

* init

* fixes

* more fixes

* fixes

* fix

* fix

* bug fix

* add objcopy update

* suppress int8

* undefined variable

---------

Co-authored-by: Michael Gschwind <[email protected]>

* Support llama3 in chat in run.cpp  (pytorch#486)

* refactor chat runner in preparation for llama3

* add sketch for llama3 prompt template and move to returning tokens

* fix tiktoken

* fixes to chat

* add default llama_ver

* Add tests for quantize json, add cuda device specification and precision to cuda.json (pytorch#519)

* remove code for no KV Cache path (pytorch#527)

* Update ADVANCED-USERS.md (pytorch#529)

Update Advanced Users description to reflect changes in the repo since the description was initially created.

* runner-aoti on cuda (pytorch#531)

* runner-aoti on cuda

* transfer results back to CPU

* transfer results back to CPU

* runner-aoti on cuda

* Update runner_build.md (pytorch#530)

Update description of runner and build process in runner_build.md

* clean up runner code a little (pytorch#532)

* clean up runner code a little

* update

* update

* pull out generate loop in chat

* updates

* edit docs

* typo

* move int8 linear class and function into qops.py (pytorch#534)

* add dtype tests for runner-aoti + runner-et (pytorch#539)

* add dtype tests for runner-aoti + runner-et

* typo

* Quantized embedding (pytorch#536)

* move int8 linear class and function into qops.py

* move Quantized Embedding to qops.py

* Move Linear int4 to qops (pytorch#537)

* move int8 linear class and function into qops.py

* move Quantized Embedding to qops.py

* move int4 linear to qops

* Revert "add dtype tests for runner-aoti + runner-et (pytorch#539)" (pytorch#548)

This reverts commit a7a24577a65be67ac9ae4dc05452f35d9c49e5d1.

* fix generate for llama3 (pytorch#538)

* fix generate for llama3

* switch more things to C

* remove C++ header

* add delegation visualization instructions (pytorch#551)

* Add dtype runner aoti (pytorch#552)

* add dtype tests for runner-aoti + runner-et

* typo

* add dtype test runner-aoti

* test sdpa with fp16 (pytorch#553)

* test sdpa with fp16

* kv cache fp32

* typo

* update (pytorch#560)

* Only support newest versions of lm-eval (pytorch#556)

Summary:
remove support for lm-eval 0.3 to reduce the options we have

Test Plan:
CI

Reviewers:

Subscribers:

Tasks:

Tags:

* split cpu eval CI by dtype (pytorch#554)

* split cpu eval CI by dtype

* fix

* differentiate names with checks

* keep one name the same as old

* fix

* Removing duplicate HF issue message from README (pytorch#559)

Co-authored-by: Michael Gschwind <[email protected]>

* doc updates (pytorch#567)

* Add VM-safe MPS check

---------

Co-authored-by: Anthony Shoumikhin <[email protected]>
Co-authored-by: metascroy <[email protected]>
Co-authored-by: Nikita Shulga <[email protected]>
Co-authored-by: lucylq <[email protected]>
Co-authored-by: Jerry Zhang <[email protected]>
Co-authored-by: Jack-Khuu <[email protected]>

* add unpacking support (pytorch#525)

* add unpacking support

* fix typos and linter

* perform parallel prefill when possible (pytorch#568)

* perform parallel prefill when possible

* typo

* disable hack

* remove print

* remove debug messages which prevent export

* fixes

* stream results in generate.py (pytorch#571)

* remove logging interfering with export

---------

Co-authored-by: Anthony Shoumikhin <[email protected]>
Co-authored-by: metascroy <[email protected]>
Co-authored-by: Nikita Shulga <[email protected]>
Co-authored-by: lucylq <[email protected]>
Co-authored-by: Jerry Zhang <[email protected]>
Co-authored-by: Jack-Khuu <[email protected]>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
CLA Signed This label is managed by the Facebook bot. Authors need to sign the CLA before a PR can be reviewed.
Projects
None yet
Development

Successfully merging this pull request may close these issues.

5 participants