search

mit-han-lab / smoothquant

[ICML 2023] SmoothQuant: Accurate and Efficient Post-Training Quantization for Large Language Models
https://arxiv.org/abs/2211.10438
MIT License
1.27k stars 150 forks source link

[BUG] Int8 inference with torch-int encounter errors #55

Open WelY1 opened 1 year ago

WelY1 commented 1 year ago

Hi! I used export_int8_model.py to smooth, quantize and export INT8 models for opt-1.3b. The model size has changed from 2509MB to 1357MB, which seems that the quantization is successful. But when I evaluated the int8 model, the following error occurred. It seems to be a problem with cutlass. How do I solve this problem? Looking forward to your reply! Thanks!

Here is my environment:

Here is the detailed error message:

RuntimeError                              Traceback (most recent call last)
Cell In[5], line 7
      6 print_model_size(model_int8)
----> 7 acc_smoothquant, lantecy_smoothquant = evaluator.evaluate(model_int8)
      8 print(f'SmoothQuant INT8 accuracy: {acc_smoothquant}, per-sample lantecy: {lantecy_smoothquant:.3f}ms')

File ~/anaconda3/envs/smoothquant/lib/python3.8/site-packages/torch/autograd/grad_mode.py:27, in _DecoratorContextManager.__call__.<locals>.decorate_context(*args, **kwargs)
     24 @functools.wraps(func)
     25 def decorate_context(*args, **kwargs):
     26     with self.clone():
---> 27         return func(*args, **kwargs)

Cell In[2], line 29, in Evaluator.evaluate(self, model)
     27 torch.cuda.synchronize()
     28 start.record()
---> 29 outputs = model(input_ids)
     30 end.record()
     31 torch.cuda.synchronize()

File ~/anaconda3/envs/smoothquant/lib/python3.8/site-packages/torch/nn/modules/module.py:1130, in Module._call_impl(self, *input, **kwargs)
   1126 # If we don't have any hooks, we want to skip the rest of the logic in
   1127 # this function, and just call forward.
...
     43 return y

RuntimeError: CUDA error: device-side assert triggered
CUDA kernel errors might be asynchronously reported at some other API call,so the stacktrace below might be incorrect.
For debugging consider passing CUDA_LAUNCH_BLOCKING=1.
------------------------------------------------------------------------------------------------------------------------------------
tlass::arch::ldsm(cutlass::Array<unsigned int, MatrixCount, true> &, const void *) [with Layout = cutlass::layout::RowMajor; int MatrixCount = 4]: block: [0,9,0], thread: [218,0,0] Assertion `0 && __PRETTY_FUNCTION__` failed.
/home/l50024761/llm/torch-int/submodules/cutlass/include/cutlass/arch/memory_sm75.h:208: void cutlass::arch::ldsm(cutlass::Array<unsigned int, MatrixCount, true> &, const void *) [with Layout = cutlass::layout::RowMajor; int MatrixCount = 4]: block: [0,9,0], thread: [219,0,0] Assertion `0 && __PRETTY_FUNCTION__` failed.
/home/l50024761/llm/torch-int/submodules/cutlass/include/cutlass/arch/memory_sm75.h:208: void cutlass::arch::ldsm(cutlass::Array<unsigned int, MatrixCount, true> &, const void *) [with Layout = cutlass::layout::RowMajor; int MatrixCount = 4]: block: [0,9,0], thread: [220,0,0] Assertion `0 && __PRETTY_FUNCTION__` failed.
/home/l50024761/llm/torch-int/submodules/cutlass/include/cutlass/arch/memory_sm75.h:208: void cutlass::arch::ldsm(cutlass::Array<unsigned int, MatrixCount, true> &, const void *) [with Layout = cutlass::layout::RowMajor; int MatrixCount = 4]: block: [0,9,0], thread: [221,0,0] Assertion `0 && __PRETTY_FUNCTION__` failed.
/home/l50024761/llm/torch-int/submodules/cutlass/include/cutlass/arch/memory_sm75.h:208: void cutlass::arch::ldsm(cutlass::Array<unsigned int, MatrixCount, true> &, const void *) [with Layout = cutlass::layout::RowMajor; int MatrixCount = 4]: block: [0,9,0], thread: [222,0,0] Assertion `0 && __PRETTY_FUNCTION__` failed.
/home/l50024761/llm/torch-int/submodules/cutlass/include/cutlass/arch/memory_sm75.h:208: void cutlass::arch::ldsm(cutlass::Array<unsigned int, MatrixCount, true> &, const void *) [with Layout = cutlass::layout::RowMajor; int MatrixCount = 4]: block: [0,9,0], thread: [223,0,0] Assertion `0 && __PRETTY_FUNCTION__` failed.
/home/l50024761/llm/torch-int/submodules/cutlass/include/cutlass/arch/memory_sm75.h:208: void cutlass::arch::ldsm(cutlass::Array<unsigned int, MatrixCount, true> &, const void *) [with Layout = cutlass::layout::RowMajor; int MatrixCount = 4]: block: [0,9,0], thread: [96,0,0] Assertion `0 && __PRETTY_FUNCTION__` failed.
/home/l50024761/llm/torch-int/submodules/cutlass/include/cutlass/arch/memory_sm75.h:208: void cutlass::arch::ldsm(cutlass::Array<unsigned int, MatrixCount, true> &, const void *) [with Layout = cutlass::layout::RowMajor; int MatrixCount = 4]: block: [0,9,0], thread: [97,0,0] Assertion `0 && __PRETTY_FUNCTION__` failed.
/home/l50024761/llm/torch-int/submodules/cutlass/include/cutlass/arch/memory_sm75.h:208: void cutlass::arch::ldsm(cutlass::Array<unsigned int, MatrixCount, true> &, const void *) [with Layout = cutlass::layout::RowMajor; int MatrixCount = 4]: block: [0,9,0], thread: [98,0,0] Assertion `0 && __PRETTY_FUNCTION__` failed.
/home/l50024761/llm/torch-int/submodules/cutlass/include/cutlass/arch/memory_sm75.h:208: void cutlass::arch::ldsm(cutlass::Array<unsigned int, MatrixCount, true> &, const void *) [with Layout = cutlass::layout::RowMajor; int MatrixCount = 4]: block: [0,9,0], thread: [99,0,0] Assertion `0 && __PRETTY_FUNCTION__` failed.
/home/l50024761/llm/torch-int/submodules/cutlass/include/cutlass/arch/memory_sm75.h:208: void cutlass::arch::ldsm(cutlass::Array<unsigned int, MatrixCount, true> &, const void *) [with Layout = cutlass::layout::RowMajor; int MatrixCount = 4]: block: [0,9,0], thread: [100,0,0] Assertion `0 && __PRETTY_FUNCTION__` failed.
/home/l50024761/llm/torch-int/submodules/cutlass/include/cutlass/arch/memory_sm75.h:208: void cutlass::arch::ldsm(cutlass::Array<unsigned int, MatrixCount, true> &, const void *) [with Layout = cutlass::layout::RowMajor; int MatrixCount = 4]: block: [0,9,0], thread: [101,0,0] Assertion `0 && __PRETTY_FUNCTION__` failed.
/home/l50024761/llm/torch-int/submodules/cutlass/include/cutlass/arch/memory_sm75.h:208: void cutlass::arch::ldsm(cutlass::Array<unsigned int, MatrixCount, true> &, const void *) [with Layout = cutlass::layout::RowMajor; int MatrixCount = 4]: block: [0,9,0], thread: [102,0,0] Assertion `0 && __PRETTY_FUNCTION__` failed.
/home/l50024761/llm/torch-int/submodules/cutlass/include/cutlass/arch/memory_sm75.h:208: void cutlass::arch::ldsm(cutlass::Array<unsigned int, MatrixCount, true> &, const void *) [with Layout = cutlass::layout::RowMajor; int MatrixCount = 4]: block: [0,9,0], thread: [103,0,0] Assertion `0 && __PRETTY_FUNCTION__` failed.
/home/l50024761/llm/torch-int/submodules/cutlass/include/cutlass/arch/memory_sm75.h:208: void cutlass::arch::ldsm(cutlass::Array<unsigned int, MatrixCount, true> &, const void *) [with Layout = cutlass::layout::RowMajor; int MatrixCount = 4]: block: [0,9,0], thread: [104,0,0] Assertion `0 && __PRETTY_FUNCTION__` failed.
/home/l50024761/llm/torch-int/submodules/cutlass/include/cutlass/arch/memory_sm75.h:208: void cutlass::arch::ldsm(cutlass::Array<unsigned int, MatrixCount, true> &, const void *) [with Layout = cutlass::layout::RowMajor; int MatrixCount = 4]: block: [0,9,0], thread: [105,0,0] Assertion `0 && __PRETTY_FUNCTION__` failed.
/home/l50024761/llm/torch-int/submodules/cutlass/include/cutlass/arch/memory_sm75.h:208: void cutlass::arch::ldsm(cutlass::Array<unsigned int, MatrixCount, true> &, const void *) [with Layout = cutlass::layout::RowMajor; int MatrixCount = 4]: block: [0,9,0], thread: [106,0,0] Assertion `0 && __PRETTY_FUNCTION__` failed.
/home/l50024761/llm/torch-int/submodules/cutlass/include/cutlass/arch/memory_sm75.h:208: void cutlass::arch::ldsm(cutlass::Array<unsigned int, MatrixCount, true> &, const void *) [with Layout = cutlass::layout::RowMajor; int MatrixCount = 4]: block: [0,9,0], thread: [107,0,0] Assertion `0 && __PRETTY_FUNCTION__` failed.
/home/l50024761/llm/torch-int/submodules/cutlass/include/cutlass/arch/memory_sm75.h:208: void cutlass::arch::ldsm(cutlass::Array<unsigned int, MatrixCount, true> &, const void *) [with Layout = cutlass::layout::RowMajor; int MatrixCount = 4]: block: [0,9,0], thread: [108,0,0] Assertion `0 && __PRETTY_FUNCTION__` failed.
/home/l50024761/llm/torch-int/submodules/cutlass/include/cutlass/arch/memory_sm75.h:208: void cutlass::arch::ldsm(cutlass::Array<unsigned int, MatrixCount, true> &, const void *) [with Layout = cutlass::layout::RowMajor; int MatrixCount = 4]: block: [0,9,0], thread: [109,0,0] Assertion `0 && __PRETTY_FUNCTION__` failed.
/home/l50024761/llm/torch-int/submodules/cutlass/include/cutlass/arch/memory_sm75.h:208: void cutlass::arch::ldsm(cutlass::Array<unsigned int, MatrixCount, true> &, const void *) [with Layout = cutlass::layout::RowMajor; int MatrixCount = 4]: block: [0,9,0], thread: [110,0,0] Assertion `0 && __PRETTY_FUNCTION__` failed.
/home/l50024761/llm/torch-int/submodules/cutlass/include/cutlass/arch/memory_sm75.h:208: void cutlass::arch::ldsm(cutlass::Array<unsigned int, MatrixCount, true> &, const void *) [with Layout = cutlass::layout::RowMajor; int MatrixCount = 4]: block: [0,9,0], thread: [111,0,0] Assertion `0 && __PRETTY_FUNCTION__` failed.
/home/l50024761/llm/torch-int/submodules/cutlass/include/cutlass/arch/memory_sm75.h:208: void cutlass::arch::ldsm(cutlass::Array<unsigned int, MatrixCount, true> &, const void *) [with Layout = cutlass::layout::RowMajor; int MatrixCount = 4]: block: [0,9,0], thread: [112,0,0] Assertion `0 && __PRETTY_FUNCTION__` failed.
/home/l50024761/llm/torch-int/submodules/cutlass/include/cutlass/arch/memory_sm75.h:208: void cutlass::arch::ldsm(cutlass::Array<unsigned int, MatrixCount, true> &, const void *) [with Layout = cutlass::layout::RowMajor; int MatrixCount = 4]: block: [0,9,0], thread: [113,0,0] Assertion `0 && __PRETTY_FUNCTION__` failed.
/home/l50024761/llm/torch-int/submodules/cutlass/include/cutlass/arch/memory_sm75.h:208: void cutlass::arch::ldsm(cutlass::Array<unsigned int, MatrixCount, true> &, const void *) [with Layout = cutlass::layout::RowMajor; int MatrixCount = 4]: block: [0,9,0], thread: [114,0,0] Assertion `0 && __PRETTY_FUNCTION__` failed.
/home/l50024761/llm/torch-int/submodules/cutlass/include/cutlass/arch/memory_sm75.h:208: void cutlass::arch::ldsm(cutlass::Array<unsigned int, MatrixCount, true> &, const void *) [with Layout = cutlass::layout::RowMajor; int MatrixCount = 4]: block: [0,9,0], thread: [115,0,0] Assertion `0 && __PRETTY_FUNCTION__` failed.
/home/l50024761/llm/torch-int/submodules/cutlass/include/cutlass/arch/memory_sm75.h:208: void cutlass::arch::ldsm(cutlass::Array<unsigned int, MatrixCount, true> &, const void *) [with Layout = cutlass::layout::RowMajor; int MatrixCount = 4]: block: [0,9,0], thread: [116,0,0] Assertion `0 && __PRETTY_FUNCTION__` failed.
/home/l50024761/llm/torch-int/submodules/cutlass/include/cutlass/arch/memory_sm75.h:208: void cutlass::arch::ldsm(cutlass::Array<unsigned int, MatrixCount, true> &, const void *) [with Layout = cutlass::layout::RowMajor; int MatrixCount = 4]: block: [0,9,0], thread: [117,0,0] Assertion `0 && __PRETTY_FUNCTION__` failed.
/home/l50024761/llm/torch-int/submodules/cutlass/include/cutlass/arch/memory_sm75.h:208: void cutlass::arch::ldsm(cutlass::Array<unsigned int, MatrixCount, true> &, const void *) [with Layout = cutlass::layout::RowMajor; int MatrixCount = 4]: block: [0,9,0], thread: [118,0,0] Assertion `0 && __PRETTY_FUNCTION__` failed.
/home/l50024761/llm/torch-int/submodules/cutlass/include/cutlass/arch/memory_sm75.h:208: void cutlass::arch::ldsm(cutlass::Array<unsigned int, MatrixCount, true> &, const void *) [with Layout = cutlass::layout::RowMajor; int MatrixCount = 4]: block: [0,9,0], thread: [119,0,0] Assertion `0 && __PRETTY_FUNCTION__` failed.
/home/l50024761/llm/torch-int/submodules/cutlass/include/cutlass/arch/memory_sm75.h:208: void cutlass::arch::ldsm(cutlass::Array<unsigned int, MatrixCount, true> &, const void *) [with Layout = cutlass::layout::RowMajor; int MatrixCount = 4]: block: [0,9,0], thread: [120,0,0] Assertion `0 && __PRETTY_FUNCTION__` failed.
/home/l50024761/llm/torch-int/submodules/cutlass/include/cutlass/arch/memory_sm75.h:208: void cutlass::arch::ldsm(cutlass::Array<unsigned int, MatrixCount, true> &, const void *) [with Layout = cutlass::layout::RowMajor; int MatrixCount = 4]: block: [0,9,0], thread: [121,0,0] Assertion `0 && __PRETTY_FUNCTION__` failed.
/home/l50024761/llm/torch-int/submodules/cutlass/include/cutlass/arch/memory_sm75.h:208: void cutlass::arch::ldsm(cutlass::Array<unsigned int, MatrixCount, true> &, const void *) [with Layout = cutlass::layout::RowMajor; int MatrixCount = 4]: block: [0,9,0], thread: [122,0,0] Assertion `0 && __PRETTY_FUNCTION__` failed.
/home/l50024761/llm/torch-int/submodules/cutlass/include/cutlass/arch/memory_sm75.h:208: void cutlass::arch::ldsm(cutlass::Array<unsigned int, MatrixCount, true> &, const void *) [with Layout = cutlass::layout::RowMajor; int MatrixCount = 4]: block: [0,9,0], thread: [123,0,0] Assertion `0 && __PRETTY_FUNCTION__` failed.
/home/l50024761/llm/torch-int/submodules/cutlass/include/cutlass/arch/memory_sm75.h:208: void cutlass::arch::ldsm(cutlass::Array<unsigned int, MatrixCount, true> &, const void *) [with Layout = cutlass::layout::RowMajor; int MatrixCount = 4]: block: [0,9,0], thread: [124,0,0] Assertion `0 && __PRETTY_FUNCTION__` failed.
/home/l50024761/llm/torch-int/submodules/cutlass/include/cutlass/arch/memory_sm75.h:208: void cutlass::arch::ldsm(cutlass::Array<unsigned int, MatrixCount, true> &, const void *) [with Layout = cutlass::layout::RowMajor; int MatrixCount = 4]: block: [0,9,0], thread: [125,0,0] Assertion `0 && __PRETTY_FUNCTION__` failed.
/home/l50024761/llm/torch-int/submodules/cutlass/include/cutlass/arch/memory_sm75.h:208: void cutlass::arch::ldsm(cutlass::Array<unsigned int, MatrixCount, true> &, const void *) [with Layout = cutlass::layout::RowMajor; int MatrixCount = 4]: block: [0,9,0], thread: [126,0,0] Assertion `0 && __PRETTY_FUNCTION__` failed.
/home/l50024761/llm/torch-int/submodules/cutlass/include/cutlass/arch/memory_sm75.h:208: void cutlass::arch::ldsm(cutlass::Array<unsigned int, MatrixCount, true> &, const void *) [with Layout = cutlass::layout::RowMajor; int MatrixCount = 4]: block: [0,9,0], thread: [127,0,0] Assertion `0 && __PRETTY_FUNCTION__` failed.