Gahow Wang
445e491123
third_party/vllm/ now tracked in git for direct patch management.
Based on vLLM v0.18.1 release with one patch applied:
vllm/v1/core/sched/scheduler.py:
Replace fatal assert with graceful skip when KV transfer callback
arrives for an already-aborted request during PD disaggregated serving.
Future vLLM modifications should be made directly in third_party/vllm/
and committed normally. The patches/ directory is kept as documentation
of what changed from upstream.
Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
43 lines
1.8 KiB
Plaintext
43 lines
1.8 KiB
Plaintext
#pragma once
|
|
|
|
#include "cutlass/gemm/collective/collective_builder.hpp"
|
|
|
|
namespace cutlass::gemm::collective {
|
|
using namespace cute;
|
|
|
|
//
|
|
// VLLMCollectiveBuilder is a wrapper around CollectiveBuilder that allows for
|
|
// for custom kernel tags, allowing you to build custom collectives. Without
|
|
// touching the cutlass library headers, using `CutlassKernelTag` will mean it
|
|
// will resort to using the standard cutlass collective builder.
|
|
//
|
|
|
|
// Use the default Cutlass collective builder, i.e. use an unmodified cutless
|
|
// collective
|
|
struct CutlassKernelTag {};
|
|
|
|
template <class KernelTag, class ArchTag, class OpClass, class ElementA,
|
|
class GmemLayoutA, int AlignmentA, class ElementB, class GmemLayoutB,
|
|
int AlignmentB, class ElementAccumulator, class TileShape_MNK,
|
|
class ClusterShape_MNK, class StageCountType,
|
|
class KernelScheduleType, class Enable = void>
|
|
struct VLLMCollectiveBuilder {
|
|
static_assert(sizeof(ElementA) == 0,
|
|
"Could not build a collective for given parameters.");
|
|
};
|
|
|
|
template <class ArchTag, class OpClass, class ElementA, class GmemLayoutA,
|
|
int AlignmentA, class ElementB, class GmemLayoutB, int AlignmentB,
|
|
class ElementAccumulator, class TileShape_MNK, class ClusterShape_MNK,
|
|
class StageCountType, class KernelScheduleType>
|
|
struct VLLMCollectiveBuilder<
|
|
CutlassKernelTag, ArchTag, OpClass, ElementA, GmemLayoutA, AlignmentA,
|
|
ElementB, GmemLayoutB, AlignmentB, ElementAccumulator, TileShape_MNK,
|
|
ClusterShape_MNK, StageCountType, KernelScheduleType> {
|
|
using CollectiveOp = typename CollectiveBuilder<
|
|
ArchTag, OpClass, ElementA, GmemLayoutA, AlignmentA, ElementB,
|
|
GmemLayoutB, AlignmentB, ElementAccumulator, TileShape_MNK,
|
|
ClusterShape_MNK, StageCountType, KernelScheduleType>::CollectiveOp;
|
|
};
|
|
|
|
}; // namespace cutlass::gemm::collective |