Add bfloat-bfloat example

codeplaysoftware · May 31, 2024 · 4e2b2e8 · 4e2b2e8
1 parent b6ae4eb
commit 4e2b2e8
Show file tree

Hide file tree

Showing 4 changed files with 14 additions and 3 deletions.
diff --git a/benchmarks/CMakeLists.txt b/benchmarks/CMakeLists.txt
@@ -71,6 +71,3 @@ if(SYCL_INTEL_TARGET)
 else(SYCL_NVIDIA_TARGET OR NOT CUTLASS_ENABLE_SYCL)
   add_subdirectory(ampere)
 endif()
-if(SYCL_NVIDIA_TARGET OR NOT CUTLASS_ENABLE_SYCL)
-  add_subdirectory(ampere)
-endif()
diff --git a/benchmarks/ampere/CMakeLists.txt b/benchmarks/ampere/CMakeLists.txt
@@ -36,3 +36,8 @@ cutlass_benchmark_add_executable(
   bench_ampere_gemm_bf16_bf16_fp32_tensor_op_fp32
   bench_ampere_gemm_bf16_bf16_fp32_tensor_op_fp32.cu
 )
+
+cutlass_benchmark_add_executable(
+  bench_ampere_gemm_bf16_bf16_bf16_tensor_op_fp32
+  bench_ampere_gemm_bf16_bf16_bf16_tensor_op_fp32.cu
+)
diff --git a/...ere_gemm_fp16_fp16_fp32_tensor_op_fp32.cu → ...ere_gemm_bf16_bf16_bf16_tensor_op_fp32.cu b/...ere_gemm_fp16_fp16_fp32_tensor_op_fp32.cu → ...ere_gemm_bf16_bf16_bf16_tensor_op_fp32.cu
diff --git a/benchmarks/common/benchmark_runner.hpp b/benchmarks/common/benchmark_runner.hpp
@@ -155,6 +155,15 @@ struct BenchmarkRunner {
     cutlass::DeviceAllocation<ElementOutput> block_D;
     cutlass::DeviceAllocation<ElementOutput> block_ref_D;
 
+    ElementOutput epsilon;
+    ElementOutput nonzero_floor;
+
+    ExampleRunner() : epsilon(static_cast<ElementOutput>(0.1f)),
+                      nonzero_floor(static_cast<ElementOutput>(0.1f)) {};
+
+    ExampleRunner(ElementOutput epsilon, ElementOutput nonzeroFloor) :
+            epsilon(epsilon), nonzero_floor(nonzeroFloor) {}
+
     //
     // Methods
     //