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.

Sign up for GitHub

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

SqueezeLLM Support #1326

Merged
merged 19 commits into from
Oct 22, 2023
Merged

SqueezeLLM Support #1326

Show file tree
Hide file tree
Changes from 2 commits
Commits
Show all changes
19 commits
Select commit Hold shift + click to select a range
f1a6703
First draft of PR
chooper1 Oct 11, 2023
a992f7e
First draft of PR
chooper1 Oct 11, 2023
427edb2
First draft of PR
chooper1 Oct 11, 2023
9e62ca8
Minor Fixes
chooper1 Oct 11, 2023
d4f77dc
Format fixes
Oct 12, 2023
6a7d807
Merge branch 'main' of https://github.com/chooper1/vllm into main
Oct 12, 2023
9494e6c
Format fixes
Oct 12, 2023
5a9546a
Format fixes
Oct 12, 2023
1094483
Merge branch 'main' into chooper1/main
WoosukKwon Oct 22, 2023
0183531
Fix output shapes
WoosukKwon Oct 22, 2023
1064b73
Add NOTE
WoosukKwon Oct 22, 2023
d1c159c
Add namespace & Refactor
WoosukKwon Oct 22, 2023
d77d3ed
Minor
WoosukKwon Oct 22, 2023
c0ef93c
Minor
WoosukKwon Oct 22, 2023
424ed63
Add packed
WoosukKwon Oct 22, 2023
249f0da
Use packed dim for sharding
WoosukKwon Oct 22, 2023
8f0e1f7
Fix Mistral
WoosukKwon Oct 22, 2023
afbe226
Fix
WoosukKwon Oct 22, 2023
4b58d9a
Err msg
WoosukKwon Oct 22, 2023
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
9 changes: 0 additions & 9 deletions csrc/quantization/squeezellm/quant_cuda_kernel.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -99,12 +99,7 @@ __global__ void NUQ4MatMulKernel(
int row = BLOCKHEIGHT4 * blockIdx.x;
int col = BLOCKWIDTH * blockIdx.y + threadIdx.x;

// __shared__ __half blockvec[BLOCKWIDTH];
// blockvec[threadIdx.x] = vec[(row / BLOCKHEIGHT4) * BLOCKWIDTH + threadIdx.x];

__shared__ half2 blockvec[blockwidth2];
// if (threadIdx.x < blockwidth2)
// blockvec[threadIdx.x] = vec[(row / BLOCKHEIGHT4) * blockwidth2 + threadIdx.x];

__shared__ __half deq2[16][BLOCKWIDTH];
int off = threadIdx.x;
Expand All @@ -114,9 +109,6 @@ __global__ void NUQ4MatMulKernel(
deq2[val][off] = lookup_table[lut_index];
}


// __syncthreads();

__half res;
half2 res2;
half2 tmp2;
Expand Down Expand Up @@ -181,7 +173,6 @@ __global__ void NUQ4MatMulKernel(
res3.y = res;
}

// atomicAdd(&mul[b * width + col], res);
atomicAdd(&mul[b * width / 2 + col / 2], res3);
}
}
3 changes: 1 addition & 2 deletions vllm/model_executor/quantization_utils/squeezellm.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -37,8 +37,7 @@ def get_supported_act_dtypes(cls) -> List[torch.dtype]:

@classmethod
def get_min_capability(cls) -> int:
# @ Coleman - TODO check this
return 80
return 70

@classmethod
def get_config_filenames(cls) -> List[str]:
Expand Down