[RM] revise builtin device

include/simsycl/system.hh

@@ -87,13 +87,11 @@ struct device_config {
     std::string name{};
     std::string vendor{};
     std::string driver_version{};
-    std::string profile{};
     std::string version{};
     std::string backend_version{};
     std::vector<sycl::aspect> aspects{};
     std::vector<std::string> extensions{};
     size_t printf_buffer_size{};
-    bool preferred_interop_user_sync{};
     std::optional<simsycl::device_id> parent_device_id{};
     uint32_t partition_max_sub_devices{};
     std::vector<sycl::info::partition_property> partition_properties{};

src/simsycl/device.cc

@@ -2,6 +2,7 @@
 #include "simsycl/sycl/range.hh"
 #include "simsycl/system.hh"
 
+#include <cassert>
 #include <iterator>
 
 namespace simsycl::detail {
@@ -387,7 +388,7 @@ std::string device::get_info<info::device::driver_version>() const {
 
 template<>
 std::string device::get_info<info::device::profile>() const {
-    return state().config.profile;
+    throw exception(errc::invalid, "not an OpenCL backend");
 }
 
 template<>
@@ -419,13 +420,15 @@ size_t device::get_info<info::device::printf_buffer_size>() const {
 
 template<>
 bool device::get_info<info::device::preferred_interop_user_sync>() const {
-    return state().config.preferred_interop_user_sync;
+    throw exception(errc::invalid, "not an OpenCL backend");
 }
 
 template<>
 sycl::device device::get_info<info::device::parent_device>() const {
-    abort();
-    // return state().config.parent_device.value();
+    const auto parent_instance = state().parent.lock();
+    assert(parent_instance.has_value() == state().config.parent_device_id.has_value());
+    if(!parent_instance.has_value()) { throw exception(errc::invalid, "not a sub-device"); }
+    return *parent_instance;
 }
 
 template<>

src/simsycl/system.cc

@@ -3,7 +3,9 @@
 #include "simsycl/detail/check.hh"
 #include "simsycl/sycl/device.hh"
 #include "simsycl/sycl/platform.hh"
+#include "simsycl/sycl/vec.hh"
 
+#include <bit>
 #include <cassert>
 #include <iostream>
 #include <limits>
@@ -303,15 +305,16 @@ const platform_config builtin_platform{
     .extensions = {},
 };
 
-const device_config builtin_device{
-    .device_type = sycl::info::device_type::gpu,
-    .vendor_id = 0,
-    .max_compute_units = 16,
-    .max_work_item_dimensions = 3,
-    .max_work_item_sizes_1 = {1024},
-    .max_work_item_sizes_2 = {1024, 1024},
-    .max_work_item_sizes_3 = {64, 1024, 1024},
-    .max_work_group_size = 1024,
+// clang-format off
+const device_config builtin_device {
+    .device_type = sycl::info::device_type::gpu, //
+    .vendor_id = 0, //
+    .max_compute_units = 16, //
+    .max_work_item_dimensions = 3, //
+    .max_work_item_sizes_1 = {1024}, //
+    .max_work_item_sizes_2 = {1024, 1024}, //
+    .max_work_item_sizes_3 = {1024, 1024, 1024}, //
+    .max_work_group_size = 1024, //
     .max_num_sub_groups = 32,
     .sub_group_sizes = {32},
     .preferred_vector_width_char = 4,
@@ -342,56 +345,75 @@ const device_config builtin_device{
     .image_max_buffer_size = 0,
     .max_samplers = 0,
     .max_parameter_size = std::numeric_limits<std::size_t>::max(),
-    .mem_base_addr_align = 8,
+    .mem_base_addr_align = 8 * sizeof(sycl::long16),
     .half_fp_config
-    = {sycl::info::fp_config::denorm, sycl::info::fp_config::inf_nan, sycl::info::fp_config::round_to_nearest,
-        sycl::info::fp_config::round_to_zero, sycl::info::fp_config::round_to_inf, sycl::info::fp_config::fma,
-        sycl::info::fp_config::correctly_rounded_divide_sqrt},
-    .single_fp_config
-    = {sycl::info::fp_config::denorm, sycl::info::fp_config::inf_nan, sycl::info::fp_config::round_to_nearest,
-        sycl::info::fp_config::round_to_zero, sycl::info::fp_config::round_to_inf, sycl::info::fp_config::fma,
-        sycl::info::fp_config::correctly_rounded_divide_sqrt},
-    .double_fp_config
-    = {sycl::info::fp_config::denorm, sycl::info::fp_config::inf_nan, sycl::info::fp_config::round_to_nearest,
-        sycl::info::fp_config::round_to_zero, sycl::info::fp_config::round_to_inf, sycl::info::fp_config::fma,
-        sycl::info::fp_config::correctly_rounded_divide_sqrt},
+#if SIMSYCL_FEATURE_HALF_TYPE
+        = {sycl::info::fp_config::denorm, sycl::info::fp_config::inf_nan,
+            sycl::info::fp_config::fma, sycl::info::fp_config::correctly_rounded_divide_sqrt},
+#else
+        = {},
+#endif
+    .single_fp_config = {sycl::info::fp_config::denorm, sycl::info::fp_config::inf_nan,
+        sycl::info::fp_config::fma, sycl::info::fp_config::correctly_rounded_divide_sqrt},
+    .double_fp_config = {sycl::info::fp_config::denorm, sycl::info::fp_config::inf_nan,
+        sycl::info::fp_config::fma, sycl::info::fp_config::correctly_rounded_divide_sqrt},
     .global_mem_cache_type = sycl::info::global_mem_cache_type::read_write,
-    .global_mem_cache_line_size = 128,
+    .global_mem_cache_line_size = 64,
     .global_mem_cache_size = 16 << 20,
     .global_mem_size = std::numeric_limits<std::size_t>::max(),
     .max_constant_buffer_size = 1 << 16,
     .max_constant_args = std::numeric_limits<uint32_t>::max(),
     .local_mem_type = sycl::info::local_mem_type::local,
     .local_mem_size = 64 << 10,
     .error_correction_support = false,
-    .host_unified_memory = false,
+    .host_unified_memory = true,
+    .atomic_memory_order_capabilities = {sycl::memory_order::relaxed, sycl::memory_order::acquire,
+        sycl::memory_order::release, sycl::memory_order::acq_rel, sycl::memory_order::seq_cst},
+    .atomic_fence_order_capabilities = {sycl::memory_order::relaxed, sycl::memory_order::acquire,
+        sycl::memory_order::release, sycl::memory_order::acq_rel, sycl::memory_order::seq_cst},
+    .atomic_memory_scope_capabilities = {sycl::memory_scope::work_item,
+        sycl::memory_scope::sub_group, sycl::memory_scope::work_group, sycl::memory_scope::device,
+        sycl::memory_scope::system },
+    .atomic_fence_scope_capabilities = {sycl::memory_scope::work_item,
+        sycl::memory_scope::sub_group, sycl::memory_scope::work_group, sycl::memory_scope::device,
+        sycl::memory_scope::system },
     .profiling_timer_resolution = 1,
-    .is_endian_little = true,
+    .is_endian_little = std::endian::native == std::endian::little,
     .is_available = true,
-    .is_compiler_available = true,
-    .is_linker_available = true,
+    .is_compiler_available = false,
+    .is_linker_available = false,
     .execution_capabilities = {sycl::info::execution_capability::exec_kernel},
     .queue_profiling = true,
     .built_in_kernels = {},
     .platform_id = "SimSYCL",
     .name = "SimSYCL virtual GPU",
     .vendor = "SimSYCL",
     .driver_version = "0.1",
-    .profile = "FULL_PROFILE",
     .version = "0.1",
-    .aspects
-    = { sycl::aspect::gpu, sycl::aspect::accelerator, sycl::aspect::fp64, sycl::aspect::atomic64,
-        sycl::aspect::queue_profiling, sycl::aspect::usm_device_allocations, sycl::aspect::usm_host_allocations,
-        sycl::aspect::usm_shared_allocations, },
-    .extensions = {},
+    .aspects = { sycl::aspect::gpu, sycl::aspect::accelerator, sycl::aspect::emulated,
+        sycl::aspect::host_debuggable,
+#if SIMSYCL_FEATURE_HALF_TYPE
+        sycl::aspect::fp16,
+#endif
+        sycl::aspect::fp64, sycl::aspect::atomic64, sycl::aspect::queue_profiling,
+        sycl::aspect::usm_device_allocations, sycl::aspect::usm_host_allocations,
+        sycl::aspect::usm_atomic_host_allocations, sycl::aspect::usm_shared_allocations,
+        sycl::aspect::usm_atomic_shared_allocations, sycl::aspect::usm_system_allocations },
+    .extensions = {
+        "cl_khr_int64_base_atomics",
+        "cl_khr_int64_extended_atomics",
+#if SIMSYCL_FEATURE_HALF_TYPE
+        "cl_khr_fp16",
+#endif
+    },
     .printf_buffer_size = std::numeric_limits<std::size_t>::max(),
-    .preferred_interop_user_sync = true,
     .partition_max_sub_devices = 0,
     .partition_properties = {},
     .partition_affinity_domains = {sycl::info::partition_affinity_domain::not_applicable},
     .partition_type_property = sycl::info::partition_property::no_partition,
     .partition_type_affinity_domain = sycl::info::partition_affinity_domain::not_applicable,
 };
+// clang-format off
 
 const system_config builtin_system{
     .platforms = {{"SimSYCL", builtin_platform}},

src/simsycl/system_config.cc

@@ -247,13 +247,11 @@ void to_json(nlohmann::json &json, const device_config &device) {
         {"name", device.name},
         {"vendor", device.vendor},
         {"driver_version", device.driver_version},
-        {"profile", device.profile},
         {"version", device.version},
         {"backend_version", device.backend_version},
         {"aspects", device.aspects},
         {"extensions", device.extensions},
         {"printf_buffer_size", device.printf_buffer_size},
-        {"preferred_interop_user_sync", device.preferred_interop_user_sync},
         {"parent_device_id", device.parent_device_id},
         {"partition_max_sub_devices", device.partition_max_sub_devices},
         {"partition_properties", device.partition_properties},
@@ -333,13 +331,11 @@ void from_json(const nlohmann::json &json, device_config &device) {
     json.at("name").get_to(device.name);
     json.at("vendor").get_to(device.vendor);
     json.at("driver_version").get_to(device.driver_version);
-    json.at("profile").get_to(device.profile);
     json.at("version").get_to(device.version);
     json.at("backend_version").get_to(device.backend_version);
     json.at("aspects").get_to(device.aspects);
     json.at("extensions").get_to(device.extensions);
     json.at("printf_buffer_size").get_to(device.printf_buffer_size);
-    json.at("preferred_interop_user_sync").get_to(device.preferred_interop_user_sync);
     json.at("parent_device_id").get_to(device.parent_device_id);
     json.at("partition_max_sub_devices").get_to(device.partition_max_sub_devices);
     json.at("partition_properties").get_to(device.partition_properties);