diff --git a/include/ck/tensor_operation/gpu/element/unary_element_wise_operation.hpp b/include/ck/tensor_operation/gpu/element/unary_element_wise_operation.hpp index 0cf81038f8..6956233a99 100644 --- a/include/ck/tensor_operation/gpu/element/unary_element_wise_operation.hpp +++ b/include/ck/tensor_operation/gpu/element/unary_element_wise_operation.hpp @@ -278,10 +278,6 @@ struct PassThroughPack2 template __host__ __device__ void operator()(Y& y, const X& x) const; - __host__ __device__ constexpr void operator()(ck::bhalf2_t& y, const ck::float2_t& x) const - { - y = bf16x2_convert_rne(x); - } __host__ __device__ constexpr void operator()(half2_t& y, const f8x2_t& x) const { diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 806badf19f..cc6b8862b3 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -288,3 +288,4 @@ add_subdirectory(scatter_gather) add_subdirectory(pass_through) add_subdirectory(tensor_slice_transfer) add_subdirectory(packed_cast) +add_subdirectory(utils) diff --git a/test/packed_cast/test_packed_cast.cpp b/test/packed_cast/test_packed_cast.cpp index f6bbbcf92f..7fb6c0a401 100644 --- a/test/packed_cast/test_packed_cast.cpp +++ b/test/packed_cast/test_packed_cast.cpp @@ -25,7 +25,7 @@ __global__ void packed_cast(float x1, float x2, ck::bhalf2_t* output) (*output)[1] = dst_buffer.template AsType()[I1]; } -void run_test(const float x1, const float x2) +void run_single_case(const float x1, const float x2, std::function launch_kernel) { const auto& get_tolerance = [](const float test_val) -> float { @@ -45,7 +45,7 @@ void run_test(const float x1, const float x2) ck::bhalf2_t value_after_cast_host; hip_check_error(hipMalloc(&value_after_cast_dev, sizeof(ck::bhalf2_t))); - packed_cast<<<1, 1>>>(x1, x2, value_after_cast_dev); + launch_kernel(x1, x2, value_after_cast_dev); hip_check_error(hipGetLastError()); hip_check_error(hipMemcpy(&value_after_cast_host, value_after_cast_dev, @@ -57,14 +57,13 @@ void run_test(const float x1, const float x2) // Convert back to floats const float x1_actual = type_convert(value_after_cast_host[0]); const float x2_actual = type_convert(value_after_cast_host[1]); - std::cout << "x1: " << x1_actual << ", x2: " << x2_actual << std::endl; ASSERT_NEAR(x1_actual, x1, get_tolerance(x1)); ASSERT_NEAR(x2_actual, x2, get_tolerance(x2)); }; -TEST(packed_cast, vectorized_float2_to_bhalf2) +void run_test(std::function launch_kernel) { - // Test packed cast from bhalf2 to float2 + // Test packed cast from bhalf2 to float2 // Use values that are representable in bhalf2 as well as values that are not constexpr int num_vals = 15; @@ -120,14 +119,21 @@ TEST(packed_cast, vectorized_float2_to_bhalf2) const float exact_in_both_value = exact_in_both[i]; const float exact_fp32_not_bf16_value = exact_fp32_not_bf16[j]; - run_test(exact_in_both_value, exact_fp32_not_bf16_value); - run_test(exact_in_both_value, -exact_fp32_not_bf16_value); - run_test(-exact_in_both_value, exact_fp32_not_bf16_value); - run_test(-exact_in_both_value, -exact_fp32_not_bf16_value); - run_test(exact_fp32_not_bf16_value, exact_in_both_value); - run_test(exact_fp32_not_bf16_value, -exact_in_both_value); - run_test(-exact_fp32_not_bf16_value, exact_in_both_value); - run_test(-exact_fp32_not_bf16_value, -exact_in_both_value); + run_single_case(exact_in_both_value, exact_fp32_not_bf16_value, launch_kernel); + run_single_case(exact_in_both_value, -exact_fp32_not_bf16_value, launch_kernel); + run_single_case(-exact_in_both_value, exact_fp32_not_bf16_value, launch_kernel); + run_single_case(-exact_in_both_value, -exact_fp32_not_bf16_value, launch_kernel); + run_single_case(exact_fp32_not_bf16_value, exact_in_both_value, launch_kernel); + run_single_case(exact_fp32_not_bf16_value, -exact_in_both_value, launch_kernel); + run_single_case(-exact_fp32_not_bf16_value, exact_in_both_value, launch_kernel); + run_single_case(-exact_fp32_not_bf16_value, -exact_in_both_value, launch_kernel); } } } + +TEST(packed_cast, vectorized_float2_to_bhalf2) +{ + run_test([](float x1, float x2, ck::bhalf2_t* output) { + packed_cast<<<1, 1>>>(x1, x2, output); + }); +} diff --git a/test/tensor_slice_transfer/CMakeLists.txt b/test/tensor_slice_transfer/CMakeLists.txt index f9d2e4555b..b245cd1162 100644 --- a/test/tensor_slice_transfer/CMakeLists.txt +++ b/test/tensor_slice_transfer/CMakeLists.txt @@ -1,32 +1,4 @@ -set(test_target test_tensor_slice_transfer) - -add_executable(${test_target} test_tensor_slice_transfer.cpp) - -target_include_directories(${test_target} PRIVATE ${CMAKE_SOURCE_DIR}/include) -target_link_libraries(${test_target} gtest gtest_main hip::device) - -target_compile_definitions(${test_target} PRIVATE - CK_EXPERIMENTAL_BIT_INT_EXTENSION_INT4 -) - -# Disable warnings that conflict with GTest -target_compile_options(${test_target} PRIVATE - -Wno-global-constructors - -Wno-unused-result - -Wno-shadow -) - -set_target_properties(${test_target} PROPERTIES - CXX_STANDARD 17 - HIP_SEPARABLE_COMPILATION ON -) - -# Set GPU architecture for compilation -if(GPU_TARGETS) - set_property(TARGET ${test_target} PROPERTY HIP_ARCHITECTURES ${GPU_TARGETS}) -else() - set_property(TARGET ${test_target} PROPERTY HIP_ARCHITECTURES "gfx950") +add_gtest_executable(test_tensor_slice_transfer test_tensor_slice_transfer.cpp) +if(result EQUAL 0) + target_link_libraries(test_tensor_slice_transfer PRIVATE utility) endif() - -# Add test to CTest -add_test(NAME ${test_target} COMMAND $) diff --git a/test/utils/CMakeLists.txt b/test/utils/CMakeLists.txt new file mode 100644 index 0000000000..57e26956bc --- /dev/null +++ b/test/utils/CMakeLists.txt @@ -0,0 +1,4 @@ +add_gtest_executable(test_sequencies test_sequencies.cpp) +if(result EQUAL 0) + target_link_libraries(test_sequencies PRIVATE utility) +endif() diff --git a/test/utils/test_sequencies.cpp b/test/utils/test_sequencies.cpp new file mode 100644 index 0000000000..b5e8071122 --- /dev/null +++ b/test/utils/test_sequencies.cpp @@ -0,0 +1,43 @@ +// SPDX-License-Identifier: MIT +// Copyright (c) 2025, Advanced Micro Devices, Inc. All rights reserved. + +#include "gtest/gtest.h" + +#include "ck/utility/common_header.hpp" +#include "ck/utility/sequence.hpp" +#include "ck/utility/sequence_helper.hpp" +#include "ck/utility/number.hpp" +#include "ck/tensor_operation/gpu/thread/threadwise_tensor_slice_transfer_util.hpp" + +using namespace ck; + + +TEST(Sequence, TestCreatingAccessSequence) +{ + constexpr int DstVectorDim = 6; + constexpr int DstScalarPerVector = 2; + constexpr index_t nDim = 8; + + constexpr auto SliceLengths = Sequence<4, 8, 1, 1, 4, 1, 2, 1>{}; + //constexpr auto DimAccessOrder = Sequence<0, 1, 2, 3, 4, 5, 7, 6>{}; + + constexpr auto ScalarPerAccess = generate_sequence( + detail::lambda_scalar_per_access{}, Number{}); + + constexpr auto Rem = SliceLengths % ScalarPerAccess; + constexpr auto Zeros = Sequence<0, 0, 0, 0, 0, 0, 0, 0>{}; + + constexpr bool is_valid = Rem == Zeros; + + std::cout << "SliceLengths: "; + SliceLengths.Print(); + std::cout << std::endl; + std::cout << "ScalarPerAccess: "; + ScalarPerAccess.Print(); + std::cout << std::endl; + std::cout << "Rem: "; + Rem.Print(); + std::cout << std::endl; + + EXPECT_EQ(is_valid, true); +}