/************************************************************************* * Copyright (c) 2022-2024, NVIDIA CORPORATION & AFFILIATES. All rights reserved. * * See LICENSE for license information. ************************************************************************/ #include #include #include #include #include #include #include #include #include #include #include "../test_common.h" using namespace transformer_engine; namespace { template void compute_ref_q(const InputType *data, OutputType *output, const size_t N, float *amax, float scale) { using compute_t = float; compute_t current_max = -1e100; for (size_t i = 0; i < N; ++i) { compute_t current = static_cast(data[i]); current_max = fmaxf(current_max, fabsf(current)); if (std::is_same::value || std::is_same::value) { output[i] = OutputType(scale * current); } else { output[i] = OutputType(current); } } *amax = current_max; } template void compute_ref_dq(const InputType *data, OutputType *output, const size_t N, float scale_inv) { using compute_t = float; for (size_t i = 0; i < N; ++i) { compute_t current = static_cast(data[i]); output[i] = OutputType(scale_inv * current); } } template void performTestQ(const size_t N) { using namespace test; DType itype = TypeInfo::dtype; DType otype = TypeInfo::dtype; Tensor input({ N }, itype); Tensor output({ N }, otype); std::unique_ptr ref_output = std::make_unique(N); fillUniform(&input); setRandomScale(&output); nvte_fp8_quantize(input.data(), output.data(), 0); float ref_amax; compute_ref_q(input.cpu_dptr(), ref_output.get(), N, &ref_amax, output.scale()); cudaDeviceSynchronize(); auto err = cudaGetLastError(); ASSERT_EQ(err, cudaSuccess) << cudaGetErrorString(err); auto [atol_amax, rtol_amax] = getTolerances(DType::kFloat32); compareResults("amax", output.amax(), ref_amax, atol_amax, rtol_amax); auto [atol, rtol] = getTolerances(otype); compareResults("output_q", output, ref_output.get(), atol, rtol); } template void performTestDQ(const size_t N) { using namespace test; DType itype = TypeInfo::dtype; DType otype = TypeInfo::dtype; Tensor input({ N }, itype); Tensor output({ N }, otype); std::unique_ptr ref_output = std::make_unique(N); fillUniform(&input); nvte_fp8_dequantize(input.data(), output.data(), 0); compute_ref_dq(input.cpu_dptr(), ref_output.get(), N, input.scale_inv()); cudaDeviceSynchronize(); auto err = cudaGetLastError(); ASSERT_EQ(err, cudaSuccess) << cudaGetErrorString(err); auto [atol, rtol] = getTolerances(otype); compareResults("output_dq", output, ref_output.get(), atol, rtol); } std::vector qdq_test_cases = {2048* 12288, 768 * 1024, 256 * 65536, 65536 * 128, 257 * 259, 128*128+1}; } //namespace class QDQTestSuite : public ::testing::TestWithParam> {}; TEST_P(QDQTestSuite, TestQ) { using namespace transformer_engine; using namespace test; const DType input_type = std::get<0>(GetParam()); const DType output_type = std::get<1>(GetParam()); const size_t N = std::get<2>(GetParam()); TRANSFORMER_ENGINE_TYPE_SWITCH_ALL(input_type, InputType, TRANSFORMER_ENGINE_TYPE_SWITCH_ALL(output_type, OutputType, performTestQ(N); ); ); } TEST_P(QDQTestSuite, TestDQ) { using namespace transformer_engine; using namespace test; const DType input_type = std::get<0>(GetParam()); const DType output_type = std::get<1>(GetParam()); const size_t N = std::get<2>(GetParam()); TRANSFORMER_ENGINE_TYPE_SWITCH_ALL(input_type, InputType, TRANSFORMER_ENGINE_TYPE_SWITCH_ALL(output_type, OutputType, performTestDQ(N); ); ); } INSTANTIATE_TEST_SUITE_P( OperatorTest, QDQTestSuite, ::testing::Combine( ::testing::Values(DType::kFloat32, DType::kBFloat16, DType::kFloat16), ::testing::Values(DType::kFloat8E4M3, DType::kFloat8E5M2), ::testing::ValuesIn(qdq_test_cases)), [](const testing::TestParamInfo& info) { std::string name = test::typeName(std::get<0>(info.param)) + "X" + test::typeName(std::get<1>(info.param)) + "X" + std::to_string(std::get<2>(info.param)); return name; });