Add sycl equivalent to cuda events for profiling

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

Hide file tree

Showing 17 changed files with 217 additions and 645 deletions.
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/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