WIP

celerity · Jan 10, 2024 · 05e48d8 · 05e48d8
1 parent a70e083
commit 05e48d8
Show file tree

Hide file tree

Showing 8 changed files with 381 additions and 33 deletions.
diff --git a/.github/workflows/simsycl_ci.yml b/.github/workflows/simsycl_ci.yml
@@ -13,33 +13,13 @@ jobs:
       fail-fast: false
 
       matrix:
-        os: [ubuntu-latest, windows-latest, macos-latest]
-        build_type: [Release, Debug]
-        c_compiler: [gcc, clang, cl]
+        os: [windows-latest]
+        build_type: [Debug]
+        c_compiler: [cl]
         include:
           - os: windows-latest
             c_compiler: cl
             cpp_compiler: cl
-          - os: ubuntu-latest
-            c_compiler: gcc
-            cpp_compiler: g++
-          - os: ubuntu-latest
-            c_compiler: clang
-            cpp_compiler: clang++
-          - os: macos-latest
-            c_compiler: gcc
-            cpp_compiler: g++-13
-        exclude:
-          - os: windows-latest
-            c_compiler: gcc
-          - os: windows-latest
-            c_compiler: clang
-          - os: ubuntu-latest
-            c_compiler: cl
-          - os: macos-latest
-            c_compiler: cl
-          - os: macos-latest
-            c_compiler: clang
 
     steps:
     - uses: actions/checkout@v3
@@ -79,8 +59,7 @@ jobs:
       # Build your program with the given configuration. Note that --config is needed because the default Windows generator is a multi-config generator (Visual Studio generator).
       run: cmake --build ${{ steps.strings.outputs.build-output-dir }} --config ${{ matrix.build_type }}
 
-    - name: Test
+    - name: Test deprecated parallel_for signature
       working-directory: ${{ steps.strings.outputs.build-output-dir }}
-      # Execute tests defined by the CMake configuration. Note that --build-config is needed because the default Windows generator is a multi-config generator (Visual Studio generator).
-      # See https://cmake.org/cmake/help/latest/manual/ctest.1.html for more detail
-      run: ctest --build-config ${{ matrix.build_type }}
+      timeout-minutes: 1
+      run: D:\a\SimSYCL\SimSYCL\build\test\Debug\tests.exe "Calls to the deprecated parallel_for signature are not ambiguous"
diff --git a/include/simsycl/detail/parallel_for.hh b/include/simsycl/detail/parallel_for.hh
@@ -68,7 +68,9 @@ void execute_parallel_for(const sycl::range<Dimensions> &range, const Offset &of
     const KernelFunc &func,
     Reducers &...reducers) //
 {
+    printf("execute_parallel_for 71\n");
     register_kernel_on_static_construction<KernelName, KernelFunc>();
+    printf("execute_parallel_for 73\n");
 
     simple_kernel<Dimensions, with_offset_v<Offset>> kernel;
     if constexpr(std::is_invocable_v<const KernelFunc, sycl::item<Dimensions, with_offset_v<Offset>>, Reducers &...,
@@ -79,7 +81,9 @@ void execute_parallel_for(const sycl::range<Dimensions> &range, const Offset &of
             std::is_invocable_v<const KernelFunc, sycl::item<Dimensions, with_offset_v<Offset>>, Reducers &...>);
         kernel = [&](const sycl::item<Dimensions> &item) { func(item, reducers...); };
     }
-    sequential_for(range, offset, kernel);
+    printf("execute_parallel_for 84\n");
+    sequential_for(range, offset, kernel);    
+    printf("execute_parallel_for 86\n");
 }
 
 template<typename KernelName, int Dimensions, typename KernelFunc, typename... Reducers>
@@ -137,6 +141,7 @@ void parallel_for(sycl::range<Dimensions> num_work_items, sycl::kernel_handler k
 template<typename KernelName, typename KernelFunc, int Dimensions>
 void parallel_for(sycl::range<Dimensions> num_work_items, sycl::id<Dimensions> work_item_offset,
     sycl::kernel_handler kh, const KernelFunc &kernel_func) {
+    printf("parallel_for 140\n");
     execute_parallel_for<KernelName>(num_work_items, work_item_offset, kh, kernel_func);
 }
 

diff --git a/include/simsycl/detail/preprocessor.hh b/include/simsycl/detail/preprocessor.hh
@@ -21,3 +21,7 @@
 #define SIMSYCL_DETAIL_DEPRECATED_IN_SYCL
 #define SIMSYCL_DETAIL_DEPRECATED_IN_SYCL_V(message)
 #endif
+
+#define SIMSYCL_STRINGIZE_DETAIL(x) #x
+#define SIMSYCL_STRINGIZE(x) SIMSYCL_STRINGIZE_DETAIL(x)
+#define SIMSYCL_LINE_STRING SIMSYCL_STRINGIZE(__LINE__)
diff --git a/include/simsycl/sycl/handler.hh b/include/simsycl/sycl/handler.hh
@@ -91,6 +91,7 @@ class handler {
     template<typename KernelName = simsycl::detail::unnamed_kernel, typename KernelType, int Dimensions>
     SIMSYCL_DETAIL_DEPRECATED_IN_SYCL void parallel_for(
         range<Dimensions> num_work_items, id<Dimensions> work_item_offset, KernelType &&kernel_func) {
+        printf("parallel_for 94\n");
         detail::parallel_for<KernelName>(num_work_items, work_item_offset, kernel_handler(this), kernel_func);
     }
 

diff --git a/src/simsycl/schedule.cc b/src/simsycl/schedule.cc
@@ -43,27 +43,39 @@ namespace simsycl::detail {
 template<int Dimensions, typename Offset>
 void sequential_for(const sycl::range<Dimensions> &range, const Offset &offset,
     const simple_kernel<Dimensions, with_offset_v<Offset>> &kernel) {
+    printf("sequential_for " SIMSYCL_LINE_STRING "\n");
     // limit the number of work items scheduled at a time to avoid allocating huge index buffers
     constexpr size_t max_schedule_chunk_size = 16 << 10;
+    printf("sequential_for " SIMSYCL_LINE_STRING "\n");
     const auto schedule_chunk_size = std::min(range.size(), max_schedule_chunk_size);
+    printf("sequential_for " SIMSYCL_LINE_STRING "\n");
     const auto &schedule = get_cooperative_schedule();
+    printf("sequential_for " SIMSYCL_LINE_STRING "\n");
     std::vector<size_t> order(schedule_chunk_size);
+    printf("sequential_for " SIMSYCL_LINE_STRING "\n");
     auto schedule_state = schedule.init(order);
-
+    printf("sequential_for " SIMSYCL_LINE_STRING "\n");
+
     for(size_t schedule_offset = 0; schedule_offset < range.size(); schedule_offset += max_schedule_chunk_size) {
+        printf("sequential_for " SIMSYCL_LINE_STRING "\n");
         for(size_t schedule_id = 0; schedule_id < schedule_chunk_size; ++schedule_id) {
+            printf("sequential_for " SIMSYCL_LINE_STRING "\n");
             const auto linear_id = schedule_offset + order[schedule_id];
             if(linear_id < range.size()) {
                 if constexpr(with_offset_v<Offset>) {
                     const auto id = offset + linear_index_to_id(range, linear_id);
-                    kernel(make_item(id, range, offset));
+                    printf("sequential_for 61\n");
+                    kernel(make_item(id, range, offset));                    
+                    printf("sequential_for 63\n");
                 } else {
                     const auto id = linear_index_to_id(range, linear_id);
                     kernel(make_item(id, range));
                 }
             }
         }
+        printf("sequential_for 68\n");
         schedule_state = schedule.update(schedule_state, order);
+        printf("sequential_for 70\n");
     }
 }
 

diff --git a/src/simsycl/system.cc b/src/simsycl/system.cc
@@ -291,12 +291,18 @@ std::shared_ptr<cooperative_schedule>
 
 void parse_environment() {
     if(g_environment_parsed) return;
+    printf("parse_environment " SIMSYCL_LINE_STRING "\n");
 
     auto prefix = env::prefix("SIMSYCL");
     const auto system = prefix.register_variable<system_config>(
-        "SYSTEM", [](const std::string_view repr) { return read_system_config(std::string(repr)); });
+        "SYSTEM", [](const std::string_view repr) { 
+        printf("parse_environment " SIMSYCL_LINE_STRING "\n");
+        return read_system_config(std::string(repr)); 
+    });
+    printf("parse_environment " SIMSYCL_LINE_STRING "\n");
     const auto schedule = prefix.register_variable<std::shared_ptr<cooperative_schedule>>(
         "SCHEDULE", [](const std::string_view repr) -> std::shared_ptr<cooperative_schedule> {
+            printf("parse_environment " SIMSYCL_LINE_STRING "\n");
             if(repr == "rr") return std::make_unique<round_robin_schedule>();
             if(repr == "shuffle") return std::make_unique<shuffle_schedule>();
             if(repr.starts_with("shuffle:")) {
@@ -306,11 +312,14 @@ void parse_environment() {
             throw env::parser_error{
                 fmt::format("Invalid schedule '{}', permitted values are 'rr', 'shuffle', and 'shuffle:<seed>'", repr)};
         });
+    printf("parse_environment " SIMSYCL_LINE_STRING "\n");
 
     if(const auto parsed = prefix.parse_and_validate(); parsed.ok()) {
+        printf("parse_environment " SIMSYCL_LINE_STRING "\n");
         g_env_system_config = parsed.get(system);
         g_env_cooperative_schedule = parsed.get_or(schedule, nullptr);
     } else {
+        printf("parse_environment " SIMSYCL_LINE_STRING "\n");
         std::cerr << parsed.warning_message() << parsed.error_message();
     }
     g_environment_parsed = true;
@@ -325,13 +334,17 @@ std::shared_ptr<cooperative_schedule> g_cooperative_schedule;
 namespace simsycl {
 
 const cooperative_schedule &get_cooperative_schedule() {
+    printf("get_cooperative_schedule " SIMSYCL_LINE_STRING "\n");
     detail::parse_environment();
+    printf("get_cooperative_schedule " SIMSYCL_LINE_STRING "\n");
     if(detail::g_cooperative_schedule == nullptr) {
+        printf("get_cooperative_schedule " SIMSYCL_LINE_STRING "\n");
         if(detail::g_env_cooperative_schedule != nullptr) {
             detail::g_cooperative_schedule = detail::g_env_cooperative_schedule;
         } else {
             detail::g_cooperative_schedule = std::make_shared<round_robin_schedule>();
         }
+        printf("get_cooperative_schedule " SIMSYCL_LINE_STRING "\n");
     }
     return *detail::g_cooperative_schedule;
 }

diff --git a/src/simsycl/system_config.cc b/src/simsycl/system_config.cc
@@ -354,6 +354,7 @@ void from_json(const nlohmann::json &json, system_config &system) {
 
 system_config read_system_config(const std::string &path_to_json_file) {
     std::ifstream ifs(path_to_json_file);
+    printf("read_system_config " SIMSYCL_LINE_STRING "\n");
     return nlohmann::json::parse(ifs).get<system_config>();
 }