Merge branch 'gfx950' into promote_ocp_fp8

include/ck/utility/data_type.hpp

@@ -1063,7 +1063,7 @@ struct non_native_vector_base<
         StaticallyIndexedArray<data_v, 1> dNx1;
     } data_;
 
-    __host__ __device__ constexpr non_native_vector_base(data_t a) : data_{data_v{a}} {}
+    __host__ __device__ constexpr non_native_vector_base(data_t a) : data_{data_v(a)} {}
     __host__ __device__ constexpr non_native_vector_base(T f)
         : non_native_vector_base(bit_cast<data_t>(f))
     {

test/data_type/test_custom_type.cpp

@@ -915,6 +915,10 @@ TEST(FP8OCP, TestAsType)
         ASSERT_EQ(left_vec.template AsType<f8_t>()(Number<i>{}),
                   ck::type_convert<f8_t>(test_vec.at(i)));
     });
+
+    ck::non_native_vector_base<ck::f8_ocp_t, 2> nnvb_f8x2(ck::type_convert<f8_t>(-10.0f));
+    ASSERT_EQ(nnvb_f8x2.template AsType<f8_t>()(Number<0>{}), ck::type_convert<f8_t>(-10.0f));
+    ASSERT_EQ(nnvb_f8x2.template AsType<f8_t>()(Number<1>{}), ck::type_convert<f8_t>(-10.0f));
 }
 
 TEST(FP8OCP, TestAsTypeReshape)
@@ -988,6 +992,10 @@ TEST(BF8OCP, TestAsType)
         ASSERT_EQ(left_vec.template AsType<bf8_t>()(Number<i>{}),
                   ck::type_convert<bf8_t>(test_vec.at(i)));
     });
+
+    ck::non_native_vector_base<bf8_t, 2> nnvb_bf8x2(ck::type_convert<bf8_t>(-10.0f));
+    ASSERT_EQ(nnvb_bf8x2.template AsType<bf8_t>()(Number<0>{}), ck::type_convert<bf8_t>(-10.0f));
+    ASSERT_EQ(nnvb_bf8x2.template AsType<bf8_t>()(Number<1>{}), ck::type_convert<bf8_t>(-10.0f));
 }
 
 TEST(BF8OCP, TestAsTypeReshape)

test/pool/test_avg_pool2d_fwd.cpp

@@ -138,7 +138,7 @@ TYPED_TEST_SUITE(AvgPool2D_BF16, AvgPool2D_BF16_Types);
 TYPED_TEST_SUITE(AvgPool2D_I8, AvgPool2D_I8_Types);
 TYPED_TEST_SUITE(AvgPool2D_F8, AvgPool2D_F8_Types);
 
-TYPED_TEST(AvgPool2D_F32, AvgPool2D_I8_Test) { this->Run(); }
+TYPED_TEST(AvgPool2D_F32, AvgPool2D_F32_Test) { this->Run(); }
 TYPED_TEST(AvgPool2D_F16, AvgPool2D_F16_Test) { this->Run(); }
 TYPED_TEST(AvgPool2D_BF16, AvgPool2D_BF16_Test) { this->Run(); }
 TYPED_TEST(AvgPool2D_I8, AvgPool2D_I8_Test) { this->Run(); }

test/pool/test_max_pool2d_fwd.cpp

@@ -143,7 +143,7 @@ TYPED_TEST_SUITE(MaxPool2D_BF16, MaxPool2D_BF16_Types);
 TYPED_TEST_SUITE(MaxPool2D_I8, MaxPool2D_I8_Types);
 TYPED_TEST_SUITE(MaxPool2D_F8, MaxPool2D_F8_Types);
 
-TYPED_TEST(MaxPool2D_F32, MaxPool2D_I8_Test) { this->Run(); }
+TYPED_TEST(MaxPool2D_F32, MaxPool2D_F32_Test) { this->Run(); }
 TYPED_TEST(MaxPool2D_F16, MaxPool2D_F16_Test) { this->Run(); }
 TYPED_TEST(MaxPool2D_BF16, MaxPool2D_BF16_Test) { this->Run(); }
 TYPED_TEST(MaxPool2D_I8, MaxPool2D_I8_Test) { this->Run(); }