Skip to content

Commit 7c2aa2c

Browse files
Oleg-Goncharovpre-commit-ci[bot]
Oleg-Goncharov
and
pre-commit-ci[bot]
authored
[Common] MOE Split dBias (NVIDIA#2674)
* Implemented the kernel with split dbias Signed-off-by: Oleg Goncharov <ogoncharov@nvidia.com> * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci Signed-off-by: Oleg Goncharov <ogoncharov@nvidia.com> * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci Signed-off-by: Oleg Goncharov <ogoncharov@nvidia.com> * Relaxed constraints on the last dimension Signed-off-by: Oleg Goncharov <ogoncharov@nvidia.com> * Added notes on group tensor restrictions into documentation Signed-off-by: Oleg Goncharov <ogoncharov@nvidia.com> * Fixes per the review Signed-off-by: Oleg Goncharov <ogoncharov@nvidia.com> * Fixed pointer Signed-off-by: Oleg Goncharov <ogoncharov@nvidia.com> * More fixes Signed-off-by: Oleg Goncharov <ogoncharov@nvidia.com> * Fixed kernel grid size Signed-off-by: Oleg Goncharov <ogoncharov@nvidia.com> --------- Signed-off-by: Oleg Goncharov <ogoncharov@nvidia.com> Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>
1 parent e6d97ff commit 7c2aa2c

10 files changed

Lines changed: 236 additions & 124 deletions

File tree

‎tests/cpp/operator/test_cast_mxfp8_grouped.cu‎

Lines changed: 78 additions & 72 deletions
Original file line numberDiff line numberDiff line change
@@ -58,8 +58,7 @@ void compute_ref(const ProcessingMethod processing_method,
5858
const size_t rows,
5959
const size_t cols,
6060
const size_t scales_stride_rowwise,
61-
const size_t scales_stride_colwise,
62-
const bool is_single_tensor)
61+
const size_t scales_stride_colwise)
6362
{
6463
const size_t tile_size_Y = 32;
6564
const size_t tile_size_X = 32;
@@ -169,10 +168,8 @@ void compute_ref(const ProcessingMethod processing_method,
169168
}
170169
}
171170

172-
if (is_single_tensor) {
173-
for (size_t j = 0; j < cols; ++j) {
174-
output_dbias[j] = static_cast<InputType>(output_dbias_fp32[j]);
175-
}
171+
for (size_t j = 0; j < cols; ++j) {
172+
output_dbias[j] = static_cast<InputType>(output_dbias_fp32[j]);
176173
}
177174
}
178175

@@ -250,19 +247,24 @@ void performTest(const ProcessingMethod processing_method,
250247
DType itype = TypeInfo<InputType>::dtype;
251248
DType otype = TypeInfo<OutputType>::dtype;
252249

250+
const bool compute_dbias = (processing_method == ProcessingMethod::CAST_DBIAS
251+
|| processing_method == ProcessingMethod::CAST_DBIAS_DACT);
252+
253253
const size_t rows = logical_shape_vec[0];
254254
const size_t cols = logical_shape_vec[1];
255255

256256
size_t elts_num = 0;
257257
size_t rowwise_sfs_num = 0;
258258
size_t colwise_sfs_num = 0;
259+
size_t sum_of_last_dims = 0;
259260

260261
std::vector<size_t> rowwise_scales_first_dim(num_tensors, 0);
261262
std::vector<size_t> rowwise_scales_last_dim(num_tensors, 0);
262263
std::vector<size_t> rowwise_scales_offset(num_tensors + 1, 0);
263264
std::vector<size_t> colwise_scales_first_dim(num_tensors, 0);
264265
std::vector<size_t> colwise_scales_last_dim(num_tensors, 0);
265266
std::vector<size_t> colwise_scales_offset(num_tensors + 1, 0);
267+
std::vector<size_t> dbias_offsets(num_tensors + 1, 0);
266268

267269
for (size_t t = 0; t < num_tensors; ++t) {
268270
const size_t M = first_dims_h[t];
@@ -285,13 +287,13 @@ void performTest(const ProcessingMethod processing_method,
285287

286288
rowwise_sfs_num += rowwise_sfs;
287289
colwise_sfs_num += colwise_sfs;
290+
sum_of_last_dims += K;
288291

289292
rowwise_scales_offset[t+1] = rowwise_sfs_num;
290293
colwise_scales_offset[t+1] = colwise_sfs_num;
294+
dbias_offsets[t+1] = sum_of_last_dims;
291295
}
292296

293-
const bool is_single_tensor = (shape_rep == SAME_BOTH_DIMS) || (shape_rep == VARYING_FIRST_DIM);
294-
295297
std::vector<size_t> scales_rowwise_shape = {rowwise_sfs_num};
296298
std::vector<size_t> scales_colwise_shape = {colwise_sfs_num};
297299

@@ -311,7 +313,7 @@ void performTest(const ProcessingMethod processing_method,
311313
std::vector<fp8e8m0> out_scales_rowwise_ref(rowwise ? rowwise_sfs_num : 0);
312314
std::vector<fp8e8m0> out_scales_colwise_ref(colwise ? colwise_sfs_num : 0);
313315

314-
std::vector<InputType> ref_output_dbias(is_single_tensor ? cols : 0);
316+
std::vector<InputType> ref_output_dbias(sum_of_last_dims, static_cast<InputType>(0.0f));
315317

316318
for (size_t i = 0; i < elts_num; ++i) {
317319
const float val = dis(gen);
@@ -336,22 +338,24 @@ void performTest(const ProcessingMethod processing_method,
336338

337339
const size_t in_data_size = elts_num * sizeof(InputType);
338340
const size_t out_data_size = elts_num * sizeof(OutputType);
341+
const size_t dbias_data_size = sum_of_last_dims * sizeof(InputType);
339342
const size_t rowwise_scales_size = rowwise_sfs_num * sizeof(fp8e8m0);
340343
const size_t colwise_scales_size = colwise_sfs_num * sizeof(fp8e8m0);
341344

342345
const size_t first_dims_size = num_tensors * sizeof(size_t);
343346
const size_t last_dims_size = num_tensors * sizeof(size_t);
344347
const size_t offsets_size = (num_tensors + 1) * sizeof(size_t);
345348

346-
InputType* grad_data_d;
347-
InputType* in_data_d;
348-
OutputType* out_data_rowwise_d;
349-
OutputType* out_data_colwise_d;
350-
fp8e8m0* out_scales_rowwise_d;
351-
fp8e8m0* out_scales_colwise_d;
352-
size_t* first_dims_d;
353-
size_t* last_dims_d;
354-
size_t* offsets_d;
349+
InputType* grad_data_d = nullptr;
350+
InputType* in_data_d = nullptr;
351+
InputType* dbias_out_data_d = nullptr;
352+
OutputType* out_data_rowwise_d = nullptr;
353+
OutputType* out_data_colwise_d = nullptr;
354+
fp8e8m0* out_scales_rowwise_d = nullptr;
355+
fp8e8m0* out_scales_colwise_d = nullptr;
356+
size_t* first_dims_d = nullptr;
357+
size_t* last_dims_d = nullptr;
358+
size_t* offsets_d = nullptr;
355359

356360
cudaMalloc((void**)&grad_data_d, in_data_size);
357361
cudaMalloc((void**)&in_data_d, in_data_size);
@@ -367,6 +371,10 @@ void performTest(const ProcessingMethod processing_method,
367371

368372
NVTEShape logical_shape_ = nvte_make_shape(logical_shape_vec.data(), logical_shape_vec.size());
369373

374+
std::vector<size_t> dbias_logical_shape_vec= {num_tensors, cols};
375+
NVTEShape dbias_logical_shape_ = nvte_make_shape(dbias_logical_shape_vec.data(),
376+
dbias_logical_shape_vec.size());
377+
370378
NVTEShape first_dims_shape_;
371379
NVTEShape last_dims_shape_;
372380
NVTEShape offsets_shape_;
@@ -382,6 +390,7 @@ void performTest(const ProcessingMethod processing_method,
382390
NVTEGroupedTensor grad_group_tensor = nvte_create_grouped_tensor(NVTE_DELAYED_TENSOR_SCALING, num_tensors, logical_shape_);
383391
NVTEGroupedTensor in_group_tensor = nvte_create_grouped_tensor(NVTE_DELAYED_TENSOR_SCALING, num_tensors, logical_shape_);
384392
NVTEGroupedTensor out_group_tensor = nvte_create_grouped_tensor(NVTE_MXFP8_1D_SCALING, num_tensors, logical_shape_);
393+
NVTEGroupedTensor output_dbias_tensor = nvte_create_grouped_tensor(NVTE_DELAYED_TENSOR_SCALING, num_tensors, dbias_logical_shape_);
385394

386395
NVTEBasicTensor grad_data_tensor = {grad_data_d, static_cast<NVTEDType>(itype), logical_shape_};
387396
NVTEBasicTensor in_data_tensor = {in_data_d, static_cast<NVTEDType>(itype), logical_shape_};
@@ -453,52 +462,41 @@ void performTest(const ProcessingMethod processing_method,
453462
&out_scales_colwise_tensor, sizeof(out_scales_colwise_tensor));
454463
}
455464

456-
Tensor output_dbias("output_dbias", std::vector<size_t>{ cols }, itype);
465+
if (compute_dbias) {
466+
cudaMalloc((void**)&dbias_out_data_d, dbias_data_size);
467+
cudaMemset(dbias_out_data_d, 0, dbias_data_size);
468+
NVTEBasicTensor output_dbias_data_tensor = {dbias_out_data_d, static_cast<NVTEDType>(itype), dbias_logical_shape_};
469+
nvte_set_grouped_tensor_param(output_dbias_tensor, NVTEGroupedTensorParam::kNVTEGroupedRowwiseData,
470+
&output_dbias_data_tensor, sizeof(output_dbias_data_tensor));
471+
}
457472

458473
// Reference (CPU)
459-
if (is_single_tensor) {
460-
461-
const size_t unpadded_rowwise_blocks_X = divide_round_up(cols, 32);
462-
const size_t unpadded_colwise_blocks_X = cols;
474+
for (size_t t = 0; t < num_tensors; ++t) {
475+
const size_t M = first_dims_h[t];
476+
const size_t K = last_dims_h[t];
463477

464-
const size_t scales_stride_rowwise = round_up_to_nearest_multiple(unpadded_rowwise_blocks_X, 4);
465-
const size_t scales_stride_colwise = round_up_to_nearest_multiple(unpadded_colwise_blocks_X, 128);
478+
const size_t scales_stride_rowwise = rowwise_scales_last_dim[t];
479+
const size_t scales_stride_colwise = colwise_scales_last_dim[t];
480+
const size_t data_offset = offsets_h[t];
481+
const size_t rowwise_sfs_offset = rowwise_scales_offset[t];
482+
const size_t colwise_sfs_offset = colwise_scales_offset[t];
483+
const size_t dbias_offset = dbias_offsets[t];
484+
485+
const InputType* const grad_ptr = grad_data.data() + data_offset;
486+
const InputType* const in_ptr = in_data.data() + data_offset;
487+
OutputType* const out_data_rowwise_ptr = out_data_rowwise_ref.data() + data_offset;
488+
OutputType* const out_data_colwise_ptr = out_data_colwise_ref.data() + data_offset;
489+
fp8e8m0* const out_scales_rowwise_ptr = out_scales_rowwise_ref.data() + rowwise_sfs_offset;
490+
fp8e8m0* const out_scales_colwise_ptr = out_scales_colwise_ref.data() + colwise_sfs_offset;
491+
InputType* const ref_output_dbias_ptr = ref_output_dbias.data() + dbias_offset;
466492

467493
compute_ref<InputType, OutputType>(
468-
processing_method, OP, rowwise, colwise, in_data.data(), grad_data.data(),
469-
out_data_rowwise_ref.data(), out_data_colwise_ref.data(),
470-
out_scales_rowwise_ref.data(), out_scales_colwise_ref.data(),
471-
ref_output_dbias.data(), rows, cols,
494+
processing_method, OP, rowwise, colwise, in_ptr, grad_ptr,
495+
out_data_rowwise_ptr, out_data_colwise_ptr,
496+
out_scales_rowwise_ptr, out_scales_colwise_ptr,
497+
ref_output_dbias_ptr, M, K,
472498
scales_stride_rowwise,
473-
scales_stride_colwise,
474-
is_single_tensor);
475-
} else {
476-
for (size_t t = 0; t < num_tensors; ++t) {
477-
const size_t M = first_dims_h[t];
478-
const size_t K = last_dims_h[t];
479-
480-
const size_t scales_stride_rowwise = rowwise_scales_last_dim[t];
481-
const size_t scales_stride_colwise = colwise_scales_last_dim[t];
482-
const size_t data_offset = offsets_h[t];
483-
const size_t rowwise_sfs_offset = rowwise_scales_offset[t];
484-
const size_t colwise_sfs_offset = colwise_scales_offset[t];
485-
486-
const InputType* const grad_ptr = grad_data.data() + data_offset;
487-
const InputType* const in_ptr = in_data.data() + data_offset;
488-
OutputType* const out_data_rowwise_ptr = out_data_rowwise_ref.data() + data_offset;
489-
OutputType* const out_data_colwise_ptr = out_data_colwise_ref.data() + data_offset;
490-
fp8e8m0* const out_scales_rowwise_ptr = out_scales_rowwise_ref.data() + rowwise_sfs_offset;
491-
fp8e8m0* const out_scales_colwise_ptr = out_scales_colwise_ref.data() + colwise_sfs_offset;
492-
493-
compute_ref<InputType, OutputType>(
494-
processing_method, OP, rowwise, colwise, in_ptr, grad_ptr,
495-
out_data_rowwise_ptr, out_data_colwise_ptr,
496-
out_scales_rowwise_ptr, out_scales_colwise_ptr,
497-
ref_output_dbias.data(), M, K,
498-
scales_stride_rowwise,
499-
scales_stride_colwise,
500-
is_single_tensor);
501-
}
499+
scales_stride_colwise);
502500
}
503501

504502
// GPU
@@ -509,9 +507,9 @@ void performTest(const ProcessingMethod processing_method,
509507
break;
510508
}
511509
case ProcessingMethod::CAST_DBIAS: {
512-
nvte_group_quantize_dbias(grad_group_tensor, out_group_tensor, output_dbias.data(), workspace.data(), 0);
510+
nvte_group_quantize_dbias(grad_group_tensor, out_group_tensor, output_dbias_tensor, workspace.data(), 0);
513511
workspace = Tensor("workspace", workspace.rowwise_shape(), workspace.dtype());
514-
nvte_group_quantize_dbias(grad_group_tensor, out_group_tensor, output_dbias.data(), workspace.data(), 0);
512+
nvte_group_quantize_dbias(grad_group_tensor, out_group_tensor, output_dbias_tensor, workspace.data(), 0);
515513
break;
516514
}
517515
case ProcessingMethod::CAST_DBIAS_DACT: {
@@ -522,10 +520,10 @@ void performTest(const ProcessingMethod processing_method,
522520
else if (OP == &dsrelu) { nvte_group_quantize_dbias_dact = &nvte_group_quantize_dbias_dsrelu; }
523521

524522
nvte_group_quantize_dbias_dact(grad_group_tensor, in_group_tensor, out_group_tensor,
525-
output_dbias.data(), workspace.data(), 0);
523+
output_dbias_tensor, workspace.data(), 0);
526524
workspace = Tensor("workspace", workspace.rowwise_shape(), workspace.dtype());
527525
nvte_group_quantize_dbias_dact(grad_group_tensor, in_group_tensor, out_group_tensor,
528-
output_dbias.data(), workspace.data(), 0);
526+
output_dbias_tensor, workspace.data(), 0);
529527
break;
530528
}
531529
case ProcessingMethod::CAST_ACT: {
@@ -586,9 +584,10 @@ void performTest(const ProcessingMethod processing_method,
586584
out_data_colwise_h.data(), rows, cols, false, mismatches_elts);
587585
}
588586

589-
if (processing_method == ProcessingMethod::CAST_DBIAS
590-
|| processing_method == ProcessingMethod::CAST_DBIAS_DACT)
591-
{
587+
if (compute_dbias) {
588+
Tensor output_dbias("output_dbias", std::vector<size_t>{ sum_of_last_dims }, itype);
589+
cudaMemcpy(output_dbias.rowwise_dptr(), dbias_out_data_d, dbias_data_size, cudaMemcpyDeviceToDevice);
590+
592591
auto [atol_dbias, rtol_dbias] = getTolerances(itype);
593592
if (itype == DType::kFloat32) {
594593
atol_dbias = 1e-4;
@@ -601,6 +600,7 @@ void performTest(const ProcessingMethod processing_method,
601600

602601
cudaFree(grad_data_d);
603602
cudaFree(in_data_d);
603+
cudaFree(dbias_out_data_d);
604604
cudaFree(first_dims_d);
605605
cudaFree(last_dims_d);
606606
cudaFree(offsets_d);
@@ -648,7 +648,8 @@ std::vector<std::vector<size_t>> input_config = {
648648
{SAME_BOTH_DIMS, 1, 128,128},
649649
{SAME_BOTH_DIMS, 2, 256,128},
650650
{VARYING_FIRST_DIM, 2, 512,128, 128,384},
651-
{VARYING_FIRST_DIM, 2, 384,160, 128,256},
651+
{VARYING_FIRST_DIM, 3, 1024,144, 128,384,512},
652+
{VARYING_FIRST_DIM, 4, 1536,160, 128,384,512,512},
652653
{VARYING_FIRST_DIM, 5, 4096,512, 128,256,384,1024,2304},
653654
{VARYING_LAST_DIM, 3, 256,896, 128,256,512},
654655
{VARYING_BOTH_DIMS, 2, 1,(128*128)+(256*256), 128,256, 128,256},
@@ -714,26 +715,31 @@ TEST_P(GroupedFusedCastMXFP8TestSuite, Test) {
714715
}
715716
}
716717
offsets[t+1] = offsets[t] + first_dims[t] * last_dims[t];
717-
// Skips tests if tensor shape is not as required by the kernel
718-
if (first_dims[t] % 128 != 0) {
718+
// Skip tests when the tensor shape is incompatible with the kernel.
719+
// The TMA engine requires strides to be 16-byte aligned.
720+
if ((first_dims[t] % 128 != 0) || (last_dims[t] % 16 != 0)) {
719721
GTEST_SKIP();
720722
}
721-
if (!is_single_tensor && (last_dims[t] % 128 != 0)) {
723+
// If a grouped tensor has a varying last dimension, it must be a multiple of 128.
724+
// Otherwise, computing the grid size adds runtime overhead in the non-persistent kernel,
725+
// since the relevant tensor metadata resides in device memory.
726+
constexpr size_t CHUNK_DIM_X = 128;
727+
if (!is_single_tensor && (last_dims[t] % CHUNK_DIM_X != 0)) {
722728
GTEST_SKIP();
723729
}
724730
}
725-
// Skips DBias tests if last dimension of tensors variates
731+
// Skip dBias tests when tensors in the group have different last dimensions.
726732
if ((processing_method == ProcessingMethod::CAST_DBIAS || processing_method == ProcessingMethod::CAST_DBIAS_DACT)
727733
&& !is_single_tensor) {
728734
GTEST_SKIP();
729735
}
730736

731-
// Skips non Act tests if the Activation type is not an identity
737+
// Skip non-activation tests when the activation type is not Identity.
732738
if ((processing_method == ProcessingMethod::CAST_ONLY || processing_method == ProcessingMethod::CAST_DBIAS)
733739
&& activation != ActivationKind::Identity) {
734740
GTEST_SKIP();
735741
}
736-
// Skips Act tests if the Activation is an identity
742+
// Skip activation tests when the activation type is Identity.
737743
if ((processing_method == ProcessingMethod::CAST_DBIAS_DACT
738744
|| processing_method == ProcessingMethod::CAST_DACT
739745
|| processing_method == ProcessingMethod::CAST_ACT) && (activation == ActivationKind::Identity)) {

‎transformer_engine/common/activation/gelu.cu‎

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -32,7 +32,7 @@ void nvte_group_dgelu(const NVTEGroupedTensor grad, const NVTEGroupedTensor inpu
3232
NVTEGroupedTensor output, cudaStream_t stream) {
3333
NVTE_API_CALL(nvte_group_dgelu);
3434
using namespace transformer_engine;
35-
NVTETensor dbias = nullptr;
35+
NVTEGroupedTensor dbias = nullptr;
3636
NVTETensor workspace = nullptr;
3737

3838
constexpr bool IS_DBIAS = false;
@@ -57,7 +57,7 @@ void nvte_quantize_dbias_dgelu(const NVTETensor input, const NVTETensor activati
5757

5858
void nvte_group_quantize_dbias_dgelu(const NVTEGroupedTensor input,
5959
const NVTEGroupedTensor activation_input,
60-
NVTEGroupedTensor output, NVTETensor dbias,
60+
NVTEGroupedTensor output, NVTEGroupedTensor dbias,
6161
NVTETensor workspace, cudaStream_t stream) {
6262
NVTE_API_CALL(nvte_group_quantize_dbias_dgelu);
6363
using namespace transformer_engine;
@@ -110,7 +110,7 @@ void nvte_group_dqgelu(const NVTEGroupedTensor grad, const NVTEGroupedTensor inp
110110
NVTEGroupedTensor output, cudaStream_t stream) {
111111
NVTE_API_CALL(nvte_group_dqgelu);
112112
using namespace transformer_engine;
113-
NVTETensor dbias = nullptr;
113+
NVTEGroupedTensor dbias = nullptr;
114114
NVTETensor workspace = nullptr;
115115

116116
constexpr bool IS_DBIAS = false;
@@ -135,7 +135,7 @@ void nvte_quantize_dbias_dqgelu(const NVTETensor input, const NVTETensor activat
135135

136136
void nvte_group_quantize_dbias_dqgelu(const NVTEGroupedTensor input,
137137
const NVTEGroupedTensor activation_input,
138-
NVTEGroupedTensor output, NVTETensor dbias,
138+
NVTEGroupedTensor output, NVTEGroupedTensor dbias,
139139
NVTETensor workspace, cudaStream_t stream) {
140140
NVTE_API_CALL(nvte_group_quantize_dbias_dqgelu);
141141
using namespace transformer_engine;

‎transformer_engine/common/activation/relu.cu‎

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -32,7 +32,7 @@ void nvte_group_drelu(const NVTEGroupedTensor grad, const NVTEGroupedTensor inpu
3232
NVTEGroupedTensor output, cudaStream_t stream) {
3333
NVTE_API_CALL(nvte_group_drelu);
3434
using namespace transformer_engine;
35-
NVTETensor dbias = nullptr;
35+
NVTEGroupedTensor dbias = nullptr;
3636
NVTETensor workspace = nullptr;
3737

3838
constexpr bool IS_DBIAS = false;
@@ -57,7 +57,7 @@ void nvte_quantize_dbias_drelu(const NVTETensor input, const NVTETensor activati
5757

5858
void nvte_group_quantize_dbias_drelu(const NVTEGroupedTensor input,
5959
const NVTEGroupedTensor activation_input,
60-
NVTEGroupedTensor output, NVTETensor dbias,
60+
NVTEGroupedTensor output, NVTEGroupedTensor dbias,
6161
NVTETensor workspace, cudaStream_t stream) {
6262
NVTE_API_CALL(nvte_group_quantize_dbias_drelu);
6363
using namespace transformer_engine;
@@ -110,7 +110,7 @@ void nvte_group_dsrelu(const NVTEGroupedTensor grad, const NVTEGroupedTensor inp
110110
NVTEGroupedTensor output, cudaStream_t stream) {
111111
NVTE_API_CALL(nvte_group_dsrelu);
112112
using namespace transformer_engine;
113-
NVTETensor dbias = nullptr;
113+
NVTEGroupedTensor dbias = nullptr;
114114
NVTETensor workspace = nullptr;
115115

116116
constexpr bool IS_DBIAS = false;
@@ -135,7 +135,7 @@ void nvte_quantize_dbias_dsrelu(const NVTETensor input, const NVTETensor activat
135135

136136
void nvte_group_quantize_dbias_dsrelu(const NVTEGroupedTensor input,
137137
const NVTEGroupedTensor activation_input,
138-
NVTEGroupedTensor output, NVTETensor dbias,
138+
NVTEGroupedTensor output, NVTEGroupedTensor dbias,
139139
NVTETensor workspace, cudaStream_t stream) {
140140
NVTE_API_CALL(nvte_group_quantize_dbias_dsrelu);
141141
using namespace transformer_engine;

0 commit comments

Comments
 (0)