Add sycl equivalent to cuda events for profiling

codeplaysoftware · May 31, 2024 · 024d0c6 · 024d0c6
1 parent 681df7a
commit 024d0c6
Show file tree

Hide file tree

Showing 19 changed files with 207 additions and 1,037 deletions.
diff --git a/CMakeLists.txt b/CMakeLists.txt
@@ -161,10 +161,6 @@ set(CUTLASS_ENABLE_PERFORMANCE ${CUTLASS_ENABLE_PROFILER} CACHE BOOL "Enable CUT
 option(CUTLASS_ENABLE_DEBUG_PRINTS "Whether or not to enable debug prints in CUTLASS kernels" OFF)
 set(CUTLASS_ENABLE_BENCHMARKS ON CACHE BOOL "Enable CUTLASS Benchmarks")
 
-if (CUTLASS_ENABLE_SYCL)
-  set(CUTLASS_ENABLE_BENCHMARKS ON CACHE BOOL "Enable CUTLASS Benchmarks")
-endif()
-
 set(CUTLASS_ENABLE_TESTS ${CUTLASS_ENABLE_TESTS_INIT} CACHE BOOL "Enable CUTLASS Tests")
 set(CUTLASS_ENABLE_GTEST_UNIT_TESTS ${CUTLASS_ENABLE_TESTS} CACHE BOOL "Enable CUTLASS GTest-based Unit Tests")
 ################################################################################

diff --git a/benchmarks/CMakeLists.txt b/benchmarks/CMakeLists.txt
@@ -68,8 +68,6 @@ endfunction()
 
 if(SYCL_INTEL_TARGET)
   add_subdirectory(pvc)
-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)

diff --git a/benchmarks/ampere/CMakeLists.txt b/benchmarks/ampere/CMakeLists.txt
@@ -36,8 +36,3 @@ 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_fp32_tensor_op_fp32
-  bench_ampere_gemm_bf16_bf16_fp32_tensor_op_fp32.cpp
-)
diff --git a/benchmarks/common/benchmark_runner.hpp b/benchmarks/common/benchmark_runner.hpp
@@ -208,6 +208,16 @@ struct BenchmarkRunner {
       auto problem_shape_MNKL = cute::append<4>(problem_size, 1);
       auto [M, N, K, L] = problem_shape_MNKL;
 
+#if defined(CUTLASS_ENABLE_SYCL)
+      sycl::property_list prop = {
+              sycl::property::queue::in_order(),
+              sycl::property::queue::enable_profiling()
+      };
+
+      auto q = sycl::queue(syclcompat::get_default_context(), syclcompat::get_current_device(), prop);
+      syclcompat::set_default_queue(q);
+#endif
+
       stride_A = cutlass::make_cute_packed_stride(StrideA{}, cute::make_shape(M, K, L));
       stride_B = cutlass::make_cute_packed_stride(StrideB{}, cute::make_shape(N, K, L));
       stride_C = cutlass::make_cute_packed_stride(StrideC{}, cute::make_shape(M, N, L));

diff --git a/examples/14_ampere_tf32_tensorop_gemm/ampere_tf32_tensorop_gemm_cute.cpp b/examples/14_ampere_tf32_tensorop_gemm/ampere_tf32_tensorop_gemm_cute.cpp
@@ -121,6 +121,13 @@ run(Gemm_Op gemm_op)
 
 void test_gemm(int m, int n, int k)
 {
+  sycl::property_list prop = {
+          sycl::property::queue::in_order(),
+          sycl::property::queue::enable_profiling()
+  };
+
+  auto q = sycl::queue(syclcompat::get_default_context(), syclcompat::get_current_device(), prop);
+  syclcompat::set_default_queue(q);
 
   std::cout << "M = " << m << std::endl;
   std::cout << "N = " << n << std::endl;

diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt
@@ -164,7 +164,6 @@ else()
     57_hopper_grouped_gemm
     58_ada_fp8_gemm
     59_ampere_gather_scatter_conv
-    sycl
   )
 
     add_subdirectory(${EXAMPLE})

diff --git a/examples/cute/tutorial/sgemm_1_sycl.cpp b/examples/cute/tutorial/sgemm_1_sycl.cpp
@@ -376,6 +376,14 @@ int main(int argc, char** argv) {
   char transB = 'T';
   if (argc >= 6) sscanf(argv[5], "%c", &transB);
 
+  sycl::property_list prop = {
+          sycl::property::queue::in_order(),
+          sycl::property::queue::enable_profiling()
+  };
+
+  auto q = sycl::queue(syclcompat::get_default_context(), syclcompat::get_current_device(), prop);
+  syclcompat::set_default_queue(q);
+
   using TA = float;
   using TB = float;
   using TC = float;

diff --git a/examples/cute/tutorial/sgemm_2_sycl.cpp b/examples/cute/tutorial/sgemm_2_sycl.cpp
@@ -397,6 +397,14 @@ int main(int argc, char** argv) {
   char transB = 'T';
   if (argc >= 6) sscanf(argv[5], "%c", &transB);
 
+  sycl::property_list prop = {
+          sycl::property::queue::in_order(),
+          sycl::property::queue::enable_profiling()
+  };
+
+  auto q = sycl::queue(syclcompat::get_default_context(), syclcompat::get_current_device(), prop);
+  syclcompat::set_default_queue(q);
+
   using TA = float;
   using TB = float;
   using TC = float;

diff --git a/examples/cute/tutorial/sgemm_sm70_sycl.cpp b/examples/cute/tutorial/sgemm_sm70_sycl.cpp
@@ -385,6 +385,14 @@ int main(int argc, char** argv) {
   char transB = 'T';
   if (argc >= 6) sscanf(argv[5], "%c", &transB);
 
+  sycl::property_list prop = {
+          sycl::property::queue::in_order(),
+          sycl::property::queue::enable_profiling()
+  };
+
+  auto q = sycl::queue(syclcompat::get_default_context(), syclcompat::get_current_device(), prop);
+  syclcompat::set_default_queue(q);
+
   using TA = float;
   using TB = float;
   using TC = float;

diff --git a/examples/cute/tutorial/sgemm_sm80_sycl.cpp b/examples/cute/tutorial/sgemm_sm80_sycl.cpp
@@ -468,6 +468,14 @@ int main(int argc, char** argv) {
   char transB = 'T';
   if (argc >= 6) sscanf(argv[5], "%c", &transB);
 
+  sycl::property_list prop = {
+          sycl::property::queue::in_order(),
+          sycl::property::queue::enable_profiling()
+  };
+
+  auto q = sycl::queue(syclcompat::get_default_context(), syclcompat::get_current_device(), prop);
+  syclcompat::set_default_queue(q);
+
   using TA = float;
   using TB = float;
   using TC = float;

diff --git a/examples/sycl/CMakeLists.txt b/examples/sycl/CMakeLists.txt
@@ -29,6 +29,4 @@
 
 if(SYCL_INTEL_TARGET)
   add_subdirectory(pvc)
-else()
-  add_subdirectory(ampere)
 endif()
diff --git a/examples/sycl/ampere/CMakeLists.txt b/examples/sycl/ampere/CMakeLists.txt
diff --git a/examples/sycl/ampere/ampere_gemm_fp16_fp16_fp32_tensor_op_fp32.cu b/examples/sycl/ampere/ampere_gemm_fp16_fp16_fp32_tensor_op_fp32.cu