Commit 512666ff authored by Chao Liu's avatar Chao Liu
Browse files

clean

parent 83511d7e
...@@ -20,22 +20,10 @@ struct PassThrough ...@@ -20,22 +20,10 @@ struct PassThrough
__host__ __device__ void operator()(double& y, const double& x) const { y = x; } __host__ __device__ void operator()(double& y, const double& x) const { y = x; }
}; };
struct Gelu
{
__host__ __device__ void operator()(float& y, const float& x) const
{
// Y=0.5*X*(1+tanh(0.797885*X+0.035677*X*X*X))
const float a = float(0.035677) * x * x;
const float b = float(0.797885) + a;
const float c = b * x;
const float d = tanh(c);
const float e = float(1.0) + d;
y = float(0.5) * x * e;
}
};
struct FastGelu struct FastGelu
{ {
// https://paperswithcode.com/method/gelu
// y = 0.5*x*(1+tanh(sqrt(2/pi)*(x+0.044715*x^3)))
__host__ __device__ void operator()(float& y, const float& x) const __host__ __device__ void operator()(float& y, const float& x) const
{ {
const float u = float(2) * x * (float(0.035677) * x * x + float(0.797885)); const float u = float(2) * x * (float(0.035677) * x * x + float(0.797885));
......
...@@ -236,7 +236,6 @@ struct GridwiseGemmBiasActivation_k0mk1_k0nk1_mn_xdl_cshuffle ...@@ -236,7 +236,6 @@ struct GridwiseGemmBiasActivation_k0mk1_k0nk1_mn_xdl_cshuffle
auto c_grid_buf = make_dynamic_buffer<AddressSpaceEnum::Global>( auto c_grid_buf = make_dynamic_buffer<AddressSpaceEnum::Global>(
p_c_grid, c_grid_desc_mblock_mperblock_nblock_nperblock.GetElementSpaceSize()); p_c_grid, c_grid_desc_mblock_mperblock_nblock_nperblock.GetElementSpaceSize());
// divide block work by [M, N] // divide block work by [M, N]
const auto block_work_idx = const auto block_work_idx =
block_2_ctile_map.CalculateBottomIndex(make_multi_index(get_block_1d_id())); block_2_ctile_map.CalculateBottomIndex(make_multi_index(get_block_1d_id()));
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment