[Bugfix] Update profile example to new add request interface + fix profiler not picking up kernels within cudagraphs (#332) This PR updates `examples/offline_profile.py` to use the new `add_requests` interface. In addition, this PR fixes issues with profiler not picking up kernels run from within a cudagraph (i.e. when profiling with `--allow-cuda-graphs`, there's is too main issues: 1) Changes from the initial profiler PR (https://github.com/neuralmagic/nm-vllm/pull/124) were wiped out by https://github.com/neuralmagic/nm-vllm/pull/224, namely the changes in `model_runner.py` converting `CUDAGraphRunner` to a `nn.Module` allowing the profiler to pick it up 2) Many kernels within the graph had the same correlation id so we were always picking the first of potentially many kernels to display, using name in addition to correlation Id sees to resolve this issue but is potentially fragile Before the PR: ``` ================================================================================ = Decode Summary Table (prompt_len=1, batch_size=1) ================================================================================ name | cuda_time_us | pct_cuda_... | invocations ================================================================================================================================ LogitsProcessor | 350.00 | 57.76 | 1.00 |- void at::native::(anonymous namespace)::indexSelectSmallIndex Device) | 18.00 | 2.97 | 9.00 |- void at::native::elementwise_kernel<128, 4, at::native::gpu_kernel_impl_no... | 4.00 | 0.66 | 1.00 |- void at::native::unrolled_elementwise_kernel Pageable) | 10.00 | 1.65 | 5.00 |- void (anonymous namespace)::elementwise_kernel_with_index(unsigned int*, u... | 1.00 | 0.17 | 1.00 |- void at::native::mbtopk::radixFindKthValues(unsig... | 8.00 | 1.32 | 4.00 |- void at::native::mbtopk::computeBlockwiseKthCounts(unsigned ... | 1.00 | 0.17 | 1.00 |- void at_cuda_detail::cub::DeviceScanByKeyInitKernel(at::cuda::deta... | 4.00 | 0.66 | 1.00 ``` After PR ``` name | cuda_time_us | pct_cuda_... | invocations ================================================================================================================================ CUDAGraphRunner | 4238.00 | 84.41 | 1.00 |- Memcpy DtoD (Device -> Device) | 5.00 | 0.10 | 5.00 |- void at::native::vectorized_elementwise_kernel<4, at::native::FillFunctor<... | 2.00 | 0.04 | 2.00 |- void at::native::(anonymous namespace)::indexSelectSmallIndex(c10::BFloat16*, c10::BFloat16 co... | 4.00 | 0.08 | 1.00 |- void vllm::scaled_fp8_quant_kernel(c10::Float8_e4m3fn*, c10... | 256.00 | 5.10 | 128.00 |- sm90_xmma_gemm_e4m3bf16_e4m3f32_f32_tn_n_tilesize64x64x128_warpgroupsize1x... | 1440.00 | 28.68 | 96.00 |- void vllm::rotary_embedding_kernel(long const*, c10::... | 96.00 | 1.91 | 32.00 |- void vllm::reshape_and_cache_flash_kernel(c10::BFloat16 con... | 64.00 | 1.27 | 32.00 |- void flash_fwd_splitkv_kernel(0)))&&vllm::_typeConvert::exists, voi... | 128.00 | 2.55 | 64.00 |- sm90_xmma_gemm_e4m3bf16_e4m3f32_f32_tn_n_tilesize64x128x128_warpgroupsize1... | 1664.00 | 33.14 | 32.00 |- void vllm::act_and_mul_kernel Device) | 18.00 | 0.36 | 9.00 |- void at::native::elementwise_kernel<128, 4, at::native::gpu_kernel_impl_no... | 4.00 | 0.08 | 1.00 |- at::native::(anonymous namespace)::fill_reverse_indices_kernel(long*, int,... | 2.00 | 0.04 | 1.00 |- void at_cuda_detail::cub::DeviceRadixSortUpsweepKernel Device) | 1.00 | 0.02 | 1.00 |- void at::native::unrolled_elementwise_kernel... | 1.00 | 0.02 | 1.00 |- void (anonymous namespace)::elementwise_kernel_with_index Pageable) | 10.00 | 0.20 | 5.00 |- void at::native::elementwise_kernel<128, 4, at::native::gpu_kernel_impl_no... | 3.00 | 0.06 | 1.00 |- void at::native::reduce_kernel<512, 1, at::native::ReduceOp(unsigned int*, u... | 1.00 | 0.02 | 1.00 |- void at::native::mbtopk::radixFindKthValues(unsig... | 8.00 | 0.16 | 4.00 |- void at::native::mbtopk::computeBlockwiseKthCounts(unsigned ... | 1.00 | 0.02 | 1.00 |- void at_cuda_detail::cub::DeviceScanByKeyInitKernel(at::cuda::deta... | 4.00 | 0.08 | 1.00 ``` --------- Co-authored-by: Lucas Wilkinson

Uh oh!

feature request: support mpt-30b #332

Description

Metadata

Metadata

Assignees

Labels

Type

Projects

Milestone

Relationships

Development

Issue actions