Add test

test/unit/epilogue/thread/linear_combination.cu

@@ -29,6 +29,8 @@
 #include "../../common/cutlass_unit_test.h"
 
 #include "cutlass/epilogue/thread/linear_combination.h"
+#include "cutlass/epilogue/thread/linear_combination_gelu.h"
+#include "cutlass/epilogue/thread/activation.h"
 
 /////////////////////////////////////////////////////////////////////////////////////////////////
 
@@ -119,3 +121,47 @@ TEST(Epilogue_thread_linear_combination, device_side_f16_f32_ptr) {
 }
 
 /////////////////////////////////////////////////////////////////////////////////////////////////
+
+TEST(Epilogue_thread_linear_combination_gelu, device_side_f16_f16_ptr) {
+
+  using Element = cutlass::half_t;
+  using ElementOutput = cutlass::half_t;
+  int const kCount = 8;
+
+  using LinearCombination = cutlass::epilogue::thread::LinearCombinationGELU<
+    ElementOutput,
+    kCount,
+    Element,
+    Element>;
+
+  Element alpha = Element(1);
+  Element beta = Element(0);
+
+  typename LinearCombination::Params params(&alpha, &beta);
+
+  LinearCombination linear_combination_op(params);
+
+  cutlass::Array<ElementOutput, kCount> source;
+  cutlass::Array<Element, kCount> accum;
+
+  for (int i = 0; i < kCount; ++i) {
+    accum[i] = Element((float)i * 0.3f);
+    source[i] = ElementOutput(0);
+  }
+
+  cutlass::Array<ElementOutput, kCount> destination = linear_combination_op(accum, source);
+
+  const float sqrt2 = sqrtf(2.0f);
+  for (int i = 0; i < kCount; ++i) {
+    float scalar = (float)accum[i];
+    ElementOutput expected = ElementOutput(
+      0.5f * scalar * (1.0f + erff(scalar / sqrt2))
+    );
+
+    ElementOutput got = destination[i];
+    ElementOutput diff(fabs((float)(expected - got)));
+    EXPECT_TRUE(diff <= std::numeric_limits<cutlass::half_t>::epsilon());
+  }
+}
+
+/////////////////////////////////////////////////////////////////////////////////////////////////