mirror of
https://github.com/ROCm/composable_kernel.git
synced 2026-05-14 02:02:46 +00:00
92b701bb16
* Add reduce2d new api * Prevent user use cross warp reduction * Fix bug of std caculation * Add rmsnorm2d * Add rmsnorm small example * Remove static assert to prevent compile fail * Add script to test performance and correctness * Add missing cmake change * refine naming * refine example of rmsnorm * Fix bug of rmsnorm * Refine naming * Fix cmake * clang format * Refine pipeline name * Add add_rmsnorm2d_rdquant kernel * Add reduce op * host verification * Fix bug of one pass pipeline * Refine tile size * Add two pass pipeline * Rename two pass to three pass * Fix bug of kSaveX == false * Add instance library * Add test script * Fix bug of x verification * Add save_x to trait * Add README * Move reduce2d into reduce folder * Fix bug of welford when number of m warp > 1 * remove reduncant comment * 1. move 06_rmsnorm2d to 10_rmsnorm2d 2. move 07_add_rmsnorm2d_rdquant to 11_add_rmsnorm2d_rdquant * clang format and add missing header * Add host validation of add + layernorm2d + rsquant * Revert "Add host validation of add + layernorm2d + rsquant" This reverts commit936cb45797. * Remove deprecated flag [ROCm/composable_kernel commit:3d60953477]
34 lines
1.1 KiB
C++
34 lines
1.1 KiB
C++
// SPDX-License-Identifier: MIT
|
|
// Copyright (c) 2018-2023, Advanced Micro Devices, Inc. All rights reserved.
|
|
|
|
#pragma once
|
|
|
|
#include "ck_tile/core.hpp"
|
|
#include "ck_tile/host/host_tensor.hpp"
|
|
#include <thread>
|
|
|
|
namespace ck_tile {
|
|
template <typename XDataType, typename ScaleDataType, typename QXDataType>
|
|
CK_TILE_HOST void reference_rowwise_quantization2d(const HostTensor<XDataType>& x_m_n,
|
|
const HostTensor<ScaleDataType>& scale_m,
|
|
HostTensor<QXDataType>& qx_m_n)
|
|
{
|
|
auto f = [&](auto m) {
|
|
const int N = x_m_n.mDesc.get_lengths()[1];
|
|
|
|
for(int n = 0; n < N; ++n)
|
|
{
|
|
auto v_x = x_m_n(m, n);
|
|
// scale = amax / 127 for int8
|
|
auto v_scale = type_convert<XDataType>(scale_m(m));
|
|
auto v_qx = v_x / v_scale;
|
|
qx_m_n(m, n) = saturates<QXDataType>{}(v_qx);
|
|
}
|
|
};
|
|
|
|
make_ParallelTensorFunctor(f,
|
|
scale_m.mDesc.get_lengths()[0])(std::thread::hardware_concurrency());
|
|
}
|
|
|
|
} // namespace ck_tile
|