[Kernel][RFC] Refactor the punica kernel based on Triton

FILL IN THE PR DESCRIPTION HERE
Motivation
LoRA is highly favored within the vLLM community, there are numerous LoRA-related issues and pull requests.  Thanks for  @Yard1 great work, we  have supported the Mutil-LoRA . However, there are several issues with the punica kernel (bgmv):

Performance drop refer to :https://github.com/vllm-project/vllm/issues/2829

We need to modify the punica config to support the new dimension(such as  https://github.com/vllm-project/vllm/pull/4007, https://github.com/vllm-project/vllm/pull/4053 , https://github.com/vllm-project/vllm/pull/4063/)

The punica kernel cannot support certain dimensions, such as 3424 (for example, in ChatGLM3, the ffn_hidden_size is 13696, and when using tp=4, the ffn_hidden_size becomes 3424), because 3424 cannot be evenly divided by 64, which will cause a static_assert failure. refer to :https://github.com/vllm-project/vllm/issues/3793 https://github.com/vllm-project/vllm/issues/4708

Only support GPU which is compute capability >= 8.0

Other issues, such as increased compile time and  the release wheel size( I'm not entirely sure if I understand correctly), as well as the potential need to support long context, may result in the second dimension of the kernel grid exceeding 65535, leading to kernel launch failure.

How to resolve
As descibed in title，I plan to resolve these issues by refactoring  the punica kernel based on Triton，the main reason is Triton provides a hardware agnostic way of programming and targeting GPUs, currently supporting both NVIDIA and AMD.  I will implement  punica triton kernel , and:

For shrink kernel, I will implement it based on GroupedGEMM and SPLIT-K
For expand kernel,  I will implement it based on GroupedGEMM

Work plan

[x]  Implement SGMV kernel(shrink+expand)
[x]  Replace punica bgmv kernel with sgmv trtion kenel
[x]  add unit test for sgmv trtion kenel

If my method is approved, I will spend one week to complete this PR,  and  tuning of the Triton kernel will be completed in the subsequent PR. Currently, I want to focus on replacing the punica bgmv kernel with the sgmv Triton kernel.
Please leave any feedback, any feedback will be appreciated.
Summary [06/27]

I have completed the development of the Triton kernel and achieved SGMV (shrink/expand) and BGMV (shrink/expand)
I have completed the  integration, and unit testing of the Triton kernel according to my approach. TP（2/4/8）/CUDA Graph/Quantization/Fully sharded LoRA/NVIDIA T4, and so forth, all yield correct results. IMHO, the  above issues  were addressed or alleviated(performance drop). the details  as follow table.
This is just a start,  there's more work to be done , such as further improving kernel performance and auto-tuning configurations,I believe these can be considered for future PR
Regardless of the final decision made by the VLLM team, I really enjoy the process of addressing these issues

cc @zhuohan123 @Yard1 @WoosukKwon  @simon-mo  @mgoin @FurtherAI

We battle it out. We can take then best of both or just select one if they are close. I do have a good portion of the integration with vLLM written, but untested, so it would be good if that still applied in the end.
@jeejeelee First, can you tell me about any limitations or assumptions? Lora weight format (paged or padded), do the kernels operate on each lora at their rank or at a max rank, do you need a minimum group size per lora (tl.dot requires shapes of 16 or greater so do you pad if there aren't enough tokens for a lora?), do you have any benchmarking, do you have any tests for correctness or anything else you think is important?

@FurtherAI  Let's cheers to wonderful similarities, I'm glad to see like-minded people  to address these issues.  Here is my response:

can you tell me about any limitations or assumptions?
My version is designed to be compatible with the current vLLM's LoRA format. If I am understand correctly,  vLLM's LoRA format is padded,   if paged format is the memory for different LoRA of the same layer is not contiguous,  my version also compatible with the paged format  by modifying some codes.

do the kernels operate on each lora at their rank or at a max rank, do you need a minimum group size per lora (tl.dot requires shapes of 16 or greater so do you pad if there aren't enough tokens for a lora?)
My version operates at maximum rank, taking into consideration vLLM's LoRA format，my version should also be compatible with operating at each lora rank.  We need ensure that the smallest BLOCK_N is 16

do you have any benchmarking
Not yet. I have not conducted a systematic benchmark, I have only compared  with the current vLLM's sgmv using NSYS,  and the performance improvement is very obvious. I will complete the relevant benchmark as soon as possible.

do you have any tests for correctness
Yes, I have not only tested with PyTorch, but I have also tested with vLLM's bgmv, for example. the above issue kernel grid exceeding 65535 was found when testing with vLLM's bgmv.

anything else you think is important?
IMHO,  Resloving these above issues is most important as soon as possilble.

I really enjoy our conversation, which deepened my understanding of multi-lora. Thanks
Cheers

@FurtherAI  I apologize for the late reply due to the weekend. I provide some code in temp_test.py

No worries, thanks for the code. I'll try to get back to you soon.

This is great! The punica kernels also take a very long time to build and take up space in our wheels. Moving them to triton can make the building process faster and reduce the wheel size.

Currently this imp still has two kernel dealing with shrink and expand separately. I wander whether we could merge them into one? So that triton could do the pipeline autotune for this fused kernel?
And take one step further, whether we could have kernel fuse x@W+x@A@B, so that could make load inference at decoding stage cost reduced?

Currently this imp still has two kernel dealing with shrink and expand separately. I wonder whether we could merge them into one? So that triton could do the pipeline autotune for this fused kernel?
And take one step further, whether we could have kernel fuse x@W+x@A@B, so that could make load inference at decoding stage cost reduced?

Hi, thanks for your attention and suggest.
Of course, we can fuse these two back-to-back GEMM kernels as done in b2b_grouped_gemm.Actually， implementing this in a single kernel is my first idea, but I have not yet figured out how to map split-k, batch_size, and M, N into grid. Therefore, I am currently implementing shrink and expand separately. In the future, I will continue to explore this. For now, my focus is on replacing the current bgmv and addressing the issues mentioned above.
Regarding "pipeline autotune for this fused kernel," I'm not quite sure what this means—did you mean to refer to the num_stage? If I understand correctly, I think if your GPU sm>8.0， trtion will invoking  pipeline optimization，regardless of whether fusion is involved.
For fuse `` x@W+x@A@B ,if we only have one LoRA, we can merge the LoRA's weight into W first and then only need to compute x@W. However, since we need to support Multi-LoRA, I think fusing everything into one kernel might not be the best choice. we also need to consider factors such as register pressure.

@jeejeelee So it looks like our kernels accomplish two partly different goals.
Yours can function as a drop in replacement for the current Punica kernels. I have done some benchmarking (which I'll attach) and assuming I understand how to bench yours correctly and it is doing what I expect, then I think yours have pretty comparable speed to Punica across the board.
Mine has two major differences, loras in a paged memory format and computing at the rank of each lora. Together these allow for more loras at a time and assuming loras are randomly distributed in rank, they get pretty fast and it makes up for the cost of loading the extra indexing information for having them in a page. Your strategy of looping along the rank dimension didn't give any speedup to mine, so I expect any differences are loading the indexing and the different format.
Based on the benchmarks and assuming lora ranks are distributed, these kernels don't offer a significant speed advantage over mine. Though they would probably be faster than mine if you add the ability to compute at the actual rank (lacking the paged loras though). It would be reasonable to use yours as a drop in replacement for Punica until I push mine or keep both and use mine if the user wants more active loras. Or it is reasonable to just wait for the kernels I've developed to be fully integrated since these kernels will likely be a marginal improvement.
Up to the vLLM team what they want to do.
bench_triton_kernels.txt

@jeejeelee So it looks like our kernels accomplish two partly different goals.
Yours can function as a drop in replacement for the current Punica kernels. I have done some benchmarking (which I'll attach) and assuming I understand how to bench yours correctly and it is doing what I expect, then I think yours have pretty comparable speed to Punica across the board.
Mine has two major differences, loras in a paged memory format and computing at the rank of each lora. Together these allow for more loras at a time and assuming loras are randomly distributed in rank, they get pretty fast and it makes up for the cost of loading the extra indexing information for having them in a page. Your strategy of looping along the rank dimension didn't give any speedup to mine, so I expect any differences are loading the indexing and the different format.
Based on the benchmarks and assuming lora ranks are distributed, these kernels don't offer a significant speed advantage over mine. Though they would probably be faster than mine if you add the ability to compute at the actual rank (lacking the paged loras though). It would be reasonable to use yours as a drop in replacement for Punica until I push mine or keep both and use mine if the user wants more active loras. Or it is reasonable to just wait for the kernels I've developed to be fully integrated since these kernels will likely be a marginal improvement.
Up to the vLLM team what they want to do. I think the expand with repeats 8 is an anomaly here, it matches the shrink kernel on other machines. bench_triton_kernels.txt

Thanks for your work benckmark. It seems that you mainly want to support S-LoRA, so you implemented these kernels.
Currently, I am following the above work plan and attempting to replace bgmv with my operator (progress 80%). Additionally, I have also added bgmv-related triton kernel(these operators have not been fully tested yet). Regardless, I will complete my work plan.
Cheers

Upload  the test result Using yi-34B:

Outstanding issues:

Full-shard LoRA support
Resolve the other LoRA's tests

The markers of the first integration method: https://github.com/jeejeelee/vllm/tree/00e007695c8cfa466f53fa74a0a601aa42a10cd7

Will you merge this soon?

Will you merge this soon?

Thank you for your attention. I'm not sure if we can merge yet, but I have completed most of the development work. You can pull my branch and manually build it for testing.

@simon-mo  Could you please check why the CI test failed? I have actually completed the unit tests locally and would like to see if there are any omissions.

To speed up the CI queue, I've cancelled the distributed tests for the latest CI run in this PR since they won't pass anyway until #5905 has been merged. Now that it has been merged, please merge main into your branch so that the CI can pass once again.

@Yard1  Thanks for your review, I will fix these asap

@jeejeelee  Hello, can your version solve the issue of “Due to limitations of the custom LoRA CUDA kernel, max_num_batched_tokens must be <= 65528 when LoRA is enabled.”? Thank you

max_num_batched_tokens must be <= 65528

I haven't tested it yet, but I believe my version does not have these limitations. You can test it yourself,Looking forward to your feedback~

/ready

Libentry test
The current version of Triton used in vLLM is 2.3.1, while the official Triton version is 3.0.0. Therefore, we tested the usage of libentry in these two versions.
Performance
triton=2.3.1

sgmv kernel test

Phi2-3B and Llama-7B

triton=3.0.0

sgmv kernel test

Phi2-3B and Llama-7B

Conclusion

For Triton 2.3.1, libentry can bring significant performance gains for small inputs, while having no effect on performance for large inputs.
When using Triton 3.0.0, adding libentry decoratorbrought almost no performance gain. This is because Triton 3.0.0 has already updated the relevant code. Please refer to: https://github.com/triton-lang/triton/pull/3503.

Test the impact
We  also added libentry  decorator to the trion kernels in vLLM and tested whether it affected the correctness of these kernels .

triton version
kernel type
test result

2.3.1
sgmv and bgmv
passed

2.3.1
fuse_moe
passed

2.3.1
prefix_prefill
passed

2.3.1
sampler
passed

3.0.0
sgmv and bgmv
passed

3.0.0
fuse_moe
passed

3.0.0
prefix_prefill
passed

3.0.0
sampler
failed

PS: The reason for the sample triton kernel failure was a compilation error, which is unrelated to libentry. This issue may need attention when updating to triton 3.0.0 in the future.
Acknowledgments
Libentry decorator was copied fromFlagGems , thanks for their great work.

@jeejeelee thanks, this is looking good. Can you add a comment to libentry code that it can be removed with triton 3.0.0?

@jeejeelee thanks, this is looking good. Can you add a comment to libentry code that it can be removed with triton 3.0.0?

Thank you for your review. I have added the comment. Additionally, the CI keeps reporting an Out of Memory (OOM) error when running test_quant_model.py. How should I resolve this issue?

@jeejeelee Let's try to merge master into your branch first and see if that resolves the CI issue

@Yard1  I have set gpu_memory_utilization to 0.2 but it still throws OOM error. I'm not very familiar with the CI test.   Could you please take a look?

An incorrect result on v100-32G
# install vllm
git clone -b refactor-punica-kernel https://github.com/jeejeelee/vllm.git
cd vllm
pip install -e .

# start vllm serve
vllm serve /home/models/Qwen1.5-14B-Chat \
    --enable-lora \
    --lora-modules sql-lora=/home/wxp/qwen1.5-14-ccb-time-lora/ \
	--dtype half \
	--tensor-parallel-size 2

# input

{
	"model": "sql-lora",
	"messages": [
		{
			"role": "user",
			"content": "你擅长从背景知识提取关键信息，已知背景知识如下：\n手机：15101234567,徐宇轩准备巡视机房；身份证：210606198206152512,下周一早上5点48分.\n你需要通过背景信息提取出name、phone、identityCard、startTime，提取的startTime需要转换为java代码。\n输出：\n"
		}
	],
	"temperature": 0.001,
	"stream": false,
	"max_tokens": 600
}

# expect:
{
    "id": "chataa41c346-4984-11ef-9930-a6690e7006bb",
    "object": "chat.completion",
    "created":<PHONE_NUMBER>,
    "model": "sql-lora",
    "choices": [
        {
            "index": 0,
            "message": {
                "role": "assistant",
                "content": "{\"name\": \"徐宇轩\", \"phone\": \"15101234567\", \"identityCard\": \"210606198206152512\", \"startTime\": \"LocalDateTime.now().plusWeeks(1).with(DayOfWeek.MONDAY).toLocalDate().atTime(5,48);\"}"
            },
            "finish_reason": "stop"
        }
    ],
    "usage": {
        "prompt_tokens": 112,
        "completion_tokens": 83,
        "total_tokens": 195
    }
}

# but get a error result:
{
    "id": "chat-4476f0fe838b4c9a832f9e050e24ae02",
    "object": "chat.completion",
    "created":<PHONE_NUMBER>,
    "model": "/home/models/Qwen1.5-14B-Chat",
    "choices": [
        {
            "index": 0,
            "message": {
                "role": "assistant",
                "content": "name: 徐宇轩\nphone:<PHONE_NUMBER>7\nidentityCard:<PHONE_NUMBER>06152512\nstartTime: \"LocalDateTime.now().plusWeeks(1).with(DayOfWeek.MONDAY).toLocalDate().atTime(5,48);\"\n\n在Java代码中，这将表示为：\n\n```java\nLocalDateTime startTime = LocalDateTime.now().plusWeeks(1).with(DayOfWeek.MONDAY).toLocalDate().atTime(5,48);\n```",
                "tool_calls": []
            },
            "logprobs": null,
            "finish_reason": "stop",
            "stop_reason": null
        }
    ],
    "usage": {
        "prompt_tokens": 112,
        "total_tokens": 237,
        "completion_tokens": 125
    }
}

@WangxuPThank you for the feedback. How is the expected result generated? Also, did you perform testing based on the vllm on the main branch?

@WangxuP Thank you for the feedback. How is the expected result generated? Also, did you perform testing based on the vllm on the main branch?

I pulled the following branch code
git clone -b refactor-punica-kernel https://github.com/jeejeelee/vllm.git

After upgrading triton==3.0.0, the issues I mentioned earlier no longer exist.
so cool!

@jeejeelee I've been testing before merging and I'm having the following problem.
Am I missing something?
model : Qwen2-72B-Instruct
LoRA : r 32, a 32
gpu node : A100 80g x 8

# install vllm
git clone -b refactor-punica-kernel https://github.com/jeejeelee/vllm.git
cd vllm
pip install -e .

# start vllm serve
vllm serve /path/to/Qwen2-72B-Instruct --served-model-name Qwen2-72B-Instruct \
--tensor-parallel-size 8 \
--enable-lora \
--lora-modules review-agent=/path/to/LoRA

The server will run normally with the Multi LoRA option.
However, the following error occurs when requesting a inference.
INFO 07-25 09:21:49 logger.py:36] Received request chat-5b51ad132e784242a0c7753f76ba9780: prompt: '<|im_start|>system\nYou are an AI programming assistant.<|im_end|>\n<|im_start|>user\nhi<|im_end|>\n<|im_start|>assistant\n', params: SamplingParams(n=1, best_of=1, presence_penalty=0.0, frequency_penalty=0.0, repetition_penalty=1.0, temperature=0.2, top_p=1.0, top_k=-1, min_p=0.0, seed=None, use_beam_search=False, length_penalty=1.0, early_stopping=False, stop=[], stop_token_ids=[], include_stop_str_in_output=False, ignore_eos=False, max_tokens=1024, min_tokens=0, logprobs=None, prompt_logprobs=None, skip_special_tokens=True, spaces_between_special_tokens=True, truncate_prompt_tokens=None), prompt_token_ids: [151644, 8948, 198, 2610, 525, 458, 15235, 15473, 17847, 13, 151645, 198, 151644, 872, 198, 6023, 151645, 198, 151644, 77091, 198], lora_request: None, prompt_adapter_request: None.
INFO:     <IP_ADDRESS>:56796 - "POST /v1/chat/completions HTTP/1.1" 200 OK
ERROR:    Exception in ASGI application
Traceback (most recent call last):
  File "/home/bc-user/anaconda3/envs/vllm-source/lib/python3.11/site-packages/starlette/responses.py", line 265, in __call__
    await wrap(partial(self.listen_for_disconnect, receive))
  File "/home/bc-user/anaconda3/envs/vllm-source/lib/python3.11/site-packages/starlette/responses.py", line 261, in wrap
    await func()
  File "/home/bc-user/anaconda3/envs/vllm-source/lib/python3.11/site-packages/starlette/responses.py", line 238, in listen_for_disconnect
    message = await receive()
              ^^^^^^^^^^^^^^^
  File "/home/bc-user/anaconda3/envs/vllm-source/lib/python3.11/site-packages/uvicorn/protocols/http/httptools_impl.py", line 553, in receive
    await self.message_event.wait()
  File "/home/bc-user/anaconda3/envs/vllm-source/lib/python3.11/asyncio/locks.py", line 213, in wait
    await fut
asyncio.exceptions.CancelledError: Cancelled by cancel scope 7fc35744c850

During handling of the above exception, another exception occurred:

  + Exception Group Traceback (most recent call last):
  |   File "/home/bc-user/anaconda3/envs/vllm-source/lib/python3.11/site-packages/uvicorn/protocols/http/httptools_impl.py", line 399, in run_asgi
  |     result = await app(  # type: ignore[func-returns-value]
  |              ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  |   File "/home/bc-user/anaconda3/envs/vllm-source/lib/python3.11/site-packages/uvicorn/middleware/proxy_headers.py", line 70, in __call__
  |     return await self.app(scope, receive, send)
  |            ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
  |   File "/home/bc-user/anaconda3/envs/vllm-source/lib/python3.11/site-packages/fastapi/applications.py", line 1054, in __call__
  |     await super().__call__(scope, receive, send)
  |   File "/home/bc-user/anaconda3/envs/vllm-source/lib/python3.11/site-packages/starlette/applications.py", line 123, in __call__
  |     await self.middleware_stack(scope, receive, send)
  |   File "/home/bc-user/anaconda3/envs/vllm-source/lib/python3.11/site-packages/starlette/middleware/errors.py", line 186, in __call__
  |     raise exc
  |   File "/home/bc-user/anaconda3/envs/vllm-source/lib/python3.11/site-packages/starlette/middleware/errors.py", line 164, in __call__
  |     await self.app(scope, receive, _send)
  |   File "/home/bc-user/anaconda3/envs/vllm-source/lib/python3.11/site-packages/starlette/middleware/cors.py", line 85, in __call__
  |     await self.app(scope, receive, send)
  |   File "/home/bc-user/anaconda3/envs/vllm-source/lib/python3.11/site-packages/starlette/middleware/exceptions.py", line 65, in __call__
  |     await wrap_app_handling_exceptions(self.app, conn)(scope, receive, send)
  |   File "/home/bc-user/anaconda3/envs/vllm-source/lib/python3.11/site-packages/starlette/_exception_handler.py", line 64, in wrapped_app
  |     raise exc
  |   File "/home/bc-user/anaconda3/envs/vllm-source/lib/python3.11/site-packages/starlette/_exception_handler.py", line 53, in wrapped_app
  |     await app(scope, receive, sender)
  |   File "/home/bc-user/anaconda3/envs/vllm-source/lib/python3.11/site-packages/starlette/routing.py", line 756, in __call__
  |     await self.middleware_stack(scope, receive, send)
  |   File "/home/bc-user/anaconda3/envs/vllm-source/lib/python3.11/site-packages/starlette/routing.py", line 776, in app
  |     await route.handle(scope, receive, send)
  |   File "/home/bc-user/anaconda3/envs/vllm-source/lib/python3.11/site-packages/starlette/routing.py", line 297, in handle
  |     await self.app(scope, receive, send)
  |   File "/home/bc-user/anaconda3/envs/vllm-source/lib/python3.11/site-packages/starlette/routing.py", line 77, in app
  |     await wrap_app_handling_exceptions(app, request)(scope, receive, send)
  |   File "/home/bc-user/anaconda3/envs/vllm-source/lib/python3.11/site-packages/starlette/_exception_handler.py", line 64, in wrapped_app
  |     raise exc
  |   File "/home/bc-user/anaconda3/envs/vllm-source/lib/python3.11/site-packages/starlette/_exception_handler.py", line 53, in wrapped_app
  |     await app(scope, receive, sender)
  |   File "/home/bc-user/anaconda3/envs/vllm-source/lib/python3.11/site-packages/starlette/routing.py", line 75, in app
  |     await response(scope, receive, send)
  |   File "/home/bc-user/anaconda3/envs/vllm-source/lib/python3.11/site-packages/starlette/responses.py", line 258, in __call__
  |     async with anyio.create_task_group() as task_group:
  |   File "/home/bc-user/anaconda3/envs/vllm-source/lib/python3.11/site-packages/anyio/_backends/_asyncio.py", line 680, in __aexit__
  |     raise BaseExceptionGroup(
  | ExceptionGroup: unhandled errors in a TaskGroup (1 sub-exception)
  +-+---------------- 1 ----------------
    | Traceback (most recent call last):
    |   File "/home/bc-user/anaconda3/envs/vllm-source/lib/python3.11/site-packages/starlette/responses.py", line 261, in wrap
    |     await func()
    |   File "/home/bc-user/anaconda3/envs/vllm-source/lib/python3.11/site-packages/starlette/responses.py", line 250, in stream_response
    |     async for chunk in self.body_iterator:
    |   File "/home/bc-user/jeejeelee-vllm/vllm/entrypoints/openai/serving_chat.py", line 228, in chat_completion_stream_generator
    |     async for res in result_generator:
    |   File "/home/bc-user/jeejeelee-vllm/vllm/engine/async_llm_engine.py", line 772, in generate
    |     async for output in self._process_request(
    |   File "/home/bc-user/jeejeelee-vllm/vllm/engine/async_llm_engine.py", line 888, in _process_request
    |     raise e
    |   File "/home/bc-user/jeejeelee-vllm/vllm/engine/async_llm_engine.py", line 884, in _process_request
    |     async for request_output in stream:
    |   File "/home/bc-user/jeejeelee-vllm/vllm/engine/async_llm_engine.py", line 93, in __anext__
    |     raise result
    |   File "/home/bc-user/jeejeelee-vllm/vllm/engine/async_llm_engine.py", line 46, in _log_task_completion
    |     return_value = task.result()
    |                    ^^^^^^^^^^^^^
    |   File "/home/bc-user/jeejeelee-vllm/vllm/engine/async_llm_engine.py", line 637, in run_engine_loop
    |     result = task.result()
    |              ^^^^^^^^^^^^^
    |   File "/home/bc-user/jeejeelee-vllm/vllm/engine/async_llm_engine.py", line 580, in engine_step
    |     request_outputs = await self.engine.step_async(virtual_engine)
    |                       ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
    |   File "/home/bc-user/jeejeelee-vllm/vllm/engine/async_llm_engine.py", line 253, in step_async
    |     output = await self.model_executor.execute_model_async(
    |              ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
    |   File "/home/bc-user/jeejeelee-vllm/vllm/executor/distributed_gpu_executor.py", line 175, in execute_model_async
    |     return await self._driver_execute_model_async(execute_model_req)
    |            ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
    |   File "/home/bc-user/jeejeelee-vllm/vllm/executor/multiproc_gpu_executor.py", line 210, in _driver_execute_model_async
    |     return await self.driver_exec_model(execute_model_req)
    |            ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
    |   File "/home/bc-user/anaconda3/envs/vllm-source/lib/python3.11/concurrent/futures/thread.py", line 58, in run
    |     result = self.fn(*self.args, **self.kwargs)
    |              ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
    |   File "/home/bc-user/jeejeelee-vllm/vllm/worker/worker_base.py", line 272, in execute_model
    |     output = self.model_runner.execute_model(
    |              ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
    |   File "/home/bc-user/anaconda3/envs/vllm-source/lib/python3.11/site-packages/torch/utils/_contextlib.py", line 115, in decorate_context
    |     return func(*args, **kwargs)
    |            ^^^^^^^^^^^^^^^^^^^^^
    |   File "/home/bc-user/jeejeelee-vllm/vllm/worker/model_runner.py", line 1314, in execute_model
    |     hidden_or_intermediate_states = model_executable(
    |                                     ^^^^^^^^^^^^^^^^^
    |   File "/home/bc-user/anaconda3/envs/vllm-source/lib/python3.11/site-packages/torch/nn/modules/module.py", line 1532, in _wrapped_call_impl
    |     return self._call_impl(*args, **kwargs)
    |            ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
    |   File "/home/bc-user/anaconda3/envs/vllm-source/lib/python3.11/site-packages/torch/nn/modules/module.py", line 1541, in _call_impl
    |     return forward_call(*args, **kwargs)
    |            ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
    |   File "/home/bc-user/jeejeelee-vllm/vllm/model_executor/models/qwen2.py", line 336, in forward
    |     hidden_states = self.model(input_ids, positions, kv_caches,
    |                     ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
    |   File "/home/bc-user/anaconda3/envs/vllm-source/lib/python3.11/site-packages/torch/nn/modules/module.py", line 1532, in _wrapped_call_impl
    |     return self._call_impl(*args, **kwargs)
    |            ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
    |   File "/home/bc-user/anaconda3/envs/vllm-source/lib/python3.11/site-packages/torch/nn/modules/module.py", line 1541, in _call_impl
    |     return forward_call(*args, **kwargs)
    |            ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
    |   File "/home/bc-user/jeejeelee-vllm/vllm/model_executor/models/qwen2.py", line 257, in forward
    |     hidden_states, residual = layer(
    |                               ^^^^^^
    |   File "/home/bc-user/anaconda3/envs/vllm-source/lib/python3.11/site-packages/torch/nn/modules/module.py", line 1532, in _wrapped_call_impl
    |     return self._call_impl(*args, **kwargs)
    |            ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
    |   File "/home/bc-user/anaconda3/envs/vllm-source/lib/python3.11/site-packages/torch/nn/modules/module.py", line 1541, in _call_impl
    |     return forward_call(*args, **kwargs)
    |            ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
    |   File "/home/bc-user/jeejeelee-vllm/vllm/model_executor/models/qwen2.py", line 209, in forward
    |     hidden_states = self.self_attn(
    |                     ^^^^^^^^^^^^^^^
    |   File "/home/bc-user/anaconda3/envs/vllm-source/lib/python3.11/site-packages/torch/nn/modules/module.py", line 1532, in _wrapped_call_impl
    |     return self._call_impl(*args, **kwargs)
    |            ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
    |   File "/home/bc-user/anaconda3/envs/vllm-source/lib/python3.11/site-packages/torch/nn/modules/module.py", line 1541, in _call_impl
    |     return forward_call(*args, **kwargs)
    |            ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
    |   File "/home/bc-user/jeejeelee-vllm/vllm/model_executor/models/qwen2.py", line 153, in forward
    |     qkv, _ = self.qkv_proj(hidden_states)
    |              ^^^^^^^^^^^^^^^^^^^^^^^^^^^^
    |   File "/home/bc-user/anaconda3/envs/vllm-source/lib/python3.11/site-packages/torch/nn/modules/module.py", line 1532, in _wrapped_call_impl
    |     return self._call_impl(*args, **kwargs)
    |            ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
    |   File "/home/bc-user/anaconda3/envs/vllm-source/lib/python3.11/site-packages/torch/nn/modules/module.py", line 1541, in _call_impl
    |     return forward_call(*args, **kwargs)
    |            ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
    |   File "/home/bc-user/jeejeelee-vllm/vllm/lora/layers.py", line 365, in forward
    |     output_parallel = self.apply(input_, bias)
    |                       ^^^^^^^^^^^^^^^^^^^^^^^^
    |   File "/home/bc-user/jeejeelee-vllm/vllm/lora/layers.py", line 762, in apply
    |     output = self.base_layer.quant_method.apply(self.base_layer, x, bias)
    |              ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
    |   File "/home/bc-user/jeejeelee-vllm/vllm/model_executor/layers/linear.py", line 122, in apply
    |     return F.linear(x, layer.weight, bias)
    |            ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
    | RuntimeError: CUDA error: CUBLAS_STATUS_ALLOC_FAILED when calling `cublasCreate(handle)`
    |
    | The above exception was the direct cause of the following exception:
    |
    | Traceback (most recent call last):
    |   File "/home/bc-user/anaconda3/envs/vllm-source/lib/python3.11/site-packages/starlette/responses.py", line 261, in wrap
    |     await func()
    |   File "/home/bc-user/anaconda3/envs/vllm-source/lib/python3.11/site-packages/starlette/responses.py", line 250, in stream_response
    |     async for chunk in self.body_iterator:
    |   File "/home/bc-user/jeejeelee-vllm/vllm/entrypoints/openai/serving_chat.py", line 228, in chat_completion_stream_generator
    |     async for res in result_generator:
    |   File "/home/bc-user/jeejeelee-vllm/vllm/engine/async_llm_engine.py", line 772, in generate
    |     async for output in self._process_request(
    |   File "/home/bc-user/jeejeelee-vllm/vllm/engine/async_llm_engine.py", line 873, in _process_request
    |     stream = await self.add_request(
    |              ^^^^^^^^^^^^^^^^^^^^^^^
    |   File "/home/bc-user/jeejeelee-vllm/vllm/engine/async_llm_engine.py", line 676, in add_request
    |     self.start_background_loop()
    |   File "/home/bc-user/jeejeelee-vllm/vllm/engine/async_llm_engine.py", line 516, in start_background_loop
    |     raise AsyncEngineDeadError(
    | vllm.engine.async_llm_engine.AsyncEngineDeadError: Background loop has errored already.
    +------------------------------------

@benedict-lee  Your steps seem fine. I'm not sure yet what's causing this issue. I'll try to reproduce your error. Can you run normally based on offline inference?

@benedict-lee Your steps seem fine. I'm not sure yet what's causing this issue. I'll try to reproduce your error. Can you run normally based on offline inference?

I cannot reproduce your error

@benedict-lee Your steps seem fine. I'm not sure yet what's causing this issue. I'll try to reproduce your error. Can you run normally based on offline inference?

I cannot reproduce your error
Can you tell which environments have been tried?

@Yard1  Could you take a look at the reason for CI OOM ? Thank you very much

@jeejeelee it looks like it's working now?

@Yard1  All CI checks have passed. Please take a look, thanks very much

I build a whl package with this branch(refactor-punica-kernel) on ubuntu22.04.
And then install this whl package on a centos7.9.
Then run `python exmaples/multilora-inference.py.
I got this:
(vllm311) [root@localhost vllm]# python examples/multilora_inference.py
WARNING 07-27 11:07:52 _custom_ops.py:14] Failed to import from vllm._C with ImportError("/lib64/libc.so.6: version `GLIBC_2.32' not found (required by /root/miniconda3/envs/vllm311/lib/python3.11/site-packages/vllm/_C.abi3.so)")
INFO 07-27 11:07:53 llm_engine.py:176] Initializing an LLM engine (v0.5.3.post1) with config: model='NousResearch/Llama-2-7b-hf', speculative_config=None, tokenizer='NousResearch/Llama-2-7b-hf', skip_tokenizer_init=False, tokenizer_mode=auto, revision=None, rope_scaling=None, rope_theta=None, tokenizer_revision=None, trust_remote_code=False, dtype=torch.float16, max_seq_len=4096, download_dir=None, load_format=LoadFormat.AUTO, tensor_parallel_size=1, pipeline_parallel_size=1, disable_custom_all_reduce=False, quantization=None, enforce_eager=False, kv_cache_dtype=auto, quantization_param_path=None, device_config=cuda, decoding_config=DecodingConfig(guided_decoding_backend='outlines'), observability_config=ObservabilityConfig(otlp_traces_endpoint=None), seed=0, served_model_name=NousResearch/Llama-2-7b-hf, use_v2_block_manager=False, enable_prefix_caching=False)
INFO 07-27 11:07:54 model_runner.py:720] Starting to load model NousResearch/Llama-2-7b-hf...
INFO 07-27 11:07:55 weight_utils.py:224] Using model weights format ['*.safetensors']
Loading safetensors checkpoint shards:   0% Completed | 0/2 [00:00<?, ?it/s]
Loading safetensors checkpoint shards:  50% Completed | 1/2 [00:01<00:01,  1.84s/it]
Loading safetensors checkpoint shards: 100% Completed | 2/2 [00:06<00:00,  3.77s/it]
Loading safetensors checkpoint shards: 100% Completed | 2/2 [00:06<00:00,  3.48s/it]

INFO 07-27 11:08:02 model_runner.py:732] Loading model weights took 12.5562 GB
/tmp/tmpbwr6jrz8/main.c:6:23: 致命错误：stdatomic.h：没有那个文件或目录
 #include <stdatomic.h>
                       ^
编译中断。

May I ask why is this happening?
here is my env_collect.py result:
(vllm311) [root@localhost vllm]# python collect_env.py
Collecting environment information...
WARNING 07-27 11:09:34 _custom_ops.py:14] Failed to import from vllm._C with ModuleNotFoundError("No module named 'vllm._C'")
PyTorch version: 2.3.1+cu121
Is debug build: False
CUDA used to build PyTorch: 12.1
ROCM used to build PyTorch: N/A

OS: CentOS Linux 7 (Core) (x86_64)
GCC version: (GCC) 4.8.5 20150623 (Red Hat 4.8.5-44)
Clang version: Could not collect
CMake version: version 3.30.1
Libc version: glibc-2.17

Python version: 3.11.9 (main, Apr 19 2024, 16:48:06) [GCC 11.2.0] (64-bit runtime)
Python platform: Linux-3.10.0-1160.119.1.el7.x86_64-x86_64-with-glibc2.17
Is CUDA available: True
CUDA runtime version: 12.4.99
CUDA_MODULE_LOADING set to: LAZY
GPU models and configuration: GPU 0: NVIDIA A800 80GB PCIe
Nvidia driver version: 550.54.14
cuDNN version: Could not collect
HIP runtime version: N/A
MIOpen runtime version: N/A
Is XNNPACK available: True

CPU:
Architecture:          x86_64
CPU op-mode(s):        32-bit, 64-bit
Byte Order:            Little Endian
CPU(s):                32
On-line CPU(s) list:   0-31
Thread(s) per core:    1
Core(s) per socket:    32
座：                 1
NUMA 节点：         1
厂商 ID：           GenuineIntel
CPU 系列：          6
型号：              106
型号名称：        Intel(R) Xeon(R) Gold 6330 CPU @ 2.00GHz
步进：              6
CPU MHz：             1995.311
BogoMIPS：            3990.62
虚拟化：           VT-x
超管理器厂商：  KVM
虚拟化类型：     完全
L1d 缓存：          32K
L1i 缓存：          32K
L2 缓存：           4096K
L3 缓存：           16384K
NUMA 节点0 CPU：    0-31
Flags:                 fpu vme de pse tsc msr pae mce cx8 apic sep mtrr pge mca cmov pat pse36 clflush dts mmx fxsr sse sse2 ss ht syscall nx pdpe1gb rdtscp lm constant_tsc arch_perfmon pebs bts rep_good nopl xtopology eagerfpu pni pclmulqdq dtes64 vmx ssse3 fma cx16 pdcm pcid sse4_1 sse4_2 x2apic movbe popcnt tsc_deadline_timer aes xsave avx f16c rdrand hypervisor lahf_lm abm 3dnowprefetch invpcid_single ssbd rsb_ctxsw ibrs ibpb stibp ibrs_enhanced tpr_shadow vnmi flexpriority ept vpid fsgsbase tsc_adjust bmi1 avx2 smep bmi2 erms invpcid rtm avx512f avx512dq rdseed adx smap avx512ifma clflushopt clwb avx512cd sha_ni avx512bw avx512vl xsaveopt xsavec xgetbv1 arat avx512vbmi umip pku ospke avx512_vbmi2 gfni vaes vpclmulqdq avx512_vnni avx512_bitalg avx512_vpopcntdq md_clear spec_ctrl intel_stibp arch_capabilities

Versions of relevant libraries:
[pip3] numpy==1.26.4
[pip3] nvidia-nccl-cu12==2.20.5
[pip3] torch==2.3.1
[pip3] torchvision==0.18.1
[pip3] transformers==4.43.3
[pip3] triton==3.0.0
[conda] numpy                     1.26.4                   pypi_0    pypi
[conda] nvidia-nccl-cu12          2.20.5                   pypi_0    pypi
[conda] torch                     2.3.1                    pypi_0    pypi
[conda] torchvision               0.18.1                   pypi_0    pypi
[conda] transformers              4.43.3                   pypi_0    pypi
[conda] triton                    3.0.0                    pypi_0    pypi
ROCM Version: Could not collect
Neuron SDK Version: N/A
vLLM Version: 0.5.3.post1
vLLM Build Flags:
CUDA Archs: Not Set; ROCm: Disabled; Neuron: Disabled
GPU Topology:
GPU0    CPU Affinity    NUMA Affinity   GPU NUMA ID
GPU0     X      0-31    0               N/A

Legend:

  X    = Self
  SYS  = Connection traversing PCIe as well as the SMP interconnect between NUMA nodes (e.g., QPI/UPI)
  NODE = Connection traversing PCIe as well as the interconnect between PCIe Host Bridges within a NUMA node
  PHB  = Connection traversing PCIe as well as a PCIe Host Bridge (typically the CPU)
  PXB  = Connection traversing multiple PCIe bridges (without traversing the PCIe Host Bridge)
  PIX  = Connection traversing at most a single PCIe bridge
  NV#  = Connection traversing a bonded set of # NVLinks

It is recommended to check the  below issues first.
WARNING 07-27 11:07:52 _custom_ops.py:14] Failed to import from vllm._C with ImportError("/lib64/libc.so.6: version `GLIBC_2.32' not found (required by /root/miniconda3/envs/vllm311/lib/python3.11/site-packages/vllm/_C.abi3.so)")

and
/tmp/tmpbwr6jrz8/main.c:6:23: 致命错误：stdatomic.h：没有那个文件或目录
 #include <stdatomic.h>

It is recommended to check the below issues first.
WARNING 07-27 11:07:52 _custom_ops.py:14] Failed to import from vllm._C with ImportError("/lib64/libc.so.6: version `GLIBC_2.32' not found (required by /root/miniconda3/envs/vllm311/lib/python3.11/site-packages/vllm/_C.abi3.so)")

and
/tmp/tmpbwr6jrz8/main.c:6:23: 致命错误：stdatomic.h：没有那个文件或目录
 #include <stdatomic.h>

@balcklive    It could also be an issue caused by Ubuntu 22.   FYI: https://github.com/vllm-project/vllm/pull/6517

hi，when  i try pip install -e . , some errors occurred
141 | struct TiledCopy : Copy_Atom
|        ^~~~~~~~~
/usr/local/cuda/include/cuda.h:2866:1: note: and the implicitly-defined constructor does not initialize ‘cuuint64_t CUtensorMap_st::opaque [16]’
2866 |     cuuint64_t opaque[CU_TENSOR_MAP_NUM_QWORDS];
| ^   ~~
/tmp/tmp4svdoami.build-temp/_deps/cutlass-src/include/cutlass/device_kernel.h: In function ‘void cutlass::device_kernel(typename Operator::Params) [with Operator = _GLOBAL__N__38c64651_16_scaled_mm_c3x_cu_22d95651::cutlass_3x_gemm<signed char, cutlass::bfloat16_t, _GLOBAL__N__38c64651_16_scaled_mm_c3x_cu_22d95651::ScaledEpilogueBias, cute::tuple<cute::C<64>, cute::C<64>, cute::C<256> >, cute::tuple<cute::C<1>, cute::C<8>, cute::C<1> >, cutlass::gemm::KernelTmaWarpSpecialized, cutlass::epilogue::TmaWarpSpecialized>::GemmKernel]’:
/tmp/tmp4svdoami.build-temp/_deps/cutlass-src/include/cutlass/device_kernel.h:104:1: note: the ABI for passing parameters with 64-byte alignment has changed in GCC 4.6
104 | void device_kernel(CUTLASS_GRID_CONSTANT typename Operator::Params const params)
| ^~~~~~~~~~~~~
[28/29] Building CUDA object CMakeFiles/_C.dir/csrc/attention/attention_kernels.cu.o
/scratch/AzureNfsServer_INPUT1/vc_data/yingmeiguo2/LLM/vllm/csrc/attention/attention_kernels.cu(709): warning #177-D: variable "thread_group_size" was declared but never referenced
int thread_group_size = ((32 / BLOCK_SIZE) > (1) ? (32 / BLOCK_SIZE) : (1));
^
  Remark: The warnings can be suppressed with "-diag-suppress <warning-number>"

  /scratch/AzureNfsServer_INPUT1/vc_data/yingmeiguo2/LLM/vllm/csrc/attention/attention_kernels.cu(868): warning #177-D: variable "thread_group_size" was declared but never referenced
      int thread_group_size = ((32 / BLOCK_SIZE) > (1) ? (32 / BLOCK_SIZE) : (1));
          ^

  ninja: build stopped: subcommand failed.
  Traceback (most recent call last):
    File "/tmp/pip-build-env-g4n22dl4/overlay/lib/python3.11/site-packages/setuptools/command/editable_wheel.py", line 153, in run
      self._create_wheel_file(bdist_wheel)
    File "/tmp/pip-build-env-g4n22dl4/overlay/lib/python3.11/site-packages/setuptools/command/editable_wheel.py", line 355, in _create_wheel_file
      files, mapping = self._run_build_commands(dist_name, unpacked, lib, tmp)
                       ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
    File "/tmp/pip-build-env-g4n22dl4/overlay/lib/python3.11/site-packages/setuptools/command/editable_wheel.py", line 278, in _run_build_commands
      self._run_build_subcommands()
    File "/tmp/pip-build-env-g4n22dl4/overlay/lib/python3.11/site-packages/setuptools/command/editable_wheel.py", line 305, in _run_build_subcommands
      self.run_command(name)
    File "/tmp/pip-build-env-g4n22dl4/overlay/lib/python3.11/site-packages/setuptools/_distutils/cmd.py", line 316, in run_command
      self.distribution.run_command(command)
    File "/tmp/pip-build-env-g4n22dl4/overlay/lib/python3.11/site-packages/setuptools/dist.py", line 945, in run_command
      super().run_command(command)
    File "/tmp/pip-build-env-g4n22dl4/overlay/lib/python3.11/site-packages/setuptools/_distutils/dist.py", line 989, in run_command
      cmd_obj.run()
    File "/tmp/pip-build-env-g4n22dl4/overlay/lib/python3.11/site-packages/setuptools/command/build_ext.py", line 93, in run
      _build_ext.run(self)
    File "/tmp/pip-build-env-g4n22dl4/overlay/lib/python3.11/site-packages/setuptools/_distutils/command/build_ext.py", line 359, in run
      self.build_extensions()
    File "<string>", line 231, in build_extensions
    File "/vc_data/yingmeiguo2/anaconda3/envs/llm/lib/python3.11/subprocess.py", line 413, in check_call
      raise CalledProcessError(retcode, cmd)
  subprocess.CalledProcessError: Command '['cmake', '--build', '.', '-j=40', '--target=_moe_C', '--target=_C']' returned non-zero exit status 1.
  /tmp/pip-build-env-g4n22dl4/overlay/lib/python3.11/site-packages/setuptools/_distutils/dist.py:989: _DebuggingTips: Problem in editable installation.
  !!

          ********************************************************************************
          An error happened while installing `vllm` in editable mode.

          The following steps are recommended to help debug this problem:

          - Try to install the project normally, without using the editable mode.
            Does the error still persist?
            (If it does, try fixing the problem before attempting the editable mode).
          - If you are using binary extensions, make sure you have all OS-level
            dependencies installed (e.g. compilers, toolchains, binary libraries, ...).
          - Try the latest version of setuptools (maybe the error was already fixed).
          - If you (or your project dependencies) are using any setuptools extension
            or customization, make sure they support the editable mode.

          After following the steps above, if the problem still persists and
          you think this is related to how setuptools handles editable installations,
          please submit a reproducible example
          (see https://stackoverflow.com/help/minimal-reproducible-example) to:

              https://github.com/pypa/setuptools/issues

          See https://setuptools.pypa.io/en/latest/userguide/development_mode.html for details.
          ********************************************************************************

  !!
    cmd_obj.run()
  Traceback (most recent call last):
    File "/vc_data/yingmeiguo2/anaconda3/envs/llm/lib/python3.11/site-packages/pip/_vendor/pyproject_hooks/_in_process/_in_process.py", line 353, in <module>
      main()
    File "/vc_data/yingmeiguo2/anaconda3/envs/llm/lib/python3.11/site-packages/pip/_vendor/pyproject_hooks/_in_process/_in_process.py", line 335, in main
      json_out['return_val'] = hook(**hook_input['kwargs'])
                               ^^^^^^^^^^^^^^^^^^^^^^^^^^^^
    File "/vc_data/yingmeiguo2/anaconda3/envs/llm/lib/python3.11/site-packages/pip/_vendor/pyproject_hooks/_in_process/_in_process.py", line 273, in build_editable
      return hook(wheel_directory, config_settings, metadata_directory)
             ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
    File "/tmp/pip-build-env-g4n22dl4/overlay/lib/python3.11/site-packages/setuptools/build_meta.py", line 453, in build_editable
      return self._build_with_temp_dir(
             ^^^^^^^^^^^^^^^^^^^^^^^^^^
    File "/tmp/pip-build-env-g4n22dl4/overlay/lib/python3.11/site-packages/setuptools/build_meta.py", line 397, in _build_with_temp_dir
      self.run_setup()
    File "/tmp/pip-build-env-g4n22dl4/overlay/lib/python3.11/site-packages/setuptools/build_meta.py", line 313, in run_setup
      exec(code, locals())
    File "<string>", line 449, in <module>
    File "/tmp/pip-build-env-g4n22dl4/overlay/lib/python3.11/site-packages/setuptools/__init__.py", line 108, in setup
      return distutils.core.setup(**attrs)
             ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
    File "/tmp/pip-build-env-g4n22dl4/overlay/lib/python3.11/site-packages/setuptools/_distutils/core.py", line 184, in setup
      return run_commands(dist)
             ^^^^^^^^^^^^^^^^^^
    File "/tmp/pip-build-env-g4n22dl4/overlay/lib/python3.11/site-packages/setuptools/_distutils/core.py", line 200, in run_commands
      dist.run_commands()
    File "/tmp/pip-build-env-g4n22dl4/overlay/lib/python3.11/site-packages/setuptools/_distutils/dist.py", line 970, in run_commands
      self.run_command(cmd)
    File "/tmp/pip-build-env-g4n22dl4/overlay/lib/python3.11/site-packages/setuptools/dist.py", line 945, in run_command
      super().run_command(command)
    File "/tmp/pip-build-env-g4n22dl4/overlay/lib/python3.11/site-packages/setuptools/_distutils/dist.py", line 989, in run_command
      cmd_obj.run()
    File "/tmp/pip-build-env-g4n22dl4/overlay/lib/python3.11/site-packages/setuptools/command/editable_wheel.py", line 153, in run
      self._create_wheel_file(bdist_wheel)
    File "/tmp/pip-build-env-g4n22dl4/overlay/lib/python3.11/site-packages/setuptools/command/editable_wheel.py", line 355, in _create_wheel_file
      files, mapping = self._run_build_commands(dist_name, unpacked, lib, tmp)
                       ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
    File "/tmp/pip-build-env-g4n22dl4/overlay/lib/python3.11/site-packages/setuptools/command/editable_wheel.py", line 278, in _run_build_commands
      self._run_build_subcommands()
    File "/tmp/pip-build-env-g4n22dl4/overlay/lib/python3.11/site-packages/setuptools/command/editable_wheel.py", line 305, in _run_build_subcommands
      self.run_command(name)
    File "/tmp/pip-build-env-g4n22dl4/overlay/lib/python3.11/site-packages/setuptools/_distutils/cmd.py", line 316, in run_command
      self.distribution.run_command(command)
    File "/tmp/pip-build-env-g4n22dl4/overlay/lib/python3.11/site-packages/setuptools/dist.py", line 945, in run_command
      super().run_command(command)
    File "/tmp/pip-build-env-g4n22dl4/overlay/lib/python3.11/site-packages/setuptools/_distutils/dist.py", line 989, in run_command
      cmd_obj.run()
    File "/tmp/pip-build-env-g4n22dl4/overlay/lib/python3.11/site-packages/setuptools/command/build_ext.py", line 93, in run
      _build_ext.run(self)
    File "/tmp/pip-build-env-g4n22dl4/overlay/lib/python3.11/site-packages/setuptools/_distutils/command/build_ext.py", line 359, in run
      self.build_extensions()
    File "<string>", line 231, in build_extensions
    File "/vc_data/yingmeiguo2/anaconda3/envs/llm/lib/python3.11/subprocess.py", line 413, in check_call
      raise CalledProcessError(retcode, cmd)
  subprocess.CalledProcessError: Command '['cmake', '--build', '.', '-j=40', '--target=_moe_C', '--target=_C']' returned non-zero exit status 1.
  [end of output]

note: This error originates from a subprocess, and is likely not a problem with pip.
ERROR: Failed building editable for vllm
Failed to build vllm
ERROR: Could not build wheels for vllm, which is required to install pyproject.toml-based projects

@qniguogym  Can you provide information about your environment?

@qniguogym Can you provide information about your environment?

Linux version 5.15.0-1038-azure (buildd@lcy02-amd64-004) (gcc (Ubuntu 9.4.0-1ubuntu1~20.04.1) 9.4.0, GNU ld (GNU Binutils for Ubuntu) 2.34) #45~20.04.1-Ubuntu SMP Tue Apr 25 18:45:15 UTC 2023
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2023 NVIDIA Corporation
Built on Mon_Apr__3_17:16:06_PDT_2023
Cuda compilation tools, release 12.1, V12.1.105
Build cuda_12.1.r12.1/compiler.32688072_0
pytorch 2.3.1
transformers 4.38.2

@qniguogym Can you provide information about your environment?

Linux version 5.15.0-1038-azure (buildd@lcy02-amd64-004) (gcc (Ubuntu 9.4.0-1ubuntu1~20.04.1) 9.4.0, GNU ld (GNU Binutils for Ubuntu) 2.34) #45~20.04.1-Ubuntu SMP Tue Apr 25 18:45:15 UTC 2023
nvcc: NVIDIA (R) Cuda compiler driver Copyright (c) 2005-2023 NVIDIA Corporation Built on Mon_Apr__3_17:16:06_PDT_2023 Cuda compilation tools, release 12.1, V12.1.105 Build cuda_12.1.r12.1/compiler.32688072_0
pytorch 2.3.1 transformers 4.38.2

What GPU are you using? I can compile it normally on my machine. I'm not sure if it's because of the GPU.

V100
郭颖美
清华大学计算机科学与技术系2018级研究生
@.***
手机：18810288163
---- Replied Message ----
| From | Jee Jee @.> |
| Date | 07/29/2024 21:03 |
| To | vllm-project/vllm @.> |
| Cc | 郭颖美 @.>,
Mention @.> |
| Subject | Re: [vllm-project/vllm] [Kernel][RFC] Refactor the punica kernel based on Triton (PR #5036) |
@qniguogym Can you provide information about your environment?
Linux version 5.15.0-1038-azure @.***) (gcc (Ubuntu 9.4.0-1ubuntu120.04.1) 9.4.0, GNU ld (GNU Binutils for Ubuntu) 2.34) #4520.04.1-Ubuntu SMP Tue Apr 25 18:45:15 UTC 2023
nvcc: NVIDIA (R) Cuda compiler driver Copyright (c) 2005-2023 NVIDIA Corporation Built on Mon_Apr__3_17:16:06_PDT_2023 Cuda compilation tools, release 12.1, V12.1.105 Build cuda_12.1.r12.1/compiler.32688072_0
pytorch 2.3.1 transformers 4.38.2
What GPU are you using? I can compile it normally on my machine. I'm not sure if it's because of the GPU.
—
Reply to this email directly, view it on GitHub, or unsubscribe.
You are receiving this because you were mentioned.Message ID: @.***>

V100 郭颖美 清华大学计算机科学与技术系2018级研究生 @.*** 手机：18810288163
…
---- Replied Message ---- | From | Jee Jee @.> | | Date | 07/29/2024 21:03 | | To | vllm-project/vllm @.> | | Cc | 郭颖美 @.>, Mention @.> | | Subject | Re: [vllm-project/vllm] [Kernel][RFC] Refactor the punica kernel based on Triton (PR #5036) | @qniguogym Can you provide information about your environment? Linux version 5.15.0-1038-azure @.) (gcc (Ubuntu 9.4.0-1ubuntu120.04.1) 9.4.0, GNU ld (GNU Binutils for Ubuntu) 2.34) #4520.04.1-Ubuntu SMP Tue Apr 25 18:45:15 UTC 2023 nvcc: NVIDIA (R) Cuda compiler driver Copyright (c) 2005-2023 NVIDIA Corporation Built on Mon_Apr__3_17:16:06_PDT_2023 Cuda compilation tools, release 12.1, V12.1.105 Build cuda_12.1.r12.1/compiler.32688072_0 pytorch 2.3.1 transformers 4.38.2 What GPU are you using? I can compile it normally on my machine. I'm not sure if it's because of the GPU. — Reply to this email directly, view it on GitHub, or unsubscribe. You are receiving this because you were mentioned.Message ID: @.>

I

V100 郭颖美 清华大学计算机科学与技术系2018级研究生 @.*** 手机：18810288163
…
---- Replied Message ---- | From | Jee Jee @.> | | Date | 07/29/2024 21:03 | | To | vllm-project/vllm @.> | | Cc | 郭颖美 @.>, Mention @.> | | Subject | Re: [vllm-project/vllm] [Kernel][RFC] Refactor the punica kernel based on Triton (PR #5036) | @qniguogym Can you provide information about your environment? Linux version 5.15.0-1038-azure @.) (gcc (Ubuntu 9.4.0-1ubuntu120.04.1) 9.4.0, GNU ld (GNU Binutils for Ubuntu) 2.34) #4520.04.1-Ubuntu SMP Tue Apr 25 18:45:15 UTC 2023 nvcc: NVIDIA (R) Cuda compiler driver Copyright (c) 2005-2023 NVIDIA Corporation Built on Mon_Apr__3_17:16:06_PDT_2023 Cuda compilation tools, release 12.1, V12.1.105 Build cuda_12.1.r12.1/compiler.32688072_0 pytorch 2.3.1 transformers 4.38.2 What GPU are you using? I can compile it normally on my machine. I'm not sure if it's because of the GPU. — Reply to this email directly, view it on GitHub, or unsubscribe. You are receiving this because you were mentioned.Message ID: @.>

Perhaps you could try switching to a newer GPU architecture to try it again

I have a v100, I can see if the issue reproduces on mine or if it's a different environment issue.

I have a v100, I can see if the issue reproduces on mine or if it's a different environment issue. But, for what it's worth, I think one of the benefits of rewriting the punica kernel to triton is enabling broader GPU platform support, so it would be awesome if we could find a way to get this working on Volta-architecture devices if that's the issue

Great, I look forward to your feedback.

@Yard1 @mgoin  The current CI is too unstable :joy:

I have a v100, I can see if the issue reproduces on mine or if it's a different environment issue. But, for what it's worth, I think one of the benefits of rewriting the punica kernel to triton is enabling broader GPU platform support, so it would be awesome if we could find a way to get this working on Volta-architecture devices if that's the issue

I also meet this error on A100, so may be it is irrelavant to GPU version.

I have a v100, I can see if the issue reproduces on mine or if it's a different environment issue. But, for what it's worth, I think one of the benefits of rewriting the punica kernel to triton is enabling broader GPU platform support, so it would be awesome if we could find a way to get this working on Volta-architecture devices if that's the issue

I also meet this error on A100, so may be it is irrelavant to GPU version.
I was able to successfully install it on the A800. My environment is: https://github.com/vllm-project/vllm/issues/6103#issue-2388613551

@Yard1 Thank you very much for your review and guidance. All CIs have passed. Could you please merge this into main?

Hi guys, can you provide details on how I can test the triton implementation (as opposed to punica) so that I can try to reproduce the issue on my v100? test script, install/build/serve command & configs, etc?

Hi @Yard1 @jeejeelee - I am currently going through the install process on my v100 and will let you know if I run into the same errors. Can you let me know how I can test the implementation at runtime (not just compile-time) so that I can verify compatibility? e.g. what model, config, and LoRAs are you using for testing, or can I run the existing tests somehow? I apologize as I'm not as familiar with the project's test suite.

Hi @Yard1 @jeejeelee - I am currently going through the install process on my v100 and will let you know if I run into the same errors. Can you let me know how I can test the implementation at runtime (not just compile-time) so that I can verify compatibility? e.g. what model, config, and LoRAs are you using for testing, or can I run the existing tests somehow? I apologize as I'm not as familiar with the project's test suite.

Thank you for being willing to spend time verifying on the V100.I f your installation is successful, you can proceed with:

Kernel testing, which can be conducted using test_script
Model testing, which can be conducted using the following  scripts:

https://github.com/vllm-project/vllm/blob/main/tests/lora/test_llama.py
https://github.com/vllm-project/vllm/blob/main/examples/multilora_inference.py

Please feel free to contact me if you have any questions.

Hi @Yard1 @jeejeelee - I am currently going through the install process on my v100 and will let you know if I run into the same errors. Can you let me know how I can test the implementation at runtime (not just compile-time) so that I can verify compatibility? e.g. what model, config, and LoRAs are you using for testing, or can I run the existing tests somehow? I apologize as I'm not as familiar with the project's test suite.

Thank you for being willing to spend time verifying on the V100. If your installation is successful, you can proceed with:

Kernel testing, which can be conducted using test_script

Model testing, which can be conducted using the following  scripts:

https://github.com/vllm-project/vllm/blob/main/tests/lora/test_llama.py
https://github.com/vllm-project/vllm/blob/main/examples/multilora_inference.py

Please feel free to contact me if you have any questions.

Install completed successfully on my v100, just waiting on recommended testing strategy

I keep running into import errors related to some of the testing stuff. Tried running from vllm/ and vllm/tests/lora but no luck. Any auggestions?

Hi @Yard1 @jeejeelee - I am currently going through the install process on my v100 and will let you know if I run into the same errors. Can you let me know how I can test the implementation at runtime (not just compile-time) so that I can verify compatibility? e.g. what model, config, and LoRAs are you using for testing, or can I run the existing tests somehow? I apologize as I'm not as familiar with the project's test suite.

Thank you for being willing to spend time verifying on the V100. If your installation is successful, you can proceed with:

Kernel testing, which can be conducted using test_script

Model testing, which can be conducted using the following  scripts:

https://github.com/vllm-project/vllm/blob/main/tests/lora/test_llama.py
https://github.com/vllm-project/vllm/blob/main/examples/multilora_inference.py

Please feel free to contact me if you have any questions.

Install completed successfully on my v100, just waiting on recommended testing strategy

I installled successfully on V100 with docker, but error occured with conda.  after installed with docker, another error：
INFO 08-01 02:24:44 model_runner.py:731] Loading model weights took 12.5562 GB
[rank0]: Traceback (most recent call last):
[rank0]:   File "/workspace/LLM/infer_lora.py", line 823, in
[rank0]:     main()
[rank0]:   File "/workspace/LLM/infer_lora.py", line 718, in main
[rank0]:     engine = initialize_engine()
[rank0]:   File "/workspace/LLM/infer_lora.py", line 739, in initialize_engine
[rank0]:     return LLMEngine.from_engine_args(engine_args)
[rank0]:   File "/workspace/LLM/vllm/vllm/engine/llm_engine.py", line 446, in from_engine_args
[rank0]:     engine = cls(
[rank0]:   File "/workspace/LLM/vllm/vllm/engine/llm_engine.py", line 264, in init
[rank0]:     self._initialize_kv_caches()
[rank0]:   File "/workspace/LLM/vllm/vllm/engine/llm_engine.py", line 363, in _initialize_kv_caches
[rank0]:     self.model_executor.determine_num_available_blocks())
[rank0]:   File "/workspace/LLM/vllm/vllm/executor/gpu_executor.py", line 94, in determine_num_available_blocks
[rank0]:     return self.driver_worker.determine_num_available_blocks()
[rank0]:   File "/usr/local/lib/python3.10/dist-packages/torch/utils/_contextlib.py", line 115, in decorate_context
[rank0]:     return func(*args, **kwargs)
[rank0]:   File "/workspace/LLM/vllm/vllm/worker/worker.py", line 179, in determine_num_available_blocks
[rank0]:     self.model_runner.profile_run()
[rank0]:   File "/usr/local/lib/python3.10/dist-packages/torch/utils/_contextlib.py", line 115, in decorate_context
[rank0]:     return func(*args, **kwargs)
[rank0]:   File "/workspace/LLM/vllm/vllm/worker/model_runner.py", line 934, in profile_run
[rank0]:     self.execute_model(model_input, kv_caches, intermediate_tensors)
[rank0]:   File "/usr/local/lib/python3.10/dist-packages/torch/utils/_contextlib.py", line 115, in decorate_context
[rank0]:     return func(*args, **kwargs)
[rank0]:   File "/workspace/LLM/vllm/vllm/worker/model_runner.py", line 1357, in execute_model
[rank0]:     hidden_or_intermediate_states = model_executable(
[rank0]:   File "/usr/local/lib/python3.10/dist-packages/torch/nn/modules/module.py", line 1532, in _wrapped_call_impl
[rank0]:     return self._call_impl(*args, **kwargs)
[rank0]:   File "/usr/local/lib/python3.10/dist-packages/torch/nn/modules/module.py", line 1541, in _call_impl
[rank0]:     return forward_call(*args, **kwargs)
[rank0]:   File "/workspace/LLM/vllm/vllm/model_executor/models/llama.py", line 423, in forward
[rank0]:     model_output = self.model(input_ids, positions, kv_caches,
[rank0]:   File "/usr/local/lib/python3.10/dist-packages/torch/nn/modules/module.py", line 1532, in _wrapped_call_impl
[rank0]:     return self._call_impl(*args, **kwargs)
[rank0]:   File "/usr/local/lib/python3.10/dist-packages/torch/nn/modules/module.py", line 1541, in _call_impl
[rank0]:     return forward_call(*args, **kwargs)
[rank0]:   File "/workspace/LLM/vllm/vllm/model_executor/models/llama.py", line 322, in forward
[rank0]:     hidden_states, residual = layer(
[rank0]:   File "/usr/local/lib/python3.10/dist-packages/torch/nn/modules/module.py", line 1532, in _wrapped_call_impl
[rank0]:     return self._call_impl(*args, **kwargs)
[rank0]:   File "/usr/local/lib/python3.10/dist-packages/torch/nn/modules/module.py", line 1541, in _call_impl
[rank0]:     return forward_call(*args, **kwargs)
[rank0]:   File "/workspace/LLM/vllm/vllm/model_executor/models/llama.py", line 241, in forward
[rank0]:     hidden_states = self.input_layernorm(hidden_states)
[rank0]:   File "/usr/local/lib/python3.10/dist-packages/torch/nn/modules/module.py", line 1532, in _wrapped_call_impl
[rank0]:     return self._call_impl(*args, **kwargs)
[rank0]:   File "/usr/local/lib/python3.10/dist-packages/torch/nn/modules/module.py", line 1541, in _call_impl
[rank0]:     return forward_call(*args, **kwargs)
[rank0]:   File "/workspace/LLM/vllm/vllm/model_executor/custom_op.py", line 13, in forward
[rank0]:     return self._forward_method(*args, **kwargs)
[rank0]:   File "/workspace/LLM/vllm/vllm/model_executor/layers/layernorm.py", line 61, in forward_cuda
[rank0]:     out = torch.empty_like(x)
[rank0]: RuntimeError: CUDA error: no kernel image is available for execution on the device
[rank0]: Compile with TORCH_USE_CUDA_DSA to enable device-side assertions.

I keep running into import errors related to some of the testing stuff. Tried running from vllm/ and vllm/tests/lora but no luck. Any auggestions?

run pytest test_punica_sizes.py rather than python test_punica_sizes.py

Hi @Yard1 @jeejeelee - I am currently going through the install process on my v100 and will let you know if I run into the same errors. Can you let me know how I can test the implementation at runtime (not just compile-time) so that I can verify compatibility? e.g. what model, config, and LoRAs are you using for testing, or can I run the existing tests somehow? I apologize as I'm not as familiar with the project's test suite.

Thank you for being willing to spend time verifying on the V100. If your installation is successful, you can proceed with:

Kernel testing, which can be conducted using test_script

Model testing, which can be conducted using the following  scripts:

https://github.com/vllm-project/vllm/blob/main/tests/lora/test_llama.py
https://github.com/vllm-project/vllm/blob/main/examples/multilora_inference.py

Please feel free to contact me if you have any questions.

Install completed successfully on my v100, just waiting on recommended testing strategy

I installled successfully on V100 with docker, but error occured with conda. after installed with docker, another error：
INFO 08-01 02:24:44 model_runner.py:731] Loading model weights took 12.5562 GB [rank0]: Traceback (most recent call last): [rank0]: File "/workspace/LLM/infer_lora.py", line 823, in [rank0]: main() [rank0]: File "/workspace/LLM/infer_lora.py", line 718, in main [rank0]: engine = initialize_engine() [rank0]: File "/workspace/LLM/infer_lora.py", line 739, in initialize_engine [rank0]: return LLMEngine.from_engine_args(engine_args) [rank0]: File "/workspace/LLM/vllm/vllm/engine/llm_engine.py", line 446, in from_engine_args [rank0]: engine = cls( [rank0]: File "/workspace/LLM/vllm/vllm/engine/llm_engine.py", line 264, in init [rank0]: self._initialize_kv_caches() [rank0]: File "/workspace/LLM/vllm/vllm/engine/llm_engine.py", line 363, in _initialize_kv_caches [rank0]: self.model_executor.determine_num_available_blocks()) [rank0]: File "/workspace/LLM/vllm/vllm/executor/gpu_executor.py", line 94, in determine_num_available_blocks [rank0]: return self.driver_worker.determine_num_available_blocks() [rank0]: File "/usr/local/lib/python3.10/dist-packages/torch/utils/_contextlib.py", line 115, in decorate_context [rank0]: return func(*args, **kwargs) [rank0]: File "/workspace/LLM/vllm/vllm/worker/worker.py", line 179, in determine_num_available_blocks [rank0]: self.model_runner.profile_run() [rank0]: File "/usr/local/lib/python3.10/dist-packages/torch/utils/_contextlib.py", line 115, in decorate_context [rank0]: return func(*args, **kwargs) [rank0]: File "/workspace/LLM/vllm/vllm/worker/model_runner.py", line 934, in profile_run [rank0]: self.execute_model(model_input, kv_caches, intermediate_tensors) [rank0]: File "/usr/local/lib/python3.10/dist-packages/torch/utils/_contextlib.py", line 115, in decorate_context [rank0]: return func(*args, **kwargs) [rank0]: File "/workspace/LLM/vllm/vllm/worker/model_runner.py", line 1357, in execute_model [rank0]: hidden_or_intermediate_states = model_executable( [rank0]: File "/usr/local/lib/python3.10/dist-packages/torch/nn/modules/module.py", line 1532, in _wrapped_call_impl [rank0]: return self._call_impl(*args, **kwargs) [rank0]: File "/usr/local/lib/python3.10/dist-packages/torch/nn/modules/module.py", line 1541, in _call_impl [rank0]: return forward_call(*args, **kwargs) [rank0]: File "/workspace/LLM/vllm/vllm/model_executor/models/llama.py", line 423, in forward [rank0]: model_output = self.model(input_ids, positions, kv_caches, [rank0]: File "/usr/local/lib/python3.10/dist-packages/torch/nn/modules/module.py", line 1532, in _wrapped_call_impl [rank0]: return self._call_impl(*args, **kwargs) [rank0]: File "/usr/local/lib/python3.10/dist-packages/torch/nn/modules/module.py", line 1541, in _call_impl [rank0]: return forward_call(*args, **kwargs) [rank0]: File "/workspace/LLM/vllm/vllm/model_executor/models/llama.py", line 322, in forward [rank0]: hidden_states, residual = layer( [rank0]: File "/usr/local/lib/python3.10/dist-packages/torch/nn/modules/module.py", line 1532, in _wrapped_call_impl [rank0]: return self._call_impl(*args, **kwargs) [rank0]: File "/usr/local/lib/python3.10/dist-packages/torch/nn/modules/module.py", line 1541, in _call_impl [rank0]: return forward_call(*args, **kwargs) [rank0]: File "/workspace/LLM/vllm/vllm/model_executor/models/llama.py", line 241, in forward [rank0]: hidden_states = self.input_layernorm(hidden_states) [rank0]: File "/usr/local/lib/python3.10/dist-packages/torch/nn/modules/module.py", line 1532, in _wrapped_call_impl [rank0]: return self._call_impl(*args, **kwargs) [rank0]: File "/usr/local/lib/python3.10/dist-packages/torch/nn/modules/module.py", line 1541, in _call_impl [rank0]: return forward_call(*args, **kwargs) [rank0]: File "/workspace/LLM/vllm/vllm/model_executor/custom_op.py", line 13, in forward [rank0]: return self._forward_method(*args, **kwargs) [rank0]: File "/workspace/LLM/vllm/vllm/model_executor/layers/layernorm.py", line 61, in forward_cuda [rank0]: out = torch.empty_like(x) [rank0]: RuntimeError: CUDA error: no kernel image is available for execution on the device [rank0]: Compile with TORCH_USE_CUDA_DSA to enable device-side assertions.

please  run python collect_env.py, and then provide the printed information.

Hi @Yard1 @jeejeelee - I am currently going through the install process on my v100 and will let you know if I run into the same errors. Can you let me know how I can test the implementation at runtime (not just compile-time) so that I can verify compatibility? e.g. what model, config, and LoRAs are you using for testing, or can I run the existing tests somehow? I apologize as I'm not as familiar with the project's test suite.

Thank you for being willing to spend time verifying on the V100. If your installation is successful, you can proceed with:

Kernel testing, which can be conducted using test_script

Model testing, which can be conducted using the following  scripts:

https://github.com/vllm-project/vllm/blob/main/tests/lora/test_llama.py
https://github.com/vllm-project/vllm/blob/main/examples/multilora_inference.py

Please feel free to contact me if you have any questions.

Install completed successfully on my v100, just waiting on recommended testing strategy

I installled successfully on V100 with docker, but error occured with conda. after installed with docker, another error：
INFO 08-01 02:24:44 model_runner.py:731] Loading model weights took 12.5562 GB [rank0]: Traceback (most recent call last): [rank0]: File "/workspace/LLM/infer_lora.py", line 823, in [rank0]: main() [rank0]: File "/workspace/LLM/infer_lora.py", line 718, in main [rank0]: engine = initialize_engine() [rank0]: File "/workspace/LLM/infer_lora.py", line 739, in initialize_engine [rank0]: return LLMEngine.from_engine_args(engine_args) [rank0]: File "/workspace/LLM/vllm/vllm/engine/llm_engine.py", line 446, in from_engine_args [rank0]: engine = cls( [rank0]: File "/workspace/LLM/vllm/vllm/engine/llm_engine.py", line 264, in init [rank0]: self._initialize_kv_caches() [rank0]: File "/workspace/LLM/vllm/vllm/engine/llm_engine.py", line 363, in _initialize_kv_caches [rank0]: self.model_executor.determine_num_available_blocks()) [rank0]: File "/workspace/LLM/vllm/vllm/executor/gpu_executor.py", line 94, in determine_num_available_blocks [rank0]: return self.driver_worker.determine_num_available_blocks() [rank0]: File "/usr/local/lib/python3.10/dist-packages/torch/utils/_contextlib.py", line 115, in decorate_context [rank0]: return func(*args, **kwargs) [rank0]: File "/workspace/LLM/vllm/vllm/worker/worker.py", line 179, in determine_num_available_blocks [rank0]: self.model_runner.profile_run() [rank0]: File "/usr/local/lib/python3.10/dist-packages/torch/utils/_contextlib.py", line 115, in decorate_context [rank0]: return func(*args, **kwargs) [rank0]: File "/workspace/LLM/vllm/vllm/worker/model_runner.py", line 934, in profile_run [rank0]: self.execute_model(model_input, kv_caches, intermediate_tensors) [rank0]: File "/usr/local/lib/python3.10/dist-packages/torch/utils/_contextlib.py", line 115, in decorate_context [rank0]: return func(*args, **kwargs) [rank0]: File "/workspace/LLM/vllm/vllm/worker/model_runner.py", line 1357, in execute_model [rank0]: hidden_or_intermediate_states = model_executable( [rank0]: File "/usr/local/lib/python3.10/dist-packages/torch/nn/modules/module.py", line 1532, in _wrapped_call_impl [rank0]: return self._call_impl(*args, **kwargs) [rank0]: File "/usr/local/lib/python3.10/dist-packages/torch/nn/modules/module.py", line 1541, in _call_impl [rank0]: return forward_call(*args, **kwargs) [rank0]: File "/workspace/LLM/vllm/vllm/model_executor/models/llama.py", line 423, in forward [rank0]: model_output = self.model(input_ids, positions, kv_caches, [rank0]: File "/usr/local/lib/python3.10/dist-packages/torch/nn/modules/module.py", line 1532, in _wrapped_call_impl [rank0]: return self._call_impl(*args, **kwargs) [rank0]: File "/usr/local/lib/python3.10/dist-packages/torch/nn/modules/module.py", line 1541, in _call_impl [rank0]: return forward_call(*args, **kwargs) [rank0]: File "/workspace/LLM/vllm/vllm/model_executor/models/llama.py", line 322, in forward [rank0]: hidden_states, residual = layer( [rank0]: File "/usr/local/lib/python3.10/dist-packages/torch/nn/modules/module.py", line 1532, in _wrapped_call_impl [rank0]: return self._call_impl(*args, **kwargs) [rank0]: File "/usr/local/lib/python3.10/dist-packages/torch/nn/modules/module.py", line 1541, in _call_impl [rank0]: return forward_call(*args, **kwargs) [rank0]: File "/workspace/LLM/vllm/vllm/model_executor/models/llama.py", line 241, in forward [rank0]: hidden_states = self.input_layernorm(hidden_states) [rank0]: File "/usr/local/lib/python3.10/dist-packages/torch/nn/modules/module.py", line 1532, in _wrapped_call_impl [rank0]: return self._call_impl(*args, **kwargs) [rank0]: File "/usr/local/lib/python3.10/dist-packages/torch/nn/modules/module.py", line 1541, in _call_impl [rank0]: return forward_call(*args, **kwargs) [rank0]: File "/workspace/LLM/vllm/vllm/model_executor/custom_op.py", line 13, in forward [rank0]: return self._forward_method(*args, **kwargs) [rank0]: File "/workspace/LLM/vllm/vllm/model_executor/layers/layernorm.py", line 61, in forward_cuda [rank0]: out = torch.empty_like(x) [rank0]: RuntimeError: CUDA error: no kernel image is available for execution on the device [rank0]: Compile with TORCH_USE_CUDA_DSA to enable device-side assertions.

please run python collect_env.py, and then provide the printed information.

PyTorch version: 2.1.2+cu121
Is debug build: False
CUDA used to build PyTorch: 12.1
ROCM used to build PyTorch: N/A
OS: Ubuntu 20.04.6 LTS (x86_64)
GCC version: (Ubuntu 9.4.0-1ubuntu1~20.04.2) 9.4.0
Clang version: Could not collect
CMake version: version 3.29.6
Libc version: glibc-2.31
Python version: 3.11.9 | packaged by conda-forge | (main, Apr 19 2024, 18:36:13) [GCC 12.3.0] (64-bit runtime)
Python platform: Linux-5.15.0-1038-azure-x86_64-with-glibc2.31
Is CUDA available: True
CUDA runtime version: 12.1.105
CUDA_MODULE_LOADING set to: LAZY
GPU models and configuration:
GPU 0: Tesla V100-SXM2-32GB
GPU 1: Tesla V100-SXM2-32GB
GPU 2: Tesla V100-SXM2-32GB
GPU 3: Tesla V100-SXM2-32GB
GPU 4: Tesla V100-SXM2-32GB
GPU 5: Tesla V100-SXM2-32GB
GPU 6: Tesla V100-SXM2-32GB
GPU 7: Tesla V100-SXM2-32GB
Nvidia driver version: 525.105.17
cuDNN version: Probably one of the following:
/usr/lib/x86_64-linux-gnu/libcudnn.so.8.9.0
/usr/lib/x86_64-linux-gnu/libcudnn_adv_infer.so.8.9.0
/usr/lib/x86_64-linux-gnu/libcudnn_adv_train.so.8.9.0
/usr/lib/x86_64-linux-gnu/libcudnn_cnn_infer.so.8.9.0
/usr/lib/x86_64-linux-gnu/libcudnn_cnn_train.so.8.9.0
/usr/lib/x86_64-linux-gnu/libcudnn_ops_infer.so.8.9.0
/usr/lib/x86_64-linux-gnu/libcudnn_ops_train.so.8.9.0
HIP runtime version: N/A
MIOpen runtime version: N/A
Is XNNPACK available: True
CPU:
Architecture:                    x86_64
CPU op-mode(s):                  32-bit, 64-bit
Byte Order:                      Little Endian
Address sizes:                   46 bits physical, 48 bits virtual
CPU(s):                          40
On-line CPU(s) list:             0-39
Thread(s) per core:              1
Core(s) per socket:              20
Socket(s):                       2
NUMA node(s):                    2
Vendor ID:                       GenuineIntel
CPU family:                      6
Model:                           85
Model name:                      Intel(R) Xeon(R) Platinum 8168 CPU @ 2.70GHz
Stepping:                        4
CPU MHz:                         2693.672
BogoMIPS:                        5387.34
Hypervisor vendor:               Microsoft
Virtualization type:             full
L1d cache:                       1.3 MiB
L1i cache:                       1.3 MiB
L2 cache:                        40 MiB
L3 cache:                        66 MiB
NUMA node0 CPU(s):               0-19
NUMA node1 CPU(s):               20-39
Vulnerability Itlb multihit:     Not affected
Vulnerability L1tf:              Mitigation; PTE Inversion
Vulnerability Mds:               Mitigation; Clear CPU buffers; SMT Host state unknown
Vulnerability Meltdown:          Mitigation; PTI
Vulnerability Mmio stale data:   Vulnerable: Clear CPU buffers attempted, no microcode; SMT Host state unknown
Vulnerability Retbleed:          Vulnerable
Vulnerability Spec store bypass: Vulnerable
Vulnerability Spectre v1:        Mitigation; usercopy/swapgs barriers and __user pointer sanitization
Vulnerability Spectre v2:        Mitigation; Retpolines, STIBP disabled, RSB filling, PBRSB-eIBRS Not affected
Vulnerability Srbds:             Not affected
Vulnerability Tsx async abort:   Mitigation; Clear CPU buffers; SMT Host state unknown
Flags:                           fpu vme de pse tsc msr pae mce cx8 apic sep mtrr pge mca cmov pat pse36 clflush mmx fxsr sse sse2 ss ht syscall nx pdpe1gb rdtscp lm constant_tsc rep_good nopl xtopology cpuid pni pclmulqdq ssse3 fma cx16 pcid sse4_1 sse4_2 movbe popcnt aes xsave avx f16c rdrand hypervisor lahf_lm abm 3dnowprefetch invpcid_single pti fsgsbase bmi1 hle avx2 smep bmi2 erms invpcid rtm avx512f avx512dq rdseed adx smap clflushopt clwb avx512cd avx512bw avx512vl xsaveopt xsavec xgetbv1 xsaves md_clear arch_capabilities
Versions of relevant libraries:
[pip3] mypy-extensions==1.0.0
[pip3] numpy==1.26.4
[pip3] nvidia-nccl-cu12==2.18.1
[pip3] onnxruntime-gpu==1.18.0
[pip3] sentence-transformers==3.0.1
[pip3] torch==2.1.2
[pip3] torchaudio==2.3.1
[pip3] torchvision==0.18.1
[pip3] transformers==4.44.0.dev0
[pip3] triton==2.1.0
[conda] blas                      1.0                         mkl
[conda] libblas                   3.9.0            16_linux64_mkl    conda-forge
[conda] libcblas                  3.9.0            16_linux64_mkl    conda-forge
[conda] liblapack                 3.9.0            16_linux64_mkl    conda-forge
[conda] libopenvino-pytorch-frontend 2024.2.0             he02047a_1    conda-forge
[conda] mkl                       2022.1.0           hc2b9512_224
[conda] numpy                     1.26.4          py311h64a7726_0    conda-forge
[conda] nvidia-nccl-cu12          2.18.1                   pypi_0    pypi
[conda] pytorch-cuda              12.1                 ha16c6d3_5    pytorch
[conda] pytorch-mutex             1.0                        cuda    pytorch
[conda] sentence-transformers     3.0.1                    pypi_0    pypi
[conda] torch                     2.1.2                    pypi_0    pypi
[conda] torchaudio                2.3.1               py311_cu121    pytorch
[conda] torchvision               0.18.1              py311_cu121    pytorch
[conda] transformers              4.44.0.dev0              pypi_0    pypi
[conda] triton                    2.1.0                    pypi_0    pypi
ROCM Version: Could not collect
Neuron SDK Version: N/A
vLLM Version: N/A
vLLM Build Flags:
CUDA Archs: Not Set; ROCm: Disabled; Neuron: Disabled
GPU Topology:
GPU0    GPU1    GPU2    GPU3    GPU4    GPU5    GPU6    GPU7    NIC0    CPU Affinity    NUMA Affinity
GPU0     X      NV2     NV1     NV2     SYS     SYS     NV1     SYS     NODE    0-19    0
GPU1    NV2      X      NV2     NV1     SYS     SYS     SYS     NV1     NODE    0-19    0
GPU2    NV1     NV2      X      NV1     NV2     SYS     SYS     SYS     NODE    0-19    0
GPU3    NV2     NV1     NV1      X      SYS     NV2     SYS     SYS     NODE    0-19    0
GPU4    SYS     SYS     NV2     SYS      X      NV1     NV1     NV2     SYS     20-39   1
GPU5    SYS     SYS     SYS     NV2     NV1      X      NV2     NV1     SYS     20-39   1
GPU6    NV1     SYS     SYS     SYS     NV1     NV2      X      NV2     SYS     20-39   1
GPU7    SYS     NV1     SYS     SYS     NV2     NV1     NV2      X      SYS     20-39   1
NIC0    NODE    NODE    NODE    NODE    SYS     SYS     SYS     SYS      X
Legend:
X    = Self
SYS  = Connection traversing PCIe as well as the SMP interconnect between NUMA nodes (e.g., QPI/UPI)
NODE = Connection traversing PCIe as well as the interconnect between PCIe Host Bridges within a NUMA node
PHB  = Connection traversing PCIe as well as a PCIe Host Bridge (typically the CPU)
PXB  = Connection traversing multiple PCIe bridges (without traversing the PCIe Host Bridge)
PIX  = Connection traversing at most a single PCIe bridge
NV#  = Connection traversing a bonded set of # NVLinks
NIC Legend:
NIC0: mlx5_0
Thanks for your help!!!

got it. Thanks!
Lots of this:

Couple be that because the v100 is a volta device it doesn't support bf16. When using vLLM elsewhere I have been able to use fp16 instead of bf16 using the --dtype flag. Could this be related? And, is it possible to support fp16 if that's the issue here?

@K-Mistele Yes, you are right.  When using vLLM, you can specify fp16 to avoid this problem. This issue also exists on GPUs with the Turing architecture. You can test only the fp16  by delete torch.bfloat16 in DTYPES

@qniguogym  It's strange, why isn't your torch version 2.3.1, and triton isn't 2.3.1 either? If you used pip install -e . to build, it should have installed the corresponding versions of the modules.What's more,  you can refer to :https://github.com/vllm-project/vllm/issues/5547. We can also wait for @K-Mistele to complete the test, as his GPU is also the V100

@chenhongyu2048  Hi, Ori is without LoRA, LoRA is punica C, and LoRA(triton) is this PR's implementation

@mgoin  I noticed you're closing some issues. Actually, I'm quite curious whether you've actually tested these locally

@chenhongyu2048 Hi, Ori is without LoRA, LoRA is punica C, and LoRA(triton) is this PR's implementation

Thanks!

@jeejeelee Yes I've been doing some smoke testing on various lora adapters for Llama and Qwen models. For some of the issues the exact model+adapter details are not available so it is difficult to reproduce, but it is clear we have more flexible support now. In any case many of these issues have grown stale and need to be refreshed with the new kernels.

@mgoin  Thank you very much for your testing. I will continue to monitor the community feedback and keep working on optimizing the Lora-related tasks. Thanks again!

hello my friend, I'm using your code but I find the performance is not as good as expected:
Average triton_version time: 0.740561 milliseconds
Average old cuda version time: 0.154018 milliseconds
My code is as below:
import torch
import time
from bgmv_shrink import bgmv_shrink
from vllm.lora.punica import bgmv

num_tokens = 128
rank = 128
hidden_size = 4096

hidden_states = torch.randn(num_tokens, hidden_size, dtype=torch.float16, device=f'cuda:{0}')
weights = torch.randn(num_tokens, 1, rank, hidden_size, dtype=torch.float16, device=f'cuda:{0}')
output_buf = torch.zeros(num_tokens, rank, dtype=torch.float16, device=f'cuda:{0}')
indices = torch.arange(0, num_tokens, dtype=torch.long, device=f'cuda:{0}')

time_list = []
for i in range(100):
    torch.cuda.synchronize()
    t0 = time.time()

    bgmv_shrink(hidden_states, weights, output_buf, indices)

    torch.cuda.synchronize()
    t1 = time.time()
    time_list.append(t1 - t0)

avg_time = sum(time_list[50:]) / 50

print(f"Average triton_version time: {avg_time * 1000:.6f} milliseconds")

time_list = []
for i in range(100):
    torch.cuda.synchronize()
    t0 = time.time()

    bgmv(output_buf, hidden_states, weights, indices, 0, 1.0)

    torch.cuda.synchronize()
    t1 = time.time()
    time_list.append(t1 - t0)

avg_time = sum(time_list[50:]) / 50
print(f"Average old cuda version time: {avg_time * 1000:.6f} milliseconds")

I run it on RTX4090GPU with cuda12.1. And I guess the reason of my unexpected performance is the LoRA rank?

@chenhongyu2048  Hi, We alse provided a benchmark repo to evaluate the performance of our implementation compared to the original Punica C and https://github.com/vllm-project/vllm/pull/5356, please see: https://github.com/jeejeelee/punica_bgmv_testing

@chenhongyu2048 Hi, We also provided a benchmark repo to evaluate the performance of our implementation compared to the original Punica C and #5356, please see: https://github.com/jeejeelee/punica_bgmv_testing. The current implementation may have a slight performance gap compared to Punica C for small operators. This is a known issue, and we will optimize it in future iterations.

Hello, I have finish the benchmark test. And I find that when hidden_size is small (I want h=4096), punica C kernels did perform better. But when h grows, the result is reversed. Thanks for your work, it did help me.
And I guess this could because of the different ways to cut matrices? Perhaps an if-else decision might be used to determine the optimal performance. But I also look forward to have a more elegant solution.

Hello, first of all, thank you very much for providing the MutiLora code. I have successfully run the code on V100, but encountered some issues. I am using the Qwen2-7b model. When performing inference using the online merged LoRa and the offline merged LoRa weights, I found that there is an inconsistency in the accuracy. Using the online merged LoRa tends to provide answers that are more similar to the base model. Why is this the case?

@valium123  What version of Triton do you have? I remember someone previously reported an issue with misaligned results on V100, but upgrading to Triton 3.3.0 resolved the alignment problem.

@valium123 What version of Triton do you have? I remember someone previously reported an issue with misaligned results on V100, but upgrading to Triton 3.3.0 resolved the alignment problem.

Hello, I am currently using Triton version 2.3.1. I noticed that the latest version of Triton on Python seems to be 3.0.0.

@valium123 What version of Triton do you have? I remember someone previously reported an issue with misaligned results on V100, but upgrading to Triton 3.3.0 resolved the alignment problem.

Hello, I am currently using Triton version 2.3.1. I noticed that the latest version of Triton on Python seems to be 3.0.0.

Yes, you can try upgrading to Triton 3.0.0 first

@valium123 What version of Triton do you have? I remember someone previously reported an issue with misaligned results on V100, but upgrading to Triton 3.3.0 resolved the alignment problem.

Hello, I am currently using Triton version 2.3.1. I noticed that the latest version of Triton on Python seems to be 3.0.0.

Yes, you can try upgrading to Triton 3.0.0 first

yeah, Upgrading Triton to 3.0.0 has aligned the results with the offline merged weights. thanks

@valium123 What version of Triton do you have? I remember someone previously reported an issue with misaligned results on V100, but upgrading to Triton 3.3.0 resolved the alignment problem.

Hello, I am currently using Triton version 2.3.1. I noticed that the latest version of Triton on Python seems to be 3.0.0.

Yes, you can try upgrading to Triton 3.0.0 first

yeah, Upgrading Triton to 3.0.0 has aligned the results with the offline merged weights. thanks
Just tested MultiLora Triton==2.3.1 on the A10 machine, the results align properly. However, with 3.0.0, the results are abnormal. It seems that the dependencies for A-series and V-series are inconsistent.

@valium123 What version of Triton do you have? I remember someone previously reported an issue with misaligned results on V100, but upgrading to Triton 3.3.0 resolved the alignment problem.

Hello, I am currently using Triton version 2.3.1. I noticed that the latest version of Triton on Python seems to be 3.0.0.

Yes, you can try upgrading to Triton 3.0.0 first

yeah, Upgrading Triton to 3.0.0 has aligned the results with the offline merged weights. thanks
Just tested MultiLora Triton==2.3.1 on the A10 machine, the results align properly. However, with 3.0.0, the results are abnormal. It seems that the dependencies for A-series and V-series are inconsistent.

That shouldn't be the case. We conducted tests on both the 3090 and A800, and the results were aligned in both cases.

@valium123 What version of Triton do you have? I remember someone previously reported an issue with misaligned results on V100, but upgrading to Triton 3.3.0 resolved the alignment problem.

Hello, I am currently using Triton version 2.3.1. I noticed that the latest version of Triton on Python seems to be 3.0.0.

Yes, you can try upgrading to Triton 3.0.0 first

yeah, Upgrading Triton to 3.0.0 has aligned the results with the offline merged weights. thanks
Just tested MultiLora Triton==2.3.1 on the A10 machine, the results align properly. However, with 3.0.0, the results are abnormal. It seems that the dependencies for A-series and V-series are inconsistent.

That shouldn't be the case. We conducted tests on both the 3090 and A800, and the results were aligned in both cases.

@jeejeelee  Just conducted an investigation, and it seems unrelated to the Triton version. I tested on A10 machines with Triton versions 2.3.1 and 3.0.0, and there is still randomness in the results when 'temperature': 0, 'topK': 1. This phenomenon does not occur when using the base model, so it seems to be caused by online LoRa merging.

@valium123 How did you install vllm? Is it the latest version? If so, you can un the unit tests to verify it, please run pytest test_punica_sizes.py

The following test results were obtained when running on a machine with an A10 GPU using triton version 3.0.0:
================================================ short test summary info =================================================
FAILED test_punica_sizes.py::test_punica_sgmv[cuda:0-0-expand-dtype0-0.5-102656-32-4-4] - RuntimeError: Comparing
FAILED test_punica_sizes.py::test_punica_sgmv[cuda:0-0-expand-dtype0-0.5-128256-32-4-4] - RuntimeError: Comparing
FAILED test_punica_sizes.py::test_punica_sgmv[cuda:0-0-expand-dtype0-0.5-128000-32-4-4] - RuntimeError: Comparing
FAILED test_punica_sizes.py::test_punica_sgmv[cuda:0-0-expand-dtype1-0.5-102656-32-4-4] - RuntimeError: Comparing
FAILED test_punica_sizes.py::test_punica_sgmv[cuda:0-0-expand-dtype1-0.5-128256-32-4-4] - RuntimeError: Comparing
FAILED test_punica_sizes.py::test_punica_sgmv[cuda:0-0-expand-dtype1-0.5-128000-32-4-4] - RuntimeError: Comparing
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype0-2-64128-32-4-4] - RuntimeError: Comparing
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype0-2-60672-32-4-4] - RuntimeError: Comparing
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype0-2-60544-32-4-4] - RuntimeError: Comparing
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype0-2-102400-32-4-4] - torch.cuda.OutOfMemoryError: CUDA out of memory. Tried to allocate 26.00 MiB. GPU
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype0-2-51328-32-4-4] - RuntimeError: Comparing
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype0-2-64256-32-4-4] - RuntimeError: Comparing
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype0-2-102656-32-4-4] - torch.cuda.OutOfMemoryError: CUDA out of memory. Tried to allocate 402.00 MiB. GPU
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype0-2-128256-32-4-4] - torch.cuda.OutOfMemoryError: CUDA out of memory. Tried to allocate 126.00 MiB. GPU
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype0-2-5888-32-4-4] - RuntimeError: Comparing
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype0-2-128000-32-4-4] - torch.cuda.OutOfMemoryError: CUDA out of memory. Tried to allocate 126.00 MiB. GPU
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype0-2-64000-32-4-4] - RuntimeError: Comparing
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype0-3-49152-32-4-4] - RuntimeError: Comparing
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype0-3-51200-32-4-4] - RuntimeError: Comparing
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype0-3-43264-32-4-4] - RuntimeError: Comparing
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype0-3-64128-32-4-4] - torch.cuda.OutOfMemoryError: CUDA out of memory. Tried to allocate 376.00 MiB. GPU
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype0-3-60672-32-4-4] - torch.cuda.OutOfMemoryError: CUDA out of memory. Tried to allocate 356.00 MiB. GPU
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype0-3-36864-32-4-4] - RuntimeError: Comparing
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype0-3-60544-32-4-4] - torch.cuda.OutOfMemoryError: CUDA out of memory. Tried to allocate 60.00 MiB. GPU
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype0-3-102400-32-4-4] - torch.cuda.OutOfMemoryError: CUDA out of memory. Tried to allocate 300.00 MiB. GPU
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype0-3-51328-32-4-4] - RuntimeError: Comparing
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype0-3-64256-32-4-4] - torch.cuda.OutOfMemoryError: CUDA out of memory. Tried to allocate 378.00 MiB. GPU
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype0-3-49408-32-4-4] - RuntimeError: Comparing
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype0-3-102656-32-4-4] - torch.cuda.OutOfMemoryError: CUDA out of memory. Tried to allocate 302.00 MiB. GPU
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype0-3-128256-32-4-4] - torch.cuda.OutOfMemoryError: CUDA out of memory. Tried to allocate 376.00 MiB. GPU
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype0-3-128000-32-4-4] - torch.cuda.OutOfMemoryError: CUDA out of memory. Tried to allocate 376.00 MiB. GPU
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype0-3-64000-32-4-4] - torch.cuda.OutOfMemoryError: CUDA out of memory. Tried to allocate 376.00 MiB. GPU
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype1-2-64128-32-4-4] - RuntimeError: Comparing
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype1-2-60672-32-4-4] - RuntimeError: Comparing
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype1-2-60544-32-4-4] - RuntimeError: Comparing
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype1-2-102400-32-4-4] - torch.cuda.OutOfMemoryError: CUDA out of memory. Tried to allocate 26.00 MiB. GPU
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype1-2-51328-32-4-4] - RuntimeError: Comparing
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype1-2-64256-32-4-4] - RuntimeError: Comparing
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype1-2-102656-32-4-4] - torch.cuda.OutOfMemoryError: CUDA out of memory. Tried to allocate 402.00 MiB. GPU
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype1-2-128256-32-4-4] - torch.cuda.OutOfMemoryError: CUDA out of memory. Tried to allocate 126.00 MiB. GPU
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype1-2-5888-32-4-4] - torch.cuda.OutOfMemoryError: CUDA out of memory. Tried to allocate 24.00 MiB. GPU
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype1-2-128000-32-4-4] - torch.cuda.OutOfMemoryError: CUDA out of memory. Tried to allocate 126.00 MiB. GPU
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype1-2-64000-32-4-4] - RuntimeError: Comparing
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype1-3-49152-32-4-4] - RuntimeError: Comparing
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype1-3-51200-32-4-4] - RuntimeError: Comparing
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype1-3-43264-32-4-4] - RuntimeError: Comparing
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype1-3-64128-32-4-4] - torch.cuda.OutOfMemoryError: CUDA out of memory. Tried to allocate 376.00 MiB. GPU
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype1-3-60672-32-4-4] - torch.cuda.OutOfMemoryError: CUDA out of memory. Tried to allocate 356.00 MiB. GPU
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype1-3-36864-32-4-4] - RuntimeError: Comparing
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype1-3-60544-32-4-4] - torch.cuda.OutOfMemoryError: CUDA out of memory. Tried to allocate 60.00 MiB. GPU
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype1-3-102400-32-4-4] - torch.cuda.OutOfMemoryError: CUDA out of memory. Tried to allocate 300.00 MiB. GPU
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype1-3-51328-32-4-4] - RuntimeError: Comparing
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype1-3-64256-32-4-4] - torch.cuda.OutOfMemoryError: CUDA out of memory. Tried to allocate 378.00 MiB. GPU
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype1-3-49408-32-4-4] - RuntimeError: Comparing
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype1-3-102656-32-4-4] - torch.cuda.OutOfMemoryError: CUDA out of memory. Tried to allocate 302.00 MiB. GPU
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype1-3-128256-32-4-4] - torch.cuda.OutOfMemoryError: CUDA out of memory. Tried to allocate 376.00 MiB. GPU
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype1-3-128000-32-4-4] - torch.cuda.OutOfMemoryError: CUDA out of memory. Tried to allocate 376.00 MiB. GPU
FAILED test_punica_sizes.py::test_punica_expand_nslices[cuda:0-0-sgmv-dtype1-3-64000-32-4-4] - torch.cuda.OutOfMemoryError: CUDA out of memory. Tried to allocate 376.00 MiB. GPU
================================= 58 failed, 4038 passed, 1 warning in 712.15s (0:11:52) =================================
There are even more errors on the V100

@valium123 You should take steps to avoid OOM errors. Also, please note that the V100 does not support bf16, see the above comments

what is the conditions that run on V100?
Only support GPU which is compute capability >= 8.0
I saw the criteria is gpu >=8.0
but post above side can run on v100.
so  what is the criteria cam make it run on v100?
Thank you

I saw the criteria is gpu >=8.0

Where do you see these criteria?
In my opinion, if you want to run LoRA on  V100, you'd better use the latest vllm (0.5.4) and use torch.float16.

I saw the criteria is gpu >=8.0

Where do you see these criteria? In my opinion, if you want to run LoRA on V100, you'd better use the latest vllm (0.5.4) and use torch.float16.

from your open post section "motivation"
Only support GPU which is compute capability >= 8.0
or do i mis-understood something?

I saw the criteria is gpu >=8.0

Where do you see these criteria? In my opinion, if you want to run LoRA on V100, you'd better use the latest vllm (0.5.4) and use torch.float16.

from your open post section "motivation" Only support GPU which is compute capability >= 8.0 or do i mis-understood something?

These were previous issues with Punica. This pull request addresses some of them

I saw the criteria is gpu >=8.0

Where do you see these criteria? In my opinion, if you want to run LoRA on V100, you'd better use the latest vllm (0.5.4) and use torch.float16.

from your open post section "motivation" Only support GPU which is compute capability >= 8.0 or do i mis-understood something?

These were previous issues with Punica. This pull request addresses some of them

ok. it is my mis-understood.  Thanks for your clarify.

Hi @Yard1 @jeejeelee - I am currently going through the install process on my v100 and will let you know if I run into the same errors. Can you let me know how I can test the implementation at runtime (not just compile-time) so that I can verify compatibility? e.g. what model, config, and LoRAs are you using for testing, or can I run the existing tests somehow? I apologize as I'm not as familiar with the project's test suite.

Thank you for being willing to spend time verifying on the V100. If your installation is successful, you can proceed with:

Kernel testing, which can be conducted using test_script

Model testing, which can be conducted using the following  scripts:

https://github.com/vllm-project/vllm/blob/main/tests/lora/test_llama.py
https://github.com/vllm-project/vllm/blob/main/examples/multilora_inference.py

Please feel free to contact me if you have any questions.

@Yard1 @jeejeelee  - I did manage to successfully run the test scripts at tests/lora/test_punica_sizes and tests/lora/test_llama.py on my Nvidia v100 Tesla 32GB (volta architecture).
To get the test_punica_sizes.py test to run, I did have to disable BF16 support on line 112. For anyone else on a v100, this data type is not supported on the volta architecture. So, it seems y'all were successful in fixing that compatibility issue :)
For anyone on a v100 who has having issues, make sure to:

disable bfloat16 on line 112 of tests/lora/text_puniza_sizes.py
pull main, create a clean conda environment (python=3.10 works great), and build from source using pip install -e . followed by installing testing dependencies at pip install -r requirements-test.py -- otherwise you may run into triton version issues.

@K-Mistele  Thank you for your testing, I really appreciate it. By the way, the current vLLM team has released version 0.5.4, and this PR was integrated into it. You can run pip install vllm directly.

@jeejeelee the latest VLLM version (0.5.4) does not support lora_rank> 64

@jeejeelee the latest VLLM version (0.5.4) does not support lora_rank> 64

#7146  had addressed this issue. However, you'll have to wait until version 0.5.5 for support, or you can build  it from the source

Got it. Thanks for your work. @jeejeelee

question1:
how to understand the following description?
1.Only support GPU which is compute capability >= 8.0
2.I have completed the integration, and unit testing of the Triton kernel according to my approach. TP（2/4/8）/CUDA Graph/Quantization/Fully sharded LoRA/NVIDIA T4,
T4's capability >= 7.5
question2:
in T4 and vllm version==0.4.3
deploy multi lora in vllm failed, error info:
RuntimeError: CUDA error: no kernel image is available for execution on the device.
my deploy command:
CUDA_VISIBLE_DEVICES=0,1,2,3 swift deploy --tensor_parallel_size 4 --dtype fp16 --model_type qwen1half-7b-chat
--model_id_or_path /cloud/user/data/data0806/llm/M2/Chat_New
--ckpt_dir /cloud/user/data/data0806/llm/M2/checkpoint-200/
--infer_back vllm -- vllm_enable_lora true --max_model_len 512 --enforce_eager
I tried to vllm to 0.5.5, there is still error

question1:
how to understand the following description?
1.Only support GPU which is compute capability >= 8.0
2.I have completed the integration, and unit testing of the Triton kernel according to my approach. TP（2/4/8）/CUDA Graph/Quantization/Fully sharded LoRA/NVIDIA T4,
T4's capability >= 7.5
question2:
in T4 and vllm version==0.4.3
deploy multi lora in vllm failed, error info:
RuntimeError: CUDA error: no kernel image is available for execution on the device.
my deploy command:
CUDA_VISIBLE_DEVICES=0,1,2,3 swift deploy --tensor_parallel_size 4 --dtype fp16 --model_type qwen1half-7b-chat
--model_id_or_path /cloud/user/data/data0806/llm/M2/Chat_New
--ckpt_dir /cloud/user/data/data0806/llm/M2/checkpoint-200/
--infer_back vllm -- vllm_enable_lora true --max_model_len 512 --enforce_eager
I tried to vllm to 0.5.5, there is still error
@jeejeelee

@askcs517

question1

These were previous issues with Punica. This pull request addresses some of them.  Please carefully read the initial description again

question2

It's unclear what error you encountered. You should also consider submitting an issue.

您好，您的邮件我已经收到，我会及时查收。麻烦您了。

Thank you for the excellent work!
I'm checking the triton kernels ,and I have two questions:

There are 6 kernels, why there are so many ones? bgmv/sgmv  expand/expand_slice, what's the difference between them?
Is the kernel already tuned? I didn't find a config file for each platform.

Thank you.

您好，您的邮件我已经收到，我会及时查收。麻烦您了。

Thank you for the excellent work! @jeejeelee
I'm checking the triton kernels ,and I have two questions:

There are 6 kernels, why there are so many ones? bgmv/sgmv expand/expand_slice, what's the difference between them?
Is the kernel already tuned? I didn't find a config file for each platform.

Thank you.

Q1:
The shrink is used for calculating lora_a, while the expand  is used for calculating lora_b.
We also differentiate between the prefill and decode stage: the prefill stage uses the sgmv kernel , while the decode stage uses the bgmv kernel. The expand_slice is applicable to qkv linear and ffn linear operations in networks similar to LLaMA.
Q2:
not yet

Thank you for the excellent work! @jeejeelee
I'm checking the triton kernels ,and I have two questions:

There are 6 kernels, why there are so many ones? bgmv/sgmv expand/expand_slice, what's the difference between them?
Is the kernel already tuned? I didn't find a config file for each platform.

Thank you.

Q1: The shrink is used for calculating lora_a, while the expand is used for calculating lora_b.
We also differentiate between the prefill and decode stage: the prefill stage uses the sgmv kernel , while the decode stage uses the bgmv kernel. The expand_slice is applicable to qkv linear and ffn linear operations in networks similar to LLaMA.
Q2: not yet

Hi, I wonder what params should be tuned?
I think the block_K is optimal, because rank is small, so we can read it into SHM in a block once.
And as for the default config, it seems the configuration is already good enough?
Do you have any suggestions for tuning the kernels ? Thank you !
        return {
            "BLOCK_N": 256,
            "SPLIT_N": _check_divisibility(hidden_size),
            "num_warps": 8
        }

    # TODO tuning this config
    N, K = lora_b_weights.shape[-2:]  # K= rank,N=hidden_size
    BLOCK_K = triton.next_power_of_2(K)
    EVEN_K = K % BLOCK_K == 0
    ADD_INPUTS = add_inputs
    CAST_TYPE = False
    if inputs.dtype == torch.float32 and lora_b_weights.dtype in [
            torch.float16,
            torch.bfloat16,
    ]:
        CAST_TYPE = True
    batches = lora_indices_tensor.size(0)
    config = get_lora_op_configs("expand", batches, N)
    grid = lambda META: (
        META["SPLIT_N"],
        batches,
    )
    _bgmv_expand_kernel[grid](
        inputs,
        lora_b_weights,
        output_tensor,
        N,
        K,
        lora_indices_tensor,
        inputs.stride(0),
        inputs.stride(1),
        lora_b_weights.stride(0),
        lora_b_weights.stride(1),
        lora_b_weights.stride(2),
        output_tensor.stride(0),
        output_tensor.stride(1),
        BLOCK_K=BLOCK_K,
        EVEN_K=EVEN_K,
        ADD_INPUTS=ADD_INPUTS,
        CAST_TYPE=CAST_TYPE,
        **config,
    )

@sleepwalker2017  See: https://github.com/jeejeelee/punica_triton_kernel/blob/main/benchmark_bgmv.py

We also differentiate between the prefill and decode stage: the prefill stage uses the sgmv kernel , while the decode stage uses the bgmv kernel. The expand_slice is applicable to qkv linear and ffn linear operations in networks similar to LLaMA.

The tuning program is fast and efficient. Does the config fit for sgmv kernels?
Or if I tune bgmv kernels for batch <= 4096, do they apply for sgmv kernels with sequence length <= 4096?  Thank you.

@sleepwalker2017 See: https://github.com/jeejeelee/punica_triton_kernel/blob/main/benchmark_bgmv.py

Sorry for so many questions, but do we need to tune for different num_stages?

The tuning program is fast and efficient. Does the config fit for sgmv kernels?

The logic of tuning is the same, but some things may need to be modified.

but do we need to tune for different num_stages?

Perhaps we need.

Hi, I have another question, for qkv_proj, why do we calculate the lora for them separately?
In my view, we can fuse the loras for qkv, and do the bgmv in a single kernel. Is that the fact?

Hi, I have another question, for qkv_proj, why do we calculate the lora for them separately? In my view, we can fuse the loras for qkv, and do the bgmv in a single kernel. Is that the fact?

Yes, you are right. In fact, we aslo can fuse lora_a and lora_b into one kernel

Hi,  why does shrink kernel uses tl.float32 as output dtype? I notice the bgmv Slora kernel uses the same dtype for input and output.

Hi, why does shrink kernel uses tl.float32 as output dtype? I notice the bgmv Slora kernel uses the same dtype for input and output.

Due to the shrink kernel utilizing atomic add  and Triton's lack of support for BF16 atomic add , see: https://github.com/triton-lang/triton/issues/1387

triton-lang/triton#1387

got it. so it's ok to use 16-bit result when using float16. Thank you.

Hi, I find we can delete the expand_slice kernels, since we can slice the tensor before calling the expand kernel.
I did the following changes, the result has no difference. Is that right?
        self.add_shrink(buffer, x, wa_t_all, scale)
        if y_offset is None and y_slice_size is None:
            self.add_expand(y, buffer, wb_t_all, add_input=True)
        else:
            tmp = torch.clone(y)
            self.add_expand_slice(y,
                                  buffer,
                                  wb_t_all,
                                  y_offset,
                                  y_slice_size,
                                  add_input=True)

            self.add_expand(**tmp[:, y_offset:y_offset+y_slice_size]**, buffer, wb_t_all, add_input=True)
            diff = torch.abs(tmp - y)
            print(diff.max().item(), "diff")
            if diff.max().item() > 0:
                import pdb; pdb.set_trace()

Hope to get your confirmation.

Hi, I find we can delete the expand_slice kernels, since we can slice the tensor before calling the expand kernel.
I did the following changes, the result has no difference. diff.max() is always 0. Is that right?
        self.add_shrink(buffer, x, wa_t_all, scale)
        if y_offset is None and y_slice_size is None:
            self.add_expand(y, buffer, wb_t_all, add_input=True)
        else:
            tmp = torch.clone(y)
            self.add_expand_slice(y,
                                  buffer,
                                  wb_t_all,
                                  y_offset,
                                  y_slice_size,
                                  add_input=True)

            self.add_expand(tmp[:, y_offset:y_offset+y_slice_size], buffer, wb_t_all, add_input=True)
            diff = torch.abs(tmp - y)
            print(diff.max().item(), "diff")
            if diff.max().item() > 0:
                import pdb; pdb.set_trace()

Hope to get your confirmation.

Yes, it seems reasonable based on initial review. However, it doesn't help much with current performance. Could we go further and only call the shrink and expand op once for the lora layers that call nslice?

Hi, I find we can delete the expand_slice kernels, since we can slice the tensor before calling the expand kernel.
I did the following changes, the result has no difference. diff.max() is always 0. Is that right?
        self.add_shrink(buffer, x, wa_t_all, scale)
        if y_offset is None and y_slice_size is None:
            self.add_expand(y, buffer, wb_t_all, add_input=True)
        else:
            tmp = torch.clone(y)
            self.add_expand_slice(y,
                                  buffer,
                                  wb_t_all,
                                  y_offset,
                                  y_slice_size,
                                  add_input=True)

            self.add_expand(tmp[:, y_offset:y_offset+y_slice_size], buffer, wb_t_all, add_input=True)
            diff = torch.abs(tmp - y)
            print(diff.max().item(), "diff")
            if diff.max().item() > 0:
                import pdb; pdb.set_trace()

Hope to get your confirmation.

Yes, it seems reasonable based on initial review. However, it doesn't help much with current performance. Could we go further and only call the shrink and expand op once for the lora layers that call nslice?

Yes, I'm working on the performance of the lora kernels, and I'll try that.