Test cases for Fp4 and fp6 coversion apis (#169)
Αυτή η υποβολή περιλαμβάνεται σε:
Venkatesh, Anavena
υποβλήθηκε από
GitHub
GitHub
γονέας
f7712bd5a8
υποβολή
6e2c871652
@@ -90,6 +90,8 @@ set(AMD_TEST_SRC
|
||||
AtomicsWithRandomActiveLanesInWavefront.cc
|
||||
fp16_ops.cc
|
||||
fp8_host.cc
|
||||
fp6_ocp.cc
|
||||
fp4_ocp.cc
|
||||
)
|
||||
|
||||
set(AMD_ARCH_SPEC_TEST_SRC
|
||||
|
||||
@@ -0,0 +1,449 @@
|
||||
/* Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#include <algorithm>
|
||||
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
#include <hip/hip_fp4.h>
|
||||
|
||||
template <typename Lambda, typename... Type>
|
||||
static __global__ void lambda_kernel_launch(Lambda l, Type... args) {
|
||||
l(args...);
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Basic test to convert given double type data to FP4 type with E2M1
|
||||
* format.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - /unit/deviceLib/fp4_ocp.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 6.5
|
||||
*/
|
||||
TEST_CASE("Unit_all_fp4_from_double") {
|
||||
SECTION("sanityx1") {
|
||||
std::vector<double> inputs{-1.0, 0.0, 1.0};
|
||||
for (const auto input : inputs) {
|
||||
__hip_fp4_e2m1 fp4(input);
|
||||
double ret = fp4;
|
||||
INFO("Original: " << input << " Return: " << ret);
|
||||
REQUIRE(ret == input);
|
||||
}
|
||||
}
|
||||
SECTION("sanityx2") {
|
||||
std::vector<double2> inputs{{-1.0, 1.0}, {-2.0, 2.0}};
|
||||
for (const auto input : inputs) {
|
||||
__hip_fp4x2_e2m1 fp4x2(input);
|
||||
double2 ret = fp4x2;
|
||||
INFO("Original: " << input.x << " Return: " << ret.x);
|
||||
INFO("Original: " << input.y << " Return: " << ret.y);
|
||||
|
||||
REQUIRE(ret.x == input.x);
|
||||
REQUIRE(ret.y == input.y);
|
||||
}
|
||||
}
|
||||
SECTION("sanityx4") {
|
||||
std::vector<double4> inputs{
|
||||
{-1.0, 0.5, 1.5, 1.0}, {-2.0, 0.5, 1.5, 2.0}, {-3.0, 0.5, 1.5, 3.0}};
|
||||
for (const auto &input : inputs) {
|
||||
__hip_fp4x4_e2m1 fp4x4(input);
|
||||
double4 ret = fp4x4;
|
||||
INFO("Original: " << input.x << ", " << input.y << ", " << input.z << ", "
|
||||
<< input.w << " Return: " << ret.x << ", " << ret.y
|
||||
<< ret.z << ", " << ret.w);
|
||||
REQUIRE(ret.x == input.x);
|
||||
REQUIRE(ret.y == input.y);
|
||||
REQUIRE(ret.z == input.z);
|
||||
REQUIRE(ret.w == input.w);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Basic test to convert given double data to FP4 type with E2M1
|
||||
* format in the device.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - /unit/deviceLib/fp4_ocp.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 6.5
|
||||
*/
|
||||
TEST_CASE("Unit_all_fp4_from_double_device") {
|
||||
SECTION("sanityx1") {
|
||||
auto fp4x1_l = [] __device__(double* inputs, float* outputs, size_t size) {
|
||||
int i = threadIdx.x;
|
||||
if (i < size) {
|
||||
__hip_fp4_e2m1 fp4(inputs[i]);
|
||||
outputs[i] = fp4;
|
||||
}
|
||||
};
|
||||
|
||||
std::vector<double> inputs{-1.0, 0.0, 1.0};
|
||||
double *d_in;
|
||||
float *d_out;
|
||||
HIP_CHECK(hipMalloc(&d_in, sizeof(double) * inputs.size()));
|
||||
HIP_CHECK(hipMalloc(&d_out, sizeof(float) * inputs.size()));
|
||||
|
||||
HIP_CHECK(hipMemcpy(d_in, inputs.data(), sizeof(double) * inputs.size(),
|
||||
hipMemcpyHostToDevice));
|
||||
lambda_kernel_launch<<<1, 32>>>(fp4x1_l, d_in, d_out, inputs.size());
|
||||
std::vector<float> outputs(inputs.size(), 0.0f);
|
||||
HIP_CHECK(hipMemcpy(outputs.data(), d_out, sizeof(float) * inputs.size(),
|
||||
hipMemcpyDeviceToHost));
|
||||
|
||||
for (size_t i = 0; i < inputs.size(); i++) {
|
||||
INFO("Original: " << inputs[i] << " Output: " << outputs[i]);
|
||||
REQUIRE(inputs[i] == outputs[i]);
|
||||
}
|
||||
|
||||
HIP_CHECK(hipFree(d_in));
|
||||
HIP_CHECK(hipFree(d_out));
|
||||
}
|
||||
|
||||
SECTION("sanityx2") {
|
||||
auto fp4x2_l = [] __device__(double2 * inputs, float2 * outputs,
|
||||
size_t size) {
|
||||
int i = threadIdx.x;
|
||||
if (i < size) {
|
||||
__hip_fp4x2_e2m1 fp4(inputs[i]);
|
||||
outputs[i] = fp4;
|
||||
}
|
||||
};
|
||||
|
||||
std::vector<double2> inputs{
|
||||
{-1.0, 0.0}, {0.0, 1.0}, {1.0, -1.0}, {1.0, 0.0}, {0.0, -1.0}};
|
||||
double2 *d_in;
|
||||
float2 *d_out;
|
||||
HIP_CHECK(hipMalloc(&d_in, sizeof(double2) * inputs.size()));
|
||||
HIP_CHECK(hipMalloc(&d_out, sizeof(float2) * inputs.size()));
|
||||
|
||||
HIP_CHECK(hipMemcpy(d_in, inputs.data(), sizeof(double2) * inputs.size(),
|
||||
hipMemcpyHostToDevice));
|
||||
lambda_kernel_launch<<<1, 32>>>(fp4x2_l, d_in, d_out, inputs.size());
|
||||
std::vector<float2> outputs(inputs.size());
|
||||
HIP_CHECK(hipMemcpy(outputs.data(), d_out, sizeof(float2) * inputs.size(),
|
||||
hipMemcpyDeviceToHost));
|
||||
|
||||
for (size_t i = 0; i < inputs.size(); i++) {
|
||||
INFO("Original: " << inputs[i].x << ", " << inputs[i].y << " Output: "
|
||||
<< outputs[i].x << ", " << outputs[i].y);
|
||||
REQUIRE(inputs[i].x == outputs[i].x);
|
||||
REQUIRE(inputs[i].y == outputs[i].y);
|
||||
}
|
||||
|
||||
HIP_CHECK(hipFree(d_in));
|
||||
HIP_CHECK(hipFree(d_out));
|
||||
}
|
||||
|
||||
SECTION("sanityx4") {
|
||||
auto fp4x4_l = [] __device__(double4 * inputs, float4 * outputs,
|
||||
size_t size) {
|
||||
int i = threadIdx.x;
|
||||
if (i < size) {
|
||||
__hip_fp4x4_e2m1 fp4(inputs[i]);
|
||||
outputs[i] = fp4;
|
||||
}
|
||||
};
|
||||
|
||||
std::vector<double4> inputs{
|
||||
{-1.0, 0.0, 1.0, 0.5}, {0.0, 1.0, -0.5, -1.0}, {1.0, 0.0, 1.0, -1.0}};
|
||||
double4 *d_in;
|
||||
float4 *d_out;
|
||||
HIP_CHECK(hipMalloc(&d_in, sizeof(double4) * inputs.size()));
|
||||
HIP_CHECK(hipMalloc(&d_out, sizeof(float4) * inputs.size()));
|
||||
|
||||
HIP_CHECK(hipMemcpy(d_in, inputs.data(), sizeof(double4) * inputs.size(),
|
||||
hipMemcpyHostToDevice));
|
||||
lambda_kernel_launch<<<1, 32>>>(fp4x4_l, d_in, d_out, inputs.size());
|
||||
std::vector<float4> outputs(inputs.size());
|
||||
HIP_CHECK(hipMemcpy(outputs.data(), d_out, sizeof(float4) * inputs.size(),
|
||||
hipMemcpyDeviceToHost));
|
||||
|
||||
for (size_t i = 0; i < inputs.size(); i++) {
|
||||
INFO("Original: " << inputs[i].x << ", " << inputs[i].y << ", "
|
||||
<< inputs[i].z << ", " << inputs[i].w << " Output: "
|
||||
<< outputs[i].x << ", " << outputs[i].y
|
||||
<< ", " << outputs[i].z << ", " << outputs[i].w);
|
||||
REQUIRE(inputs[i].x == outputs[i].x);
|
||||
REQUIRE(inputs[i].y == outputs[i].y);
|
||||
REQUIRE(inputs[i].z == outputs[i].z);
|
||||
REQUIRE(inputs[i].w == outputs[i].w);
|
||||
}
|
||||
|
||||
HIP_CHECK(hipFree(d_in));
|
||||
HIP_CHECK(hipFree(d_out));
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Basic test to convert given signed interger data to FP4 type with E2M1
|
||||
* format.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - /unit/deviceLib/fp4_ocp.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 6.5
|
||||
*/
|
||||
TEMPLATE_TEST_CASE("Unit_all_fp4_from_interger_data", "", int, long int,
|
||||
long long int, short int) {
|
||||
SECTION("Fp4 with e2m1") {
|
||||
std::vector<TestType> input{-1, 0, 1};
|
||||
for (const auto val : input) {
|
||||
__hip_fp4_e2m1 fp4(val);
|
||||
float ret = fp4;
|
||||
INFO("In: " << val);
|
||||
INFO("Out: " << ret);
|
||||
REQUIRE(ret == val);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Basic test to convert given unsigned integer data to FP4 type with E2M1
|
||||
* format.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - /unit/deviceLib/fp4_ocp.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 6.5
|
||||
*/
|
||||
TEMPLATE_TEST_CASE("Unit_all_fp4_from__unsigned_integer_data", "", unsigned int,
|
||||
unsigned long int, unsigned long long int,
|
||||
unsigned short int) {
|
||||
SECTION("Fp4 with e2m1") {
|
||||
std::vector<TestType> input{1, 2, 3};
|
||||
for (const auto val : input) {
|
||||
__hip_fp4_e2m1 fp4(val);
|
||||
float ret = fp4;
|
||||
INFO("In: " << val);
|
||||
INFO("Out: " << ret);
|
||||
REQUIRE(ret == val);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Basic test to convert given signed interger data to FP4 type in device with E2M1
|
||||
* format.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - /unit/deviceLib/fp4_ocp.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 6.5
|
||||
*/
|
||||
|
||||
TEMPLATE_TEST_CASE("Unit_all_fp4_from_integer_data_device", "", int,
|
||||
long int, long long int,
|
||||
short int) {
|
||||
std::vector<float> all_fp4{-6.0f, -4.0f, -3.0f, -2.0f, -1.5f,
|
||||
-1.0f, -0.5f, 0.0f, 0.5f, 1.0f,
|
||||
1.5f, 2.0f, 3.0f, 4.0f, 6.0f};
|
||||
auto fp4x1_l = [] __device__(TestType *inputs, float *outputs, size_t size) {
|
||||
int i = threadIdx.x;
|
||||
if (i < size) {
|
||||
__hip_fp4_e2m1 fp4(inputs[i]);
|
||||
outputs[i] = fp4;
|
||||
}
|
||||
};
|
||||
|
||||
std::vector<TestType> inputs;
|
||||
inputs.reserve(30);
|
||||
for (int i = 0; i <= 6; i += 1) {
|
||||
inputs.push_back(i);
|
||||
}
|
||||
TestType *d_in;
|
||||
float *d_out;
|
||||
HIP_CHECK(hipMalloc(&d_in, sizeof(TestType) * inputs.size()));
|
||||
HIP_CHECK(hipMalloc(&d_out, sizeof(float) * inputs.size()));
|
||||
|
||||
HIP_CHECK(hipMemcpy(d_in, inputs.data(), sizeof(TestType) * inputs.size(),
|
||||
hipMemcpyHostToDevice));
|
||||
lambda_kernel_launch<<<1, 32>>>(fp4x1_l, d_in, d_out, inputs.size());
|
||||
|
||||
std::vector<float> outputs(inputs.size(), 0.0f);
|
||||
HIP_CHECK(hipMemcpy(outputs.data(), d_out, sizeof(float) * inputs.size(),
|
||||
hipMemcpyDeviceToHost));
|
||||
|
||||
for (size_t i = 0; i < inputs.size(); i++) {
|
||||
auto lbound = std::lower_bound(all_fp4.begin(), all_fp4.end(), outputs[i]);
|
||||
INFO("Original: " << inputs[i] << " Output: " << *lbound);
|
||||
REQUIRE(*lbound == outputs[i]);
|
||||
}
|
||||
|
||||
HIP_CHECK(hipFree(d_in));
|
||||
HIP_CHECK(hipFree(d_out));
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Basic test to convert given unsigned interger data to FP4 type in device with E2M1
|
||||
* format.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - /unit/deviceLib/fp4_ocp.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 6.5
|
||||
*/
|
||||
TEMPLATE_TEST_CASE("Unit_all_fp4_from__unsigned_integer_data_device", "",
|
||||
unsigned int, unsigned long int, unsigned long long int,
|
||||
unsigned short int) {
|
||||
std::vector<float> all_fp4{-6.0f, -4.0f, -3.0f, -2.0f, -1.5f,
|
||||
-1.0f, -0.5f, 0.0f, 0.5f, 1.0f,
|
||||
1.5f, 2.0f, 3.0f, 4.0f, 6.0f};
|
||||
auto fp4x1_l = [] __device__(TestType *inputs, float *outputs, size_t size) {
|
||||
int i = threadIdx.x;
|
||||
if (i < size) {
|
||||
__hip_fp4_e2m1 fp4(inputs[i]);
|
||||
outputs[i] = fp4;
|
||||
}
|
||||
};
|
||||
|
||||
std::vector<TestType> inputs;
|
||||
inputs.reserve(30);
|
||||
for (int i = -6; i <= 6; i += 1) {
|
||||
inputs.push_back(i);
|
||||
}
|
||||
TestType *d_in;
|
||||
float *d_out;
|
||||
HIP_CHECK(hipMalloc(&d_in, sizeof(TestType) * inputs.size()));
|
||||
HIP_CHECK(hipMalloc(&d_out, sizeof(float) * inputs.size()));
|
||||
|
||||
HIP_CHECK(hipMemcpy(d_in, inputs.data(), sizeof(TestType) * inputs.size(),
|
||||
hipMemcpyHostToDevice));
|
||||
lambda_kernel_launch<<<1, 32>>>(fp4x1_l, d_in, d_out, inputs.size());
|
||||
|
||||
std::vector<float> outputs(inputs.size(), 0.0f);
|
||||
HIP_CHECK(hipMemcpy(outputs.data(), d_out, sizeof(float) * inputs.size(),
|
||||
hipMemcpyDeviceToHost));
|
||||
|
||||
for (size_t i = 0; i < inputs.size(); i++) {
|
||||
auto lbound = std::lower_bound(all_fp4.begin(), all_fp4.end(), outputs[i]);
|
||||
INFO("Original: " << inputs[i] << " Output: " << *lbound);
|
||||
REQUIRE(*lbound == outputs[i]);
|
||||
}
|
||||
|
||||
HIP_CHECK(hipFree(d_in));
|
||||
HIP_CHECK(hipFree(d_out));
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Basic test to convert given double type data to FP4 type with E2M3 and
|
||||
* E3M2 formats.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - /unit/deviceLib/fp4_ocp.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 6.5
|
||||
*/
|
||||
TEST_CASE("Unit_ocp_fp4_from_double_full_range_host") {
|
||||
std::vector<double> in;
|
||||
in.reserve(30);
|
||||
for (double i = -6.0; i <= 6.0; i += 0.5) {
|
||||
in.push_back(i);
|
||||
}
|
||||
|
||||
std::vector<float> expected{-6.0f, -6.0f, -4.0f, -4.0f, -4.0f, -4.0f, -3.0f,
|
||||
-2.0f, -2.0f, -1.5f, -1.0f, -0.5f, 0.0f, 0.5f,
|
||||
1.0f, 1.5f, 2.0f, 2.0f, 3.0f, 4.0f, 4.0f,
|
||||
4.0f, 4.0f, 6.0f, 6.0f};
|
||||
|
||||
for (size_t i = 0; i < in.size(); i++) {
|
||||
__hip_fp4_e2m1 fp4(in[i]);
|
||||
float fp32 = fp4;
|
||||
INFO("Original: " << in[i] << " Output: " << fp32
|
||||
<< " Expected: " << expected[i]);
|
||||
REQUIRE(expected[i] == fp32);
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Basic test to convert given double type data to FP4 type in device with
|
||||
* E2M3 and E3M2 formats.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - /unit/deviceLib/fp4_ocp.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 6.5
|
||||
*/
|
||||
|
||||
TEST_CASE("Unit_ocp_fp4_from_double_full_range_device") {
|
||||
std::vector<float> all_fp4{-6.0f, -4.0f, -3.0f, -2.0f, -1.5f,
|
||||
-1.0f, -0.5f, 0.0f, 0.5f, 1.0f,
|
||||
1.5f, 2.0f, 3.0f, 4.0f, 6.0f};
|
||||
auto fp4x1_l = [] __device__(double *inputs, float *outputs, size_t size) {
|
||||
int i = threadIdx.x;
|
||||
if (i < size) {
|
||||
__hip_fp4_e2m1 fp4(inputs[i]);
|
||||
outputs[i] = fp4;
|
||||
}
|
||||
};
|
||||
|
||||
std::vector<double> inputs;
|
||||
inputs.reserve(30);
|
||||
for (double i = -6.0; i <= 6.0; i += 0.5) {
|
||||
inputs.push_back(i);
|
||||
}
|
||||
double *d_in;
|
||||
float *d_out;
|
||||
HIP_CHECK(hipMalloc(&d_in, sizeof(double) * inputs.size()));
|
||||
HIP_CHECK(hipMalloc(&d_out, sizeof(float) * inputs.size()));
|
||||
|
||||
HIP_CHECK(hipMemcpy(d_in, inputs.data(), sizeof(double) * inputs.size(),
|
||||
hipMemcpyHostToDevice));
|
||||
lambda_kernel_launch<<<1, 32>>>(fp4x1_l, d_in, d_out, inputs.size());
|
||||
|
||||
std::vector<float> outputs(inputs.size(), 0.0f);
|
||||
HIP_CHECK(hipMemcpy(outputs.data(), d_out, sizeof(float) * inputs.size(),
|
||||
hipMemcpyDeviceToHost));
|
||||
|
||||
for (size_t i = 0; i < inputs.size(); i++) {
|
||||
auto lbound = std::lower_bound(all_fp4.begin(), all_fp4.end(), outputs[i]);
|
||||
INFO("Original: " << inputs[i] << " Output: " << *lbound);
|
||||
REQUIRE(*lbound == outputs[i]);
|
||||
}
|
||||
|
||||
HIP_CHECK(hipFree(d_in));
|
||||
HIP_CHECK(hipFree(d_out));
|
||||
}
|
||||
@@ -0,0 +1,331 @@
|
||||
/* Copyright (c) 2025 Advanced Micro Devices, Inc. All rights reserved.
|
||||
|
||||
Permission is hereby granted, free of charge, to any person obtaining a copy
|
||||
of this software and associated documentation files (the "Software"), to deal
|
||||
in the Software without restriction, including without limitation the rights
|
||||
to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
||||
copies of the Software, and to permit persons to whom the Software is
|
||||
furnished to do so, subject to the following conditions:
|
||||
|
||||
The above copyright notice and this permission notice shall be included in
|
||||
all copies or substantial portions of the Software.
|
||||
|
||||
THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
#include <hip/hip_fp6.h>
|
||||
#include <hip_test_common.hh>
|
||||
|
||||
#include <bitset>
|
||||
#include <type_traits>
|
||||
#include <vector>
|
||||
|
||||
std::vector<__hip_fp6_storage_t> get_all_fp6_ocp_nums() {
|
||||
std::vector<__hip_fp6_storage_t> ret;
|
||||
constexpr unsigned short max_fp6_num = 0b0011'1111;
|
||||
ret.reserve(max_fp6_num + 1);
|
||||
|
||||
for (unsigned short i = 0; i <= max_fp6_num; i++) {
|
||||
ret.push_back(static_cast<__hip_fp6_storage_t>(i));
|
||||
}
|
||||
return ret;
|
||||
}
|
||||
|
||||
template <bool is_e2m3, typename T>
|
||||
__global__ void Type_to_fp6(T* f, __hip_fp6_storage_t* res, size_t size) {
|
||||
auto i = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (i < size) {
|
||||
if constexpr (is_e2m3) {
|
||||
__hip_fp6_e2m3 tmp(f[i]);
|
||||
res[i] = tmp.__x;
|
||||
} else {
|
||||
__hip_fp6_e3m2 tmp(f[i]);
|
||||
res[i] = tmp.__x;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Basic test to convert given interger values to FP6 type in the host
|
||||
* with E2M3 and E3M2 formats.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - /unit/deviceLib/fp6_ocp.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 6.5
|
||||
*/
|
||||
TEMPLATE_TEST_CASE("Unit_all_fp6_ocp_vector_cvt_interger_data", "", int,
|
||||
long int, long long int, short int) {
|
||||
SECTION("Fp6 with e2m3") {
|
||||
std::vector<TestType> input = {0, 1, 2, 3, 4, 5, 6, 7, -0,
|
||||
-1, -2, -3, -4, -5, -6, -7 };
|
||||
for (const auto val : input) {
|
||||
__hip_fp6_e2m3 fp6(val);
|
||||
float ret = fp6;
|
||||
INFO("In: " << val);
|
||||
INFO("Out: " << ret);
|
||||
REQUIRE(ret == val);
|
||||
}
|
||||
}
|
||||
SECTION("Fp6 with e3m2 ") {
|
||||
std::vector<TestType> input = {0, 1, 2, 3, 4, 5, 6, 7, 8, 10, 12, 14,
|
||||
16, 20, 24, 28, -0, -1, -2, -3, -4, -5, -6,
|
||||
-7, -8, -10, -12, -14, -16, -20, -24, -28};
|
||||
for (const auto val : input) {
|
||||
__hip_fp6_e3m2 fp6(val);
|
||||
float ret = fp6;
|
||||
INFO("In: " << val);
|
||||
INFO("Out: " << ret);
|
||||
REQUIRE(ret == val);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Basic test to convert given unsigned interger values to FP6 type in the host
|
||||
* with E2M3 and E3M2 formats.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - /unit/deviceLib/fp6_ocp.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 6.5
|
||||
*/
|
||||
TEMPLATE_TEST_CASE("Unit_all_fp6_ocp_vector_cvt_unsigned_interger_data", "",
|
||||
int, long int, long long int, short int) {
|
||||
SECTION("Fp6 with e2m3") {
|
||||
std::vector<TestType> input = {0, 1, 2, 3, 4, 5, 6, 7};
|
||||
for (const auto val : input) {
|
||||
__hip_fp6_e2m3 fp6(val);
|
||||
float ret = fp6;
|
||||
INFO("In: " << val);
|
||||
INFO("Out: " << ret);
|
||||
REQUIRE(ret == val);
|
||||
}
|
||||
}
|
||||
SECTION("Fp6 with e3m2 ") {
|
||||
std::vector<TestType> input = {0, 1, 2, 3, 4, 5, 6, 7, 8, 10, 12, 14,
|
||||
16, 20, 24, 28};
|
||||
for (const auto val : input) {
|
||||
__hip_fp6_e3m2 fp6(val);
|
||||
float ret = fp6;
|
||||
INFO("In: " << val);
|
||||
INFO("Out: " << ret);
|
||||
REQUIRE(ret == val);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Basic test to convert given unsigned interger values to FP6 type in the device
|
||||
* with E2M3 and E3M2 formats.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - /unit/deviceLib/fp6_ocp.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 6.5
|
||||
*/
|
||||
TEMPLATE_TEST_CASE("Unit_all_fp6_ocp_vector_cvt_unsigned_integer_device", "",
|
||||
unsigned int, unsigned long int, unsigned long long int,
|
||||
unsigned short int) {
|
||||
bool is_e2m3 = GENERATE(true, false);
|
||||
std::vector<TestType> f_vals;
|
||||
std::vector<__hip_fp6_storage_t> all_vals;
|
||||
constexpr TestType lhs = 0;
|
||||
constexpr TestType rhs = 64;
|
||||
constexpr TestType step = 1;
|
||||
|
||||
f_vals.reserve(500);
|
||||
all_vals.reserve(500);
|
||||
|
||||
for (TestType fval = lhs; fval <= rhs; fval += step) {
|
||||
if (is_e2m3) {
|
||||
__hip_fp6_e2m3 tmp(fval);
|
||||
all_vals.push_back(tmp.__x);
|
||||
} else {
|
||||
__hip_fp6_e3m2 tmp(fval);
|
||||
all_vals.push_back(tmp.__x);
|
||||
}
|
||||
f_vals.push_back(fval);
|
||||
}
|
||||
|
||||
TestType* d_f_vals;
|
||||
__hip_fp6_storage_t* d_res;
|
||||
|
||||
HIP_CHECK(hipMalloc(&d_f_vals, sizeof(TestType) * f_vals.size()));
|
||||
HIP_CHECK(hipMalloc(&d_res, sizeof(__hip_fp6_storage_t) * f_vals.size()));
|
||||
|
||||
HIP_CHECK(hipMemcpy(d_f_vals, f_vals.data(), sizeof(TestType) *
|
||||
f_vals.size(), hipMemcpyHostToDevice));
|
||||
|
||||
auto fp6_kernel = is_e2m3 ? Type_to_fp6<true, TestType> :
|
||||
Type_to_fp6<false, TestType>;
|
||||
fp6_kernel<<<(f_vals.size() / 64) + 1, 64>>>(d_f_vals, d_res, f_vals.size());
|
||||
|
||||
std::vector<__hip_fp6_storage_t>
|
||||
final_res(f_vals.size(), static_cast<__hip_fp6_storage_t>(0));
|
||||
|
||||
HIP_CHECK(hipMemcpy(final_res.data(), d_res,
|
||||
sizeof(__hip_fp6_storage_t) * final_res.size(), hipMemcpyDeviceToHost));
|
||||
|
||||
for (size_t i = 0; i < final_res.size(); i++) {
|
||||
INFO("Checking: " << f_vals[i] << " for: " << (is_e2m3 ? "e2m3" : "e3m2")
|
||||
<< " original: " << (int)all_vals[i]
|
||||
<< " convert back: " << (int)final_res[i] << " Idx : "
|
||||
<< i);
|
||||
TestType gpu_cvt_res = 0, cpu_cvt_res = 0;
|
||||
float gpu_cvt_res_ = 0.0, cpu_cvt_res_ = 0.0;
|
||||
if (is_e2m3) {
|
||||
__hip_fp6_e2m3 gtmp;
|
||||
gtmp.__x = final_res[i];
|
||||
gpu_cvt_res_ = gtmp;
|
||||
gpu_cvt_res = gpu_cvt_res_;
|
||||
__hip_fp6_e2m3 ctmp;
|
||||
ctmp.__x = all_vals[i];
|
||||
cpu_cvt_res_ = ctmp;
|
||||
cpu_cvt_res = cpu_cvt_res_;
|
||||
} else {
|
||||
__hip_fp6_e3m2 gtmp;
|
||||
gtmp.__x = final_res[i];
|
||||
gpu_cvt_res_ = gtmp;
|
||||
gpu_cvt_res = gpu_cvt_res_;
|
||||
__hip_fp6_e3m2 ctmp;
|
||||
ctmp.__x = all_vals[i];
|
||||
cpu_cvt_res_ = ctmp;
|
||||
cpu_cvt_res = cpu_cvt_res_;
|
||||
}
|
||||
INFO("cpu cvt val: " << cpu_cvt_res << " gpu cvt val: " << gpu_cvt_res);
|
||||
REQUIRE(cpu_cvt_res == gpu_cvt_res);
|
||||
}
|
||||
|
||||
HIP_CHECK(hipFree(d_f_vals));
|
||||
HIP_CHECK(hipFree(d_res));
|
||||
}
|
||||
|
||||
/**
|
||||
* Test Description
|
||||
* ------------------------
|
||||
* - Basic test to convert given interger values to FP6 type in device with
|
||||
* E2M3 and E3M2 formats.
|
||||
* Test source
|
||||
* ------------------------
|
||||
* - /unit/deviceLib/fp6_ocp.cc
|
||||
* Test requirements
|
||||
* ------------------------
|
||||
* - HIP_VERSION >= 6.5
|
||||
*/
|
||||
|
||||
TEMPLATE_TEST_CASE("Unit_all_fp6_ocp_vector_cvt_interger_data_device", "", int,
|
||||
long int, long long int, short int) {
|
||||
bool is_e2m3 = GENERATE(true, false);
|
||||
std::vector<TestType> f_vals;
|
||||
std::vector<__hip_fp6_storage_t> all_vals;
|
||||
SECTION("all representable numbers") {
|
||||
all_vals = get_all_fp6_ocp_nums();
|
||||
f_vals.reserve(all_vals.size());
|
||||
|
||||
for (const auto& fp6 : all_vals) {
|
||||
TestType f = 0;
|
||||
float f_ = 0.0;
|
||||
if (is_e2m3) {
|
||||
__hip_fp6_e2m3 tmp;
|
||||
tmp.__x = fp6;
|
||||
f_ = tmp;
|
||||
f = f_;
|
||||
} else {
|
||||
__hip_fp6_e3m2 tmp;
|
||||
tmp.__x = fp6;
|
||||
f_ = tmp;
|
||||
f = f_;
|
||||
}
|
||||
f_vals.push_back(f);
|
||||
}
|
||||
}
|
||||
SECTION("Range stepped numbers") {
|
||||
constexpr TestType lhs = -30;
|
||||
constexpr TestType rhs = 30;
|
||||
constexpr TestType step = 1;
|
||||
|
||||
f_vals.reserve(500);
|
||||
all_vals.reserve(500);
|
||||
|
||||
for (TestType fval = lhs; fval <= rhs; fval += step) {
|
||||
if (is_e2m3) {
|
||||
__hip_fp6_e2m3 tmp(fval);
|
||||
all_vals.push_back(tmp.__x);
|
||||
} else {
|
||||
__hip_fp6_e3m2 tmp(fval);
|
||||
all_vals.push_back(tmp.__x);
|
||||
}
|
||||
f_vals.push_back(fval);
|
||||
}
|
||||
}
|
||||
|
||||
TestType* d_f_vals;
|
||||
__hip_fp6_storage_t* d_res;
|
||||
|
||||
HIP_CHECK(hipMalloc(&d_f_vals, sizeof(TestType) * f_vals.size()));
|
||||
HIP_CHECK(hipMalloc(&d_res, sizeof(__hip_fp6_storage_t) * f_vals.size()));
|
||||
|
||||
HIP_CHECK(hipMemcpy(d_f_vals, f_vals.data(),
|
||||
sizeof(TestType) *f_vals.size(), hipMemcpyHostToDevice));
|
||||
|
||||
auto fp6_kernel = is_e2m3 ? Type_to_fp6<true, TestType> :
|
||||
Type_to_fp6<false, TestType>;
|
||||
fp6_kernel<<<(f_vals.size() / 64) + 1, 64>>>(d_f_vals, d_res, f_vals.size());
|
||||
|
||||
std::vector<__hip_fp6_storage_t> final_res(f_vals.size(),
|
||||
static_cast<__hip_fp6_storage_t>(0));
|
||||
|
||||
HIP_CHECK(hipMemcpy(final_res.data(), d_res, sizeof(__hip_fp6_storage_t) *
|
||||
final_res.size(), hipMemcpyDeviceToHost));
|
||||
|
||||
for (size_t i = 0; i < final_res.size(); i++) {
|
||||
INFO("Checking: " << f_vals[i] << " for: " << (is_e2m3 ? "e2m3" : "e3m2")
|
||||
<< " original: " << (int)all_vals[i]
|
||||
<< " convert back: " << (int)final_res[i] << " Idx : "
|
||||
<< i);
|
||||
TestType gpu_cvt_res = 0.0f, cpu_cvt_res = 0.0f;
|
||||
float gpu_cvt_res_ = 0.0, cpu_cvt_res_ = 0.0;
|
||||
if (is_e2m3) {
|
||||
__hip_fp6_e2m3 gtmp;
|
||||
gtmp.__x = final_res[i];
|
||||
gpu_cvt_res_ = gtmp;
|
||||
gpu_cvt_res = gpu_cvt_res_;
|
||||
__hip_fp6_e2m3 ctmp;
|
||||
ctmp.__x = all_vals[i];
|
||||
cpu_cvt_res_ = ctmp;
|
||||
cpu_cvt_res = cpu_cvt_res_;
|
||||
} else {
|
||||
__hip_fp6_e3m2 gtmp;
|
||||
gtmp.__x = final_res[i];
|
||||
gpu_cvt_res_ = gtmp;
|
||||
gpu_cvt_res = gpu_cvt_res_;
|
||||
__hip_fp6_e3m2 ctmp;
|
||||
ctmp.__x = all_vals[i];
|
||||
cpu_cvt_res_ = ctmp;
|
||||
cpu_cvt_res = cpu_cvt_res_;
|
||||
}
|
||||
|
||||
INFO("cpu cvt val: " << cpu_cvt_res << " gpu cvt val: " << gpu_cvt_res);
|
||||
REQUIRE(cpu_cvt_res == gpu_cvt_res);
|
||||
}
|
||||
|
||||
HIP_CHECK(hipFree(d_f_vals));
|
||||
HIP_CHECK(hipFree(d_res));
|
||||
}
|
||||
|
||||
Αναφορά σε νέο ζήτημα
Block a user