ROCm/composable_kernel
1
0
Fork 0
mirror of https://github.com/ROCm/composable_kernel.git synced 2026-05-15 02:27:57 +00:00
Code Issues Packages Projects Releases Wiki Activity

Merge commit 'daf71fb8e4442352b1c5bb0a7c5a3ecc9f7f0c5a' into develop

Browse Source
This commit is contained in:
assistant-librarian[bot]
2025-06-25 13:27:51 +00:00
parent c35d1ccde6
commit 605ce2526d
3 changed files with 0 additions and 15 deletions
Download Patch File Download Diff File Expand all files Collapse all files

6
include/ck/ck.hpp
View File

@@ -244,12 +244,6 @@
// workaround: compiler issue on gfx908
#define CK_WORKAROUND_SWDEV_388832 1
// workaround: compiler issue on gfx950
#define CK_WORKAROUND_FP32_TO_FP4_SR_CONVERSION 1
// workaround: compiler issue on gfx950
#define CK_TEMP_DISABLE_FP4_TESTS 1
// workaround: compiler issue on gfx950
#define CK_WORKAROUND_FP16_TO_FP8_CONVERSION 1

7
include/ck/utility/type_convert.hpp
View File

@@ -1500,16 +1500,9 @@ inline __host__ __device__ f4x2_t f4_convert_sr(float2_t x, float scale = 1.0f)
uint32_t bitwise;
f4x2_t f4x2_array[4];
} value{0};
// apply a temporary workaround for gfx950
#if CK_WORKAROUND_FP32_TO_FP4_SR_CONVERSION
uint8_t l = utils::sat_convert_to_type_sr<f4_t>(x[1] / scale, rng);
uint8_t h = utils::sat_convert_to_type_sr<f4_t>(x[0] / scale, rng);
value.bitwise = (h << 4) | l;
#else
// permute high bits and low bits to match the order of the original vector
value.bitwise = __builtin_amdgcn_cvt_scalef32_sr_pk_fp4_f32(
value.bitwise, float2_t{x[1], x[0]}, rng, scale, 0);
#endif // CK_WORKAROUND_FP32_TO_FP4_SR_CONVERSION
return value.f4x2_array[0];
#else
constexpr int seed = 1254739;

2
test/data_type/test_mx_fp4.cpp
View File

@@ -240,7 +240,6 @@ TEST(MXFP4, HostScaledConvert)
EXPECT_EQ(test_size, i);
}
#if !CK_TEMP_DISABLE_FP4_TESTS
__global__ void test_mx_fp4_device_scaled_convert(uint64_t N, float* p_test, uint64_t* p_completed)
{
test_mx_fp4_scaled_convert(N, p_test, p_completed);
@@ -540,4 +539,3 @@ TEST(MXFP4, DeviceF4x32ToF32x32ScaledConvert)
EXPECT_EQ(N, completed);
EXPECT_EQ(N, i);
}
#endif // CK_TEMP_DISABLE_FP4_TESTS