Add source code and build / env documentation

.gitignore

@@ -1,4 +1,5 @@
 .cache/
 .vscode/
 /build*
-/compile_commands.json
+/compile_commands.json
+/hdoc-output

README.md

@@ -1,17 +1,16 @@
 <p align="center">
-<img src="resources/logo.png" alt="SimSYCL">
+<img src="resources/logo.png" width="600px" alt="SimSYCL">
+<br/>
 <i>“Technically correct is the best kind of correct”</i><br/>&nbsp;
 </p>
 
-# What and why is this?
-
 SimSYCL is a single-threaded, synchronous, library-only implementation of the SYCL 2020 specification. It enables you to test your SYCL applications against simulated hardware of different characteristics and discover bugs with its extensive verification capabilities.
 
 SimSYCL is in a very early stage of development - try it at your own risk!
 
 ## Requirements
 
-SimSYCL requires the Boost `context` libary.
+SimSYCL requires CMake, a C++20 compiler and the Boost `context` libary to be installed.
 
 ## Supported Platforms
 
@@ -25,6 +24,55 @@ The following platform and compiler combinations are currently tested in CI:
 Other platforms and compilers should also work, as long as they have sufficient C++20 support.
 Note that Clang versions prior to 17 do not currently work due to their incomplete CTAD support.
 
+## Installing SimSYCL
+
+To build SimSYCL and install it next to the source directory, run e.g.:
+
+```sh
+cmake -B build -DCMAKE_BUILD_TYPE=Debug -DCMAKE_INSTALL_PREFIX=$(pwd)-install
+cmake --build build --target install
+```
+
+### Build Options
+
+These can be set on the CMake command line via `-DOPTION=VALUE`.
+
+| option | values | effect |
+|---|---|---|
+| `SIMSYCL_ANNOTATE_SYCL_DEPRECATIONS` | `OFF`,`ON` | Mark deprecated SYCL APIs with `[[deprecated]]` (default `ON`) |
+| `SIMSYCL_ENABLE_ASAN`| `OFF`,`ON` | Build SimSYCL and the user code with AddressSanitizer (default `OFF`) |
+| `SIMSYCL_CHECK_MODE` | `SIMSYCL_CHECK_{NONE,LOG,THROW,ABORT}` | How to report verification errors (default `ABORT`) |
+
+## Using SimSYCL
+
+To get started, simply copy the `examples` folder to a separate location and edit its files as you see fit. To build against SimSYCL installed as above:
+
+```sh
+cmake -B build -DCMAKE_BUILD_TYPE=Debug -DCMAKE_PREFIX_PATH=/path/where/you/installed/SimSYCL
+cmake --build build
+```
+
+### Environment Variables
+
+These are available to all SimSYCL applications.
+
+| variable | values | effect |
+|---|---|---|
+| `SIMSYCL_SYSTEM` | `system.json` | Simulate the system defined in `system.json` |
+| `SIMSYCL_SCHEDULE` | `rr`, `shuffle`, `shuffle:<seed>` | Choose a schedule for work item order in kernels |
+
+### System Definition Files
+
+Systems are defined by listing all devices, platforms and their runtime properties in a JSON file.
+As a starting point, you can export the built-in system definition via
+```c++
+simsycl::write_system_config("system.json", simsycl::builtin_system);
+```
+and then use your (modified) system definition via
+```sh
+SIMSYCL_SYSTEM=system.json build/matmul
+```
+
 ## Research
 
 For a detailed introduction and evaluation of SimSYCL, please refer to the [the IWOCL'24 paper](https://dl.acm.org/doi/pdf/10.1145/3648115.3648136).

include/simsycl/detail/lock.hh

@@ -7,6 +7,9 @@
 
 namespace simsycl::detail {
 
+/// SimSYCL can be used from multiple threads simultaneously, but will serialize all actual SYCL work as well as access
+/// to structures with reference semantics using the (singleton) system lock. It is safe to lock recursively from within
+/// a single thread.
 class system_lock {
   public:
     system_lock();
@@ -18,6 +21,8 @@ class system_lock {
     std::lock_guard<std::recursive_mutex> m_lock;
 };
 
+/// Mutable state that is potentially shared between threads should be wrapped in a `shared_value` to ensure it can only
+/// be accessed when a `system_lock` is in scope.
 template<typename T>
 class shared_value {
   public:

include/simsycl/schedule.hh

@@ -7,6 +7,11 @@
 
 namespace simsycl {
 
+/// A schedule generates execution orders for work items within the constraints of group synchronization.
+///
+/// The kernel function is invoked once for each work item as prescribed by the schedule. For ND-range kernels, they can
+/// be suspended and resumed multiple times around group collective functions, in which case the schedule can instruct a
+/// different order for each iteration.
 class cooperative_schedule {
   public:
     using state = uint64_t;
@@ -18,16 +23,29 @@ class cooperative_schedule {
     cooperative_schedule &operator=(cooperative_schedule &&) = delete;
     virtual ~cooperative_schedule() = default;
 
+    /// Fill a vector of linear work item indices for the initial round of kernel invocations.
+    ///
+    /// The retured `state` is to be carried into the first invocation of `update()`, if any.
     [[nodiscard]] virtual state init(std::vector<size_t> &order) const = 0;
+
+    /// After work items have been suspended on an ND-range kernel collective function, fill the index vector with the
+    /// next sequence of work item indices.
+    ///
+    /// The retured `state` is to be carried into the next invocation of `update()`.
     [[nodiscard]] virtual state update(state state_before, std::vector<size_t> &order) const = 0;
 };
 
+/// A schedule executing threads in-order by linear thread id.
 class round_robin_schedule final : public cooperative_schedule {
   public:
     [[nodiscard]] state init(std::vector<size_t> &order) const override;
     [[nodiscard]] state update(state state_before, std::vector<size_t> &order) const override;
 };
 
+/// A schedule executing threads in randomly shuffled order. After each collective barrier, indices are re-shuffled.
+///
+/// SYCL programs need to ensure their kernels are correct under any order of thread items, so using a
+/// `shuffle_schedule` (potentially with multiple seeds) allows fuzzing that assumption.
 class shuffle_schedule final : public cooperative_schedule {
   public:
     shuffle_schedule() = default;
@@ -40,8 +58,12 @@ class shuffle_schedule final : public cooperative_schedule {
     uint64_t m_seed = 1234567890;
 };
 
-// effect is thread-local
+/// Return the thread-locally active schedule.
 const cooperative_schedule &get_cooperative_schedule();
+
+/// Set the thread-locally active schedule for future kernel invocations.
+///
+/// Must not be called from within a kernel.
 void set_cooperative_schedule(std::shared_ptr<const cooperative_schedule> schedule);
 
 } // namespace simsycl

include/simsycl/system.hh

@@ -12,10 +12,15 @@ namespace simsycl {
 
 class cooperative_schedule;
 
+/// Identifier for `sycl::platform`s within a `system_config`.
 using platform_id = std::string;
+
+/// Identifier for `sycl::device`s within a `system_config`.
 using device_id = std::string;
-using system_id = std::string;
 
+/// Configuration for a single `sycl::device`.
+///
+/// All data members will be reflected in the `sycl::info::device` properties of the created device.
 struct device_config {
     sycl::info::device_type device_type{};
     uint32_t vendor_id{};
@@ -97,6 +102,9 @@ struct device_config {
     sycl::info::partition_affinity_domain partition_type_affinity_domain{};
 };
 
+/// Configuration for a single `sycl::platform`.
+///
+/// All data members will be reflected in the `sycl::info::platform` properties of the created platform.
 struct platform_config {
     std::string profile{};
     std::string version{};
@@ -105,20 +113,44 @@ struct platform_config {
     std::vector<std::string> extensions{};
 };
 
+/// Configuration for the entire system simulated by SimSYCL.
 struct system_config {
+    /// All platforms returned by `sycl::platform::get_platforms()`, in order.
+    ///
+    /// The `platform_id` is only relevant as a reference within this struct.
     std::unordered_map<platform_id, platform_config> platforms{};
+
+    /// All devices returned by `sycl::device::get_devices()`, in order.
+    ///
+    /// The `device_id` is only relevant as a reference within this struct.
     std::unordered_map<device_id, device_config> devices{};
 };
 
+/// Configuration of the builtin platform as returned through `sycl::platform::get_platforms()` by default.
 extern const platform_config builtin_platform;
+
+/// Configuration of the builtin device as returned through `sycl::device::get_devices()` by default.
 extern const device_config builtin_device;
+
+/// Default system configuration when not overriden by the environment.
 extern const system_config builtin_system;
 
+/// Return the system configuration specified by the environment via `SIMSYCL_SYSTEM=system.json`, or `builtin_system`
+/// as a fallback.
 const system_config &get_default_system_config();
+
+/// Read a `system_config` from a JSON file.
 system_config read_system_config(const std::string &path_to_json_file);
+
+/// Write a `system_config` to a JSON file.
 void write_system_config(const std::string &path_to_json_file, const system_config &config);
+
+/// Replace the active `system_config` so that all future calls to device and platform selection functions return
+/// members of this configuration.
 void configure_system(const system_config &system);
 
+/// Return the schedule for threads of a kernel specified by the environment via `SIMSYCL_SCHEDULE`, or a
+/// `round_robin_schedule` as a fallback.
 std::shared_ptr<const cooperative_schedule> get_default_cooperative_schedule();
 
 } // namespace simsycl