Partially address review feedback

Author: aelovikov-intel

sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler.asciidoc

@@ -1199,6 +1199,44 @@ this extension.
 
 == Non-normative implementation notes for {dpcpp}
 
+=== C/{cpp} header files and limitations when the language is sycl
+
+By default, the SYCL runtime compiler uses a self-contained set of C and {cpp}
+header files when compiling kernels in the `sycl` language. This means that SYCL
+applications using this feature can be run even on a system that does not have
+these headers installed. However, the self contained header files may not be the
+same as the C and {cpp} header files that were used to build the host part of
+the SYCL application. As a result, there are additional limitations around data
+that is shared between the host part of the application and the kernel. These
+limitations apply to arguments that are passed to the kernel and also to data
+shared through USM or through accessors. Additionally, other header files might
+be installed in the same location as system C headers (e.g., `/usr/include/`).
+Those will not be available as well.
+
+Types that are defined by the compiler (e.g. fundamental types like `int` and
+`float`) are guaranteed to have the same representation and alignment
+requirements in both the host compiler and in the compiler used to compile the
+kernel. Therefore data using these types can be safely shared. However, types
+defined by the C or {cpp} library (e.g. types in the `std` namespace) are not
+guaranteed to be the same, so data defined using these types cannot be safely
+shared. There are a few specific exceptions to this limitation. The following C
+/ {cpp} types are guaranteed to have the same representation and alignment
+requirements, so data defined as these types can be safely shared:
+
+* The following fixed-width integer types: `int8_t`, `int16_t`, `int32_t`,
+`int64_t`, `uint8_t`, `uint16_t`, `uint32_t`, `uint64_t`.
+* The types `size_t` and `ptrdiff_t`.
+
+Array and class types defined by your application are safe to share so long as
+the element types are safe. Pointer types are safe to share so long as the
+pointed-at type is safe. Enumeration types defined by your application are also
+safe.
+
+It is also possible to use the system C and C++ headers instead of the
+self-contained versions. See the description of the
+`--sycl-rtc-use-system-includes` option for more details.
+
+
 === Supported `build_options` when the language is `sycl`
 
 The SYCL runtime compiler supports the following {dpcpp} options to be passed in
@@ -1359,44 +1397,6 @@ default.
 
 === Known issues and limitations when the language is `sycl`
 
-==== C/{cpp} header files
-
-By default, the SYCL runtime compiler uses a self-contained set of C and {cpp}
-header files when compiling kernels in the `sycl` language (unless target has
-other dependencies in system includes). This means that SYCL applications using
-this feature can be run even on a system that does not have these headers
-installed. However, the self contained header files may not be the same as the C
-and {cpp} header files that were used to build the host part of the SYCL
-application. As a result, there are additional limitations around data that is
-shared between the host part of the application and the kernel. These
-limitations apply to arguments that are passed to the kernel and also to data
-shared through USM or through accessors. Additionally, other header files might
-be installed in the same location as system C headers (e.g., `/usr/include/`).
-Those will not be available as well.
-
-Types that are defined by the compiler (e.g. fundamental types like `int` and
-`float`) are guaranteed to have the same representation and alignment
-requirements in both the host compiler and in the compiler used to compile the
-kernel. Therefore data using these types can be safely shared. However, types
-defined by the C or {cpp} library (e.g. types in the `std` namespace) are not
-guaranteed to be the same, so data defined using these types cannot be safely
-shared. There are a few specific exceptions to this limitation. The following C
-/ {cpp} types are guaranteed to have the same representation and alignment
-requirements, so data defined as these types can be safely shared:
-
-* The following fixed-width integer types: `int8_t`, `int16_t`, `int32_t`,
-`int64_t`, `uint8_t`, `uint16_t`, `uint32_t`, `uint64_t`.
-* The types `size_t` and `ptrdiff_t`.
-
-Array and class types defined by your application are safe to share so long as
-the element types are safe. Pointer types are safe to share so long as the
-pointed-at type is safe. Enumeration types defined by your application are also
-safe.
-
-It is also possible to use the system C and C++ headers instead of the
-self-contained versions. See the description of the
-`--sycl-rtc-use-system-includes` option for more details.
-
 ==== Changing the compiler action or output
 
 As the {dpcpp} frontend is integrated tightly in the runtime compilation

sycl/test-e2e/KernelCompiler/sycl_host_device_data_layout.cpp

@@ -17,6 +17,22 @@
 #define STRINGIFY(x) #x
 #define EXPAND_AND_STRINGIFY(x) STRINGIFY(x)
 
+// Needs to be duplicated between host/device. @{
+
+// Comma would make preprocessor macro trickier.
+using mint3 = sycl::marray<int, 3>;
+
+enum E {
+  V0 = 0x12345689,
+};
+static_assert(sizeof(E) == 4);
+enum class ScopedE {
+  ScopedV0 = 0x12345689,
+};
+static_assert(sizeof(ScopedE) == 4);
+
+// }@
+
 namespace syclexp = sycl::ext::oneapi::experimental;
 int main() {
   sycl::queue q;
@@ -35,6 +51,15 @@ namespace syclexp = sycl::ext::oneapi::experimental;
 
 using mint3 = sycl::marray<int, 3>;
 
+enum E {
+  V0 = 0x12345689,
+};
+static_assert(sizeof(E) == 4);
+enum class ScopedE {
+  ScopedV0 = 0x12345689,
+};
+static_assert(sizeof(ScopedE) == 4);
+
 extern "C"
 SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::single_task_kernel))
 void foo(TYPE *in, TYPE *out, size_t *align_out, size_t *size_out, bool *equal_out) {
@@ -150,12 +175,35 @@ void foo(TYPE *in, TYPE *out, size_t *align_out, size_t *size_out, bool *equal_o
   TEST3(sycl::range<3>, 0x1122334455667788, 0x1223344556677889,
         0x132435465768798A)
 
+  TEST(sycl::id<1>, 0x1122334455667788)
+  TEST2(sycl::id<2>, 0x1122334455667788, 0x1223344556677889)
+  TEST3(sycl::id<3>, 0x1122334455667788, 0x1223344556677889,
+        0x132435465768798A)
+
+  // Making these work with macros would be too much work:
+  Test(sycl::nd_range<1>{{0x1122334455667788}, {0x1223344556677889}},
+       "sycl::nd_range<1>", "{0x1122334455667788}, {0x1223344556677889}");
+  Test(sycl::nd_range<2>{{0x1122334455667788, 0x2132435465768798},
+                         {0x1223344556677889, 0x2233445586778899}},
+       "sycl::nd_range<2>",
+       "{0x1122334455667788, 0x2132435465768798}, {0x1223344556677889, "
+       "0x2233445586778899}");
+  Test(
+      sycl::nd_range<3>{
+          {0x1122334455667788, 0x2132435465768798, 0x31525364758697A8},
+          {0x1223344556677889, 0x2233445586778899, 0x32435465768798A9}},
+      "sycl::nd_range<3>",
+      "{0x1122334455667788, 0x2132435465768798, 0x31525364758697A8}, "
+      "{0x1223344556677889, 0x2233445586778899, 0x32435465768798A9}");
+
   TEST2(sycl::short2, 0x1234, 0x2345)
   TEST3(sycl::short3, 0x1234, 0x2345, 0x3456)
 
-  using mint3 = sycl::marray<int, 3>;
   TEST3(mint3, 0x1234, 0x2345, 0x3456)
 
+  TEST(E, V0)
+  TEST(ScopedE, ScopedE::ScopedV0)
+
   sycl::free(align, q);
   sycl::free(size, q);
   sycl::free(equal, q);