update trtllm gemm

Signed-off-by: Xiwen Yu <13230610+VALLIS-NERIA@users.noreply.github.com>
This commit is contained in:
Xiwen Yu 2025-09-05 23:56:24 +08:00
parent 2c3f4cbeee
commit f8864b9061
150 changed files with 3534 additions and 1244 deletions

View File

@ -449,7 +449,7 @@ using namespace cutlass::epilogue;
using KernelScheduleSM10x = std::conditional_t<IsSM103FP4, KernelScheduleSM103, KernelScheduleSM100>; \
\
using KernelScheduleSM120 = cutlass ::gemm ::collective::KernelScheduleAuto; \
using KernelScheduleBW = std::conditional_t<IsSM120, KernelScheduleSM120, KernelScheduleSM100>; \
using KernelScheduleBW = std::conditional_t<IsSM120, KernelScheduleSM120, KernelScheduleSM10x>; \
\
using KernelSchedule = std::conditional_t<IsBlackwell, KernelScheduleBW, KernelScheduleSM90>; \
\

View File

@ -148,8 +148,8 @@ void dispatchMoeGemmFinalDispatchTmaWarpSpecialized(TmaWarpSpecializedGroupedGem
"passing 103-real as an arch to build_wheel.py."*/);
first_time = false;
}
return dispatchMoeGemmSelectBiasTmaWarpSpecialized<cutlass::arch::Sm100, T, WeightType, OutputType, EpilogueTag,
FUSION, TileShape, ClusterShape>(
return dispatchMoeGemmFinalDispatchTmaWarpSpecialized<cutlass::arch::Sm100, T, WeightType, OutputType,
EpilogueTag, FUSION, TileShape, ClusterShape>(
hopper_input, num_experts, multi_processor_count, stream, occupancy, workspace_size);
}
// #endif

View File

@ -18,6 +18,7 @@
#include "KernelRunner.h"
#include "tensorrt_llm/common/assert.h"
#include "tensorrt_llm/common/cudaUtils.h"
#include "tensorrt_llm/common/envUtils.h"
#include "trtllmGen_gemm_export/GemmInterface.h"
#include "trtllmGen_gemm_export/GemmOptions.h"
@ -33,6 +34,23 @@ using namespace gemm::gemm;
static GemmInterface::ModuleCache globalTrtllmGenGemmModuleCache;
constexpr bool isSMCompatible(int gpuSM, SmVersion kernelSM)
{
if (gpuSM == 103)
{
return kernelSM == SmVersion::Sm103a || kernelSM == SmVersion::Sm100f;
}
else if (gpuSM == 100)
{
return kernelSM == SmVersion::Sm100a || kernelSM == SmVersion::Sm100f;
}
else if (gpuSM == 90)
{
return kernelSM == SmVersion::Sm90a;
}
return true;
}
TrtllmGenGemmRunner::TrtllmGenGemmRunner(TrtllmGenGemmRunnerOptions const& options_)
: mOptions(options_)
{
@ -41,7 +59,7 @@ TrtllmGenGemmRunner::TrtllmGenGemmRunner(TrtllmGenGemmRunnerOptions const& optio
auto const configs = gemm.getGemmConfigs();
mPassingConfigIndices.clear();
int gpuNativeSmVersion = tensorrt_llm::common::getSMVersion();
for (size_t i = 0; i < gemm.getNumGemmConfigs(); ++i)
{
auto const options = configs[i].mOptions;
@ -50,7 +68,8 @@ TrtllmGenGemmRunner::TrtllmGenGemmRunner(TrtllmGenGemmRunnerOptions const& optio
if (options.mDtypeA == mOptions.eltTypeA && options.mDtypeC == mOptions.outputType
&& options.mUseDeepSeekFp8 == mOptions.deepSeekFp8
&& options.mTransposeMmaOutput == mOptions.transposeMmaOutput
&& (mOptions.eltTypeB == gemm::trtllm::gen::Dtype::Void || options.mDtypeB == mOptions.eltTypeB))
&& (mOptions.eltTypeB == gemm::trtllm::gen::Dtype::Void || options.mDtypeB == mOptions.eltTypeB)
&& isSMCompatible(gpuNativeSmVersion, configs[i].mSm))
{
mPassingConfigIndices.push_back(i);
}

View File

@ -109,12 +109,11 @@ struct GemmOptions
MatrixLayout layoutA, MatrixLayout layoutB, int m, int mmaK, tg::MmaKind mmaKind, int mmaM, int mmaN,
bool mockAllReduce, int n, int numSlicesForSplitK, int numSlicesForSliceK, int numStages, int numStagesMma,
int numStagesMmaWithinWorkTile, int numStagesMmaAcrossWorkTile, int numStagesWorkId, bool outputDebugTensors,
bool patchF2fp, bool useShuffledMatrixA, bool sliceK, SplitK splitK, bool transposeMmaOutput, int tileM,
int tileN, int tileK, bool useUnrollLoop2xForMma, bool useCustomMmaSchedule,
bool useHoistTryWaitForCustomMmaSchedule, bool useDeepSeekFp8, bool usePerTokenSfA, bool usePerTokenSfB,
bool useTmaStore, bool useTwoTmaLoadWarps, bool useTwoMmaWarps, std::optional<int32_t> sfBlockSizeA,
tg::SfLayout sfLayoutA, tg::SfLayout sfLayoutB, tg::SfLayout sfLayoutC, int sfReshapeFactor,
TileScheduler tileScheduler)
bool patchF2fp, std::optional<int32_t> sfBlockSizeA, tg::SfLayout sfLayoutA, tg::SfLayout sfLayoutB,
tg::SfLayout sfLayoutC, int sfReshapeFactor, bool sliceK, SplitK splitK, int tileK, int tileM, int tileN,
TileScheduler tileScheduler, bool transposeMmaOutput, bool useCustomMmaSchedule, bool useDeepSeekFp8,
bool useHoistTryWaitForCustomMmaSchedule, bool usePerTokenSfA, bool usePerTokenSfB, bool useShuffledMatrixA,
bool useTmaStore, bool useTwoTmaLoadWarps, bool useTwoMmaWarps, bool useUnrollLoop2xForMma, int worldSize)
: mAllReduceAlgo{allReduceAlgo}
, mBiasType{biasType}
, mBlockK(blockK)
@ -161,27 +160,16 @@ struct GemmOptions
, mNumStagesWorkId{numStagesWorkId}
, mOutputDebugTensors{outputDebugTensors}
, mPatchF2fp{patchF2fp}
, mUseShuffledMatrixA{useShuffledMatrixA}
, mSliceK{sliceK}
, mSplitK{splitK}
, mTransposeMmaOutput{transposeMmaOutput}
, mTileM{tileM}
, mTileN{tileN}
, mTileK{tileK}
, mUseUnrollLoop2xForMma{useUnrollLoop2xForMma}
, mUseCustomMmaSchedule{useCustomMmaSchedule}
, mUseHoistTryWaitForCustomMmaSchedule{useHoistTryWaitForCustomMmaSchedule}
, mUseDeepSeekFp8{useDeepSeekFp8}
, mUsePerTokenSfA{usePerTokenSfA}
, mUsePerTokenSfB{usePerTokenSfB}
, mUseTmaStore{useTmaStore}
, mUseTwoTmaLoadWarps{useTwoTmaLoadWarps}
, mUseTwoMmaWarps{useTwoMmaWarps}
, mSfBlockSizeA{sfBlockSizeA}
, mSfLayoutA{sfLayoutA}
, mSfLayoutB{sfLayoutB}
, mSfLayoutC{sfLayoutC}
, mSfReshapeFactor{sfReshapeFactor}
, mSliceK{sliceK}
, mSplitK{splitK}
, mTileK{tileK}
, mTileM{tileM}
, mTileN{tileN}
, mTileScheduler{tileScheduler}
, mTransposeMmaOutput{transposeMmaOutput}
, mUseCustomMmaSchedule{useCustomMmaSchedule}
@ -302,40 +290,6 @@ struct GemmOptions
bool mOutputDebugTensors{false};
// Patch float conversions.
bool mPatchF2fp{false};
// Reorder rows/cols in the A matrix for the better memory accesses in the M-major epilogue.
bool mUseShuffledMatrixA{false};
// Slice-K implementation to use TileM dimension for TileK.
bool mSliceK{false};
// The location of the exchange for split-K (it's None when split-K is disabled).
SplitK mSplitK{SplitK::None};
// Save output of MMA in M-major format.
bool mTransposeMmaOutput{false};
// M tile dimension of GEMM.
int mTileM{128};
// N tile dimension of GEMM.
int mTileN{32};
// K tile dimension of GEMM.
int mTileK{16};
// Whether to unroll the loop by 2x.
bool mUseUnrollLoop2xForMma{true};
// Use custom MMA schedule optimized for low-latency.
bool mUseCustomMmaSchedule{false};
// The purpose of hoisting trywaits is to opportunistically peek at the availability of the next
// k-block. It benefits when the next k-block is already available and thus sustaining the
// momentum, but it adds latency to the first k-block for smaller k-loop.
bool mUseHoistTryWaitForCustomMmaSchedule{false};
// Use DeepSeek Fp8.
bool mUseDeepSeekFp8{false};
// Apply per-token scales from A
bool mUsePerTokenSfA{false};
// Apply per-token scales from B
bool mUsePerTokenSfB{false};
// Use TMA to store the result.
bool mUseTmaStore{true};
// Use two different warps for A and B matrix load.
bool mUseTwoTmaLoadWarps{false};
// Use two different warps for MMA tasks. Applicable only to DeepSeek FP8.
bool mUseTwoMmaWarps{false};
// Block size of A. For dtypeA == E2m1 and dtypeB == E4m3.
std::optional<int32_t> mSfBlockSizeA{std::nullopt};
// Scale factors layout for A.
@ -350,6 +304,16 @@ struct GemmOptions
// But it reduces the number of L2 requests under the hood and potentially improves perf.
// Applies to layout 8x4 only.
int mSfReshapeFactor{1};
// Slice-K implementation to use TileM dimension for TileK.
bool mSliceK{false};
// The location of the exchange for split-K (it's None when split-K is disabled).
SplitK mSplitK{SplitK::None};
// K tile dimension of GEMM.
int mTileK{16};
// M tile dimension of GEMM.
int mTileM{128};
// N tile dimension of GEMM.
int mTileN{32};
// Tile scheduler type.
TileScheduler mTileScheduler{TileScheduler::Static};
// Save output of MMA in M-major format.
@ -520,24 +484,6 @@ inline std::string dumpOptions(GemmOptions const& options)
ss << "mNumStagesWorkId=" << options.mNumStagesWorkId << "," << std::endl;
ss << "mOutputDebugTensors=" << options.mOutputDebugTensors << "," << std::endl;
ss << "mPatchF2fp=" << options.mPatchF2fp << "," << std::endl;
ss << "mUseShuffledMatrixA=" << options.mUseShuffledMatrixA << "," << std::endl;
ss << "mSliceK=" << options.mSliceK << "," << std::endl;
ss << "mSplitK="
<< "gemm::SplitK(" << static_cast<int32_t>(options.mSplitK) << ")"
<< "," << std::endl;
ss << "mTransposeMmaOutput=" << options.mTransposeMmaOutput << "," << std::endl;
ss << "mTileM=" << options.mTileM << "," << std::endl;
ss << "mTileN=" << options.mTileN << "," << std::endl;
ss << "mTileK=" << options.mTileK << "," << std::endl;
ss << "mUseUnrollLoop2xForMma=" << options.mUseUnrollLoop2xForMma << "," << std::endl;
ss << "mUseCustomMmaSchedule=" << options.mUseCustomMmaSchedule << "," << std::endl;
ss << "mUseHoistTryWaitForCustomMmaSchedule=" << options.mUseHoistTryWaitForCustomMmaSchedule << "," << std::endl;
ss << "mUseDeepSeekFp8=" << options.mUseDeepSeekFp8 << "," << std::endl;
ss << "mUsePerTokenSfA=" << options.mUsePerTokenSfA << "," << std::endl;
ss << "mUsePerTokenSfB=" << options.mUsePerTokenSfB << "," << std::endl;
ss << "mUseTmaStore=" << options.mUseTmaStore << "," << std::endl;
ss << "mUseTwoTmaLoadWarps=" << options.mUseTwoTmaLoadWarps << "," << std::endl;
ss << "mUseTwoMmaWarps=" << options.mUseTwoMmaWarps << "," << std::endl;
if (options.mSfBlockSizeA.has_value())
{
ss << "mSfBlockSizeA=" << options.mSfBlockSizeA.value() << "," << std::endl;
@ -558,6 +504,13 @@ inline std::string dumpOptions(GemmOptions const& options)
<< "trtllm::gen::SfLayout(" << static_cast<int32_t>(options.mSfLayoutC) << ")"
<< "," << std::endl;
ss << "mSfReshapeFactor=" << options.mSfReshapeFactor << "," << std::endl;
ss << "mSliceK=" << options.mSliceK << "," << std::endl;
ss << "mSplitK="
<< "gemm::SplitK(" << static_cast<int32_t>(options.mSplitK) << ")"
<< "," << std::endl;
ss << "mTileK=" << options.mTileK << "," << std::endl;
ss << "mTileM=" << options.mTileM << "," << std::endl;
ss << "mTileN=" << options.mTileN << "," << std::endl;
ss << "mTileScheduler="
<< "gemm::TileScheduler(" << static_cast<int32_t>(options.mTileScheduler) << ")"
<< "," << std::endl;
@ -609,6 +562,7 @@ inline int32_t getShuffleBlockSize(int epilogueTileM)
// Check if the options are valid or not.
inline bool checkAndUpdateGemmOptions(GemmOptions& options, bool isBlackwell, int tpGrpSize, bool updateOptions = true)
{
options.mWorldSize = tpGrpSize;
if (options.mDtypeB == tg::Dtype::Void)
{
@ -1265,23 +1219,17 @@ inline bool checkAndUpdateGemmOptions(GemmOptions& options, bool isBlackwell, in
int const clampedAndPaddedPerCtaK = divUpMul(perCtaK - paddingForK, options.mTileK);
if (options.mUseUnrollLoop2xForMma)
{
// Number of iterations in K dimension after padding.
// Note the perCtaK in each CTA in the splitK group are padded to the same number of iterations.
// E.g., K = 512, TileK = 128, numSlicesForSplitK = 3. Then the padded K is
// Check that the padded K and clamped padded K (K rounded to next multiple of tileK) is a
// multiple of 2*TileK when UnrollLoop2x is enabled. This is to avoid deadlock when mma runs
// even-numbered loop while the other warps run odd-numbered loop.
//
// ceil(512 / (128*3)) * (128*3) = 768
//
int paddedK = divUpMul(options.mK, options.mTileK * options.mNumSlicesForSplitK);
// Check that the padded K (K rounded to next multiple of tileK) is a multiple of 2*TileK when
// UnrollLoop2x is enabled. This is to avoid deadlock when mma runs even-numbered loop while the
// other warps run odd-numbered loop.
//
bool notSupported = (paddedK / options.mNumSlicesForSplitK) % (options.mTileK * 2) != 0;
bool notSupported
= (perCtaK % (options.mTileK * 2) != 0) || (clampedAndPaddedPerCtaK % (options.mTileK * 2) != 0);
if (notSupported)
{
TLLM_LOG_WARNING("Size K / splitK must be a multiple of TileK * 2. Found TileK=", options.mTileK,
" and K=", options.mK, " (paddedK=", paddedK, ") and numSlicesForSplitK=", options.mNumSlicesForSplitK,
". Disabling unrollLoop2xForMma.");
" and K=", options.mK, " (paddedK=", paddedK, " clampedAndPaddedPerCtaK=", clampedAndPaddedPerCtaK,
") and numSlicesForSplitK=", options.mNumSlicesForSplitK, ". Disabling unrollLoop2xForMma.");
if (updateOptions)
{
options.mUseUnrollLoop2xForMma = false;

View File

@ -1,6 +1,8 @@
{
"templates": {
"GemmFp4LowLatency": {
"smVersion": "100f",
"patchF2fp": false,
"dtypeA": "e2m1",
"dtypeC": "e2m1",
"mmaM": 128,
@ -32,6 +34,8 @@
"usePdl": true
},
"GemmFp4Throughput": {
"smVersion": "100f",
"patchF2fp": false,
"dtypeA": "e2m1",
"dtypeC": "e2m1",
"mmaM": 128,
@ -64,6 +68,8 @@
"usePdl": true
},
"GemmFp8DeepSeekLowLatency": {
"smVersion": "100f",
"patchF2fp": false,
"dtypeA": "e4m3",
"dtypeC": "e4m3",
"mmaM": 64,
@ -94,6 +100,8 @@
},
"GemmPerTensorScalingFp8Throughput": {
"smVersion": "100f",
"patchF2fp": false,
"dtypeA": "e4m3",
"dtypeC": "e4m3",
"mmaM": 128,
@ -124,6 +132,8 @@
"usePdl": true
},
"GemmPerTensorScalingFp8LowLatency": {
"smVersion": "100f",
"patchF2fp": false,
"dtypeA": "e4m3",
"dtypeC": "e4m3",
"mmaM": 128,
@ -153,6 +163,8 @@
"usePdl": true
},
"GemmDeepSeekFp8LowLatency": {
"smVersion": "100f",
"patchF2fp": false,
"dtypeA": "e4m3",
"dtypeC": "e4m3",
"mmaM": 64,
@ -185,6 +197,8 @@
"usePdl": true
},
"GemmDeepSeekFp8Throughput": {
"smVersion": "100f",
"patchF2fp": false,
"dtypeA": "e4m3",
"dtypeC": "e4m3",
"mmaM": 64,
@ -218,6 +232,8 @@
"gridWaitForPrimaryB": false
},
"GemmMxE2m1MxE4m3LowLatency": {
"smVersion": "100f",
"patchF2fp": false,
"dtypeA": "mxe2m1",
"dtypeB": "mxe4m3",
"dtypeC": "mxe4m3",
@ -250,6 +266,8 @@
"usePdl": true
},
"GemmFp4xFp8": {
"smVersion": "100f",
"patchF2fp": false,
"dtypeA": "e2m1",
"dtypeMmaA": "e4m3",
"dtypeB": "e4m3",
@ -354,6 +372,7 @@
"_comment": "Tile 8 to 128",
"dtypeC": ["bf16", "fp16", "e4m3"],
"useUnrollLoop2xForMma": [true, false],
"smVersion": ["100a", "103a"],
"mmaN,tileN,epilogueTileN,tileK,numSlicesForSplitK,clusterDimZ": [[8, 8, 8, 512, 2, 2], [16, 16, 16, 512, 2, 2], [32, 32, 32, 512, 2, 2], [64, 64, 64, 512, 2, 2], [128, 128, 128, 256, 1, 1]]
}
]

View File

@ -0,0 +1,3 @@
version https://git-lfs.github.com/spec/v1
oid sha256:57baebc1bebdcc4690b3c8061538ad28f43a1d978fe60ea20e7e7f41b7bbdf68
size 386058

View File

@ -0,0 +1,3 @@
version https://git-lfs.github.com/spec/v1
oid sha256:cecaa0a37da10747bcd837087ad219d44b1b8e10b131405f22f2e63caac0627b
size 386058

View File

@ -0,0 +1,3 @@
version https://git-lfs.github.com/spec/v1
oid sha256:0a2358e0f5ac493f86111bf55d990d6215e7d18e7bfec9670aeacc3ccbe10cf4
size 401106

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:7f811d33d80842160a2ee26716d3e5bf4ad7fd2a497e8761d067a72f1608053b
size 402540

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:643f83e0b2d1c140d4ac7d8adc01bdd219a9d275ec12baff52e2b8696c117bd3
size 402540

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:3c31e771760a040cdccdecc2cb115b4ac88a216ed1dacd36a571beaafce00eef
size 435250

View File

@ -0,0 +1,3 @@
version https://git-lfs.github.com/spec/v1
oid sha256:7844f9bdb0f1979e72de458914b16564e7db0d8bc2c5767e9f4e72d52ec1add1
size 442306

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:b5f970c1d7af7eb0dae865d637e3ceabc741541569ef080a021d3aa0dc2fde54
size 441322

View File

@ -0,0 +1,3 @@
version https://git-lfs.github.com/spec/v1
oid sha256:a231df3e4e7cec4a64d6d0ccdf72cc7f64de871db9c15f9737611279d5a9d5b0
size 450252

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:f5468140541cba6b53f5c5233ddd8d451dba42c66fe7d66c6580ccff36b7fec1
size 539502

View File

@ -0,0 +1,3 @@
version https://git-lfs.github.com/spec/v1
oid sha256:2a45070b882db932cd763a16c3f331d42aa62f2d1db78ed79818bb8f688fc80c
size 547296

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:883067370db4a8dcf2a52c6e148a8c59e8d74d4491d0b94670e29869fb9a1e7e
size 610644

View File

@ -0,0 +1,3 @@
version https://git-lfs.github.com/spec/v1
oid sha256:82a242b53ae8e0751b780c9b7c30bf2dd62a443ee33840dea3862ff079eaefe6
size 619525

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:0985fb3d4c681ff1ad3c13d20f340c3a148eaab5f4df15d0bed35094f92ce113
size 561012

View File

@ -0,0 +1,3 @@
version https://git-lfs.github.com/spec/v1
oid sha256:ec9a6d672df79bfef38e0a3463a381b2c79991286e0e28f002483e17eeef88f9
size 568806

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:8a9e112e3175ee6d64c1ff76cc3b404dd75c3b2141a7146160ecbb9a57bbe4c6
size 632155

View File

@ -0,0 +1,3 @@
version https://git-lfs.github.com/spec/v1
oid sha256:e8b5f812e202a85f979969ac56a8958f73dd18bce5915daccfc76a4a88737853
size 641035

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:772428f18b342613c83eeef97d6b38f09cf399997f218f4440654d511f2cc8d5
size 615920

View File

@ -0,0 +1,3 @@
version https://git-lfs.github.com/spec/v1
oid sha256:e5709895e9ff9322a0bdec4287bf9971766759fc84c8e74e65cd07dfafe6c63a
size 622925

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:1df03bdccc4f41b298be4d2df28a764cda834c6ae9ed8ff91f1fd5b45b2e8d16
size 686965

View File

@ -0,0 +1,3 @@
version https://git-lfs.github.com/spec/v1
oid sha256:33168a4e0a1f46a1bc0bfda709b26d967d793bc8c686df93d46b4734f2948f73
size 695845

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:80af264b83492c79130ce08a95c9cbd1d58fa0d36da7759dd0a495209f83485d
size 528742

View File

@ -0,0 +1,3 @@
version https://git-lfs.github.com/spec/v1
oid sha256:60ee6676a92ba6b30f8d3c253883e65a51822fe26814bce699c1b81cb3eb0cf0
size 536438

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:f0fe35dc86d38bcef3c3c689cbb561199cf8cf98597f9ff4ae9895eb58e2a369
size 599884

View File

@ -0,0 +1,3 @@
version https://git-lfs.github.com/spec/v1
oid sha256:0cf701be1d5f702f14a837bcacf2d54eceae395d293835c67abf70b0c56875e2
size 608764

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:903142bcccbe54f2d7b494eb99d121880c22335051d45faf0b95104b874eb463
size 511260

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:5f0fadcfe6a193d113bcb19062514c82d2995f186d60695565e3f483f83f67b2
size 542788

View File

@ -0,0 +1,3 @@
version https://git-lfs.github.com/spec/v1
oid sha256:facb3f578a783c69d110096c735beb9e7ff56ac2a1336b054581358824125d5f
size 281718

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:c1c05018c4da492492d9c8bab66c424917c46117f8b7a67208043f8fc7fba704
size 291096

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:ef59df3fb79d5f5c21fa24b315f256a7182b61ad303143a52bba813198b9f9af
size 291096

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:2c6d934483b0c258a701b7d865e40fe7aa7c477336634007cb7dca7fae95871d
size 388906

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:3a1aef3173b48379b2bc6088c7b3171b335da30373340f181e6b3128c51691b1
size 423592

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:c4ea06d167a693a818dd590b56401af415f897c092ced523d0fbfa6b0c5906eb
size 405976

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:805c814d02aca70430eaec401317f931b3854d7fe4f07d655ca958ca42533205
size 441450

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:aebbf5a4b7e2c5551b9da7f161a7e4e595034ddb85c7edbca1cb4f163c2ce108
size 446034

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:d9145c573e096aefb3bbb56c1cadd31eba4b6cd3ea8db4c2f219be2ff3b5bf35
size 480720

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:5102b2d333b6351fd81c87fd3b6d565b0a11487c18ac757c5fcfb4bdf07ab6e2
size 414998

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:2331b555c625adbcc552f000423a1563943199c4cf83743d1e0dff865fbe77b1
size 449684

View File

@ -0,0 +1,3 @@
version https://git-lfs.github.com/spec/v1
oid sha256:43f5ccef2bf74c02ee600eb32087f9af0061ff78d96af3afcaf9709352c0f757
size 268504

View File

@ -0,0 +1,3 @@
version https://git-lfs.github.com/spec/v1
oid sha256:74516021158480dd91bcbdad6baca90772c424a6d04f7493b7ccfbe80481ca70
size 273404

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:810dd29b0147752e090272e476795327299eefe3594b07ad6a916b997b7715dc
size 291990

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:4364bc6ccd1bffa6dc32548f830dc1aa72e22563e3dfc93819540e1dd160a912
size 291990

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:9cfa36aa48820d645e8baf93101c186d6a1d47c959ef0f1858be993477046449
size 296100

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:63c171d81fdf24812ca43b5ded22a9f717602839e1b468c9504da4d19cdc4859
size 296100

View File

@ -0,0 +1,3 @@
version https://git-lfs.github.com/spec/v1
oid sha256:43822b8373319e1aacf183c17409ca74ed5d21fc0fa1258fda2289eb7d6c6c4a
size 451636

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:fad520c7710f04f82b728c641b17747f07dd2dfc430aeab12bd294cde2a295ac
size 482720

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:c4103e063e856bd93f518518b59370a79421c62035f3e1077e1b7e88b75b365c
size 482720

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:849b6eb121c1125dcd8944ebd5fa0bbdb5c25ae0cd083aff8942623d3b1ee6cc
size 435242

View File

@ -0,0 +1,3 @@
version https://git-lfs.github.com/spec/v1
oid sha256:3f6025b8a5eaafb2aaf79511b1e18da169f6aed24bdff59645dc27c8c3cefcdf
size 443086

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:e0def6d94d10c80e70f7c57deed553f8847bf7b32fb02d8d4de8d4e744897ee6
size 442104

View File

@ -0,0 +1,3 @@
version https://git-lfs.github.com/spec/v1
oid sha256:e59f5b13904ef4df7e0762d3ef5a575205e50e7c024c74853c4607e13ba7bfc9
size 450244

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:1b25ce03959310c1ceb7b738e23b42314ae77185d874dcb3136a54bdb3afb8b8
size 538112

View File

@ -0,0 +1,3 @@
version https://git-lfs.github.com/spec/v1
oid sha256:343cc14d909b7b88c92081b8367afdde5243f1169eefba6f8df184ee70a5453b
size 545908

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:98270c39f23adcdc847e40ec4af0906bce0eadfadc0478cd80b6c5b01b67d4d2
size 610044

View File

@ -0,0 +1,3 @@
version https://git-lfs.github.com/spec/v1
oid sha256:4026662f036e5d34cc204e69a488f2309d5aea18a14ed2694b3cb7b544df50df
size 618925

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:4cfe52e1a652b3c98a7036e8183dd54c19eb656b93c6cf1c4bb8515aaa0cb618
size 559128

View File

@ -0,0 +1,3 @@
version https://git-lfs.github.com/spec/v1
oid sha256:a2bd28541ed6e2404736cc41b320a4c7f9affadb03d33678984edb5c5f7dd90a
size 567712

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:5b3bdf48cff6be67299350858d2e179db6a168b8b4499df4024bf8abec748534
size 631851

View File

@ -0,0 +1,3 @@
version https://git-lfs.github.com/spec/v1
oid sha256:4962e405081cb73e85dba85ee67afb17642ac446f37e0f8cb078099a1ff894de
size 640731

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:8237552efb5f053c7b93118145411ea1c440d7303da2d9596f113ae9bca687a3
size 614332

View File

@ -0,0 +1,3 @@
version https://git-lfs.github.com/spec/v1
oid sha256:d56ff0911d2163a88a6e289a7655ef43c33e7794c2e019acc094f5f2590da68f
size 622917

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:857632aea04ed78ead18b560bde571e98d1b91d2e255ff596333273f8106a1a6
size 686907

View File

@ -0,0 +1,3 @@
version https://git-lfs.github.com/spec/v1
oid sha256:b414b9f078fd18141a6dd945292003cdcc7b2564f1464f09f709621fc75a4058
size 695837

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:a335e06b9157f5a1d3662a31fca4db8eadcd523dddf5a4acd4aabd35ad83d3fc
size 526464

View File

@ -0,0 +1,3 @@
version https://git-lfs.github.com/spec/v1
oid sha256:0d40fb7859add20033f9f25a74ec35844624101737bbc09e8b3ed1180b4868dd
size 534950

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:864aaafa89f55cc352096dd0fab344b2c26e4ce1bc838a5ddf2310126b9294e2
size 598346

View File

@ -0,0 +1,3 @@
version https://git-lfs.github.com/spec/v1
oid sha256:517192044528f06c12fd0137b6330b0c4a8c820556094fa2f3c81519de679c4d
size 606486

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:f8d770bfe3bd1bbfc7ba7e1d998356e7c5fa30128e531475630beddbf6a89518
size 674989

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:7592d2fcdd5334457804c60e997caab7b55bca0b98e49b704338851b1f26fcc5
size 707257

View File

@ -0,0 +1,3 @@
version https://git-lfs.github.com/spec/v1
oid sha256:3e4d9abb2428146f9caae7cbb5acbf1cfc13d25d32861f5d917a062e9ee45dc2
size 289554

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:f20d9bf427d6963d127d8fd7ffd455164e97dafe74bc129187cb13aa2c1afb7b
size 304112

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:cfb3f952d55edb894d93b8b022884865371fb3bf63033d74ce34a05e1f7e8820
size 304112

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:f4e42191612ceffb7de8b7d112433f54a2d136f215c39782aa46384fdef72a04
size 411986

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:4aea2a209820f76e2408c7204a09b51d49b98691b1bc42ab2ed41e6ba279741d
size 446672

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:34b2788fd26d7f32f515f331d68bc1b1d9f52ab7e58795f74f568e3ea7114cf3
size 449134

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:5e4858f8eb635c01dedb7a10c17a948aa7d792438052e278f12083c504ef5344
size 483820

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:239d1614204c6114a67842a21de11fa88f2da2c838c819609428f409c1a692f0
size 525206

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:7f920f02c0aa0dcdf604ce1b6a9f6f3e076a0974ec301cbb49f325ddd4639e54
size 561470

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:151143fe03078b16180e32fb803c0cbb67e3e4cda70d7e080971cc590ca14cf7
size 425992

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:35d609db47a06ef1df2365983a4f6f4aeee5326731d42d21cd3f763c35ed8dfe
size 460676

View File

@ -0,0 +1,3 @@
version https://git-lfs.github.com/spec/v1
oid sha256:300caabdc160e52ecc1e6d06c9cbb19d99da2a754517c7a90e777212f29a7973
size 266916

View File

@ -0,0 +1,3 @@
version https://git-lfs.github.com/spec/v1
oid sha256:872b83bf44b95a4fce348d71ed632a1f9923fbc3ffad912a8f61bcc6727fbe70
size 271768

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:f4ad4bf1c626ebc7772502f08694368ff1f40f05902595dd245fa2b7afa9c087
size 289614

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:09ecee09c11f19eca97c5406b81a2f8abd87fc72d5194798accbd58a5f7a8411
size 289614

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:d5a182c886f003dbfe1ad202c81e5be1463778b39479bd173727a1cb0a2b951f
size 294514

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:df3685b97caa5f7df6fcba3b3da4b74ee50cf868357ae41b815d3f37650e2550
size 294514

View File

@ -0,0 +1,3 @@
version https://git-lfs.github.com/spec/v1
oid sha256:34eb8e5d484e36740b10a8b5a536733cd5dd6ddc3a0f965cac796566c1614ab5
size 391626

View File

@ -0,0 +1,3 @@
version https://git-lfs.github.com/spec/v1
oid sha256:4696fe15f42eb29905c9b143bb17a60b89b69de07e796551b047531dfa37cabe
size 391626

View File

@ -0,0 +1,3 @@
version https://git-lfs.github.com/spec/v1
oid sha256:a028115d2dfad91a372fa4db7078c466126d6df5286465c23a173cda4100c214
size 406672

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:03723b80c6a02fe986e0a7c64447fbd7233e63d49f5638ededf5ecaf079e9f0d
size 408106

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:4f71e297a4440d974e6815927b8a4c927385160672363f4cc59cc8a70d2b755f
size 408106

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:910dc558f3954db07df445f2c8655eb8762e5cef44eb402ac15ac3da61422a1b
size 433664

View File

@ -0,0 +1,3 @@
version https://git-lfs.github.com/spec/v1
oid sha256:4583d906e291e5408461bd7e6904fcb7ccd02cdd3812c7f8fe3766044b2d8cde
size 441508

View File

@ -1,3 +0,0 @@
version https://git-lfs.github.com/spec/v1
oid sha256:e8d96b98b3533bd19e0124278967d196210549a599628fc5f52cb6e0d66398fa
size 439736

Some files were not shown because too many files have changed in this diff Show More