Initial version of the doc

Author: AlexeySachkov

sycl/doc/design/MappingHostAddressesToDeviceEntities.md

@@ -0,0 +1,240 @@
+# Mapping host variables to compiler-generated info
+
+[SYCL 2020][sycl-2020-spec] specification and some extensions such as
+[SYCL_INTEL_device_global][device-global-ext-spec] imply that implementation is
+capable to somehow map addresses of a host objects to their counterparts in
+device programs.
+
+For example, in order to implement specialization constants on top of SPIR-V, we
+need to be able to map addresses of `specialization_id` variables into numeric
+IDs of corresponding specialization constants at SPIR-V level.
+
+Another example is device global [implementation][device-global-design], where
+in order to communicate a value of `device_global` variable between host and
+device we need to map its host address to a symbolic name/identifier and some
+other info, which is used at PI layer and below.
+
+This design document describes a generic way how to map address of any SYCL
+object defined in a namespace scope to its unique symbolic ID. Please note that
+this document doesn't try to map the address to something other than a unique
+symbolic ID: other required information is usually generated by the device
+compiler and communicated to the runtime by device image properties. Unique
+symbolic ID which can be obtained from mapping mechanism described in this
+design document could be used as a key in those properties to propagate
+additional information using existing mechanisms.
+
+So, overall the picture looks like:
+- device compiler generates property set/s which provide mapping
+  "unique symbolic ID" -> "various information required by DPC++ RT"
+- device or host compiler generates mapping
+  "address of a host variable" -> "unique symbolic ID" (as described below by
+  this document)
+- DPC++ RT uses these to mappings to obtain required information and somehow
+  uses it
+
+This design document describes two approaches of how the 
+"address of a host variable" -> "unique symbolic ID" mapping can be generated:
+the first one with integration footer and another one with modification of the
+host compiler.
+
+Both approaches have their pros and cons and they are expected to be implemented
+and exists in the implementation at the same time, but only one of them will be
+used at a time depending on whether 3rd-party host compiler is used or not.
+
+Integration footer can be used with 3rd-party host compilers, but it requires
+appending to a translation unit provided by user, which could affect debug
+information: since there are no compilers that support appending a file at the
+end (similar to `-include`), appending is done by generating a temporary input
+file using concatenation of the original input and integration footer.
+
+Such replacement of the main translation unit causes the following issues:
+- debug information about the source file might be incorrect, leading to
+  problems with gdb `l` command and code coverage tools
+- checksum of host and device source files becomes different which causes device
+  code debugging to be completely broken in some environments (such as MS Visual
+  Studio, for example)
+
+Customizing host compiler allows to avoid issues with debuggers and code
+coverage tools, but that is not an option if user wants to compile host part
+of an app with a 3rd-party host compiler.
+
+Further sections describe the implementation design of both approaches in more
+details, note that there are few components which should be modified regardless
+of which approach is in use.
+
+## Common front-end part
+
+DPC++ FE should support the following attribute:
+`[[__sycl_detail__::uniquely_identifiable_object(kind)]]`. This attribute accepts
+a string literal and should be applied to types (like `device_global` or
+`specialization_id`).
+
+Presence of the attribute instructs the compiler to perform the following
+things:
+- emit `sycl-unique-id` LLVM IR attribute on each definition of a variable of
+  type marked with `[[__sycl_detail__::uniquely_identifiable_object(kind)]]`
+  attribute. `sycl-unique-id` LLVM IR attribute should be accompanied by a
+  unique string identifier of a variable it is attached to. The rules for
+  creating this string are the same as for `__builtin_sycl_unique_stable_id` and
+  the same algorithm can be used when generating the string for the attribute
+- emit `sycl-uid-type` LLVM IR attribute alongside `sycl-unique-id`, which
+  contains the `kind` string passed to
+  `[[__sycl_detail__::uniquely_identifiable_object(kind)]]` attribute
+
+**TODO**: we have `[[__sycl_detail__::device_global]]` attribute documented in
+[device global design doc][device-global-design], which instructs front-end to
+emit some additional semantic checking. Shall we leave it in place or that
+request for semantic checking should also be documented by
+`[[__sycl_detail__::uniquely_identifiable_object(kind)]]` attribute when `kind`
+is set to a certain value?
+
+**TODO**: alternatively, we could completely re-use existing
+`[[__sycl_detail__::device_global]]` attribute and introduce another one for
+specialization constants, i.e. it is a question of whether or not we want to
+generalize unique IDs generation in form of a generic attribute or not.
+
+When DPC++ compiler is used as both host and device compiler, then the attribute
+should be respected by both host and device compiler passes and LLVM IR
+attributes should appear in LLVM IR for both host and device code. When DPC++
+compiler is only used as a device compiler, then we don't expect the attribute
+to be handled on host, apparently.
+
+Another thing we need from DPC++ FE compiler is to define a special macro, which
+will allow to distinguish it from other compilers. That is needed to apply the
+aforementioned attribute conditionally to avoid spamming users with warnings
+about unknown attributes.
+
+The suggested macro name is `__INTEL_SYCL_HOST_COMPILER__`. It should be defined
+when the compiler is invoked in SYCL host mode (`-fsycl-is-host` `-cc1` flag).
+
+## Common headers part
+
+Header files should be modified by adding the new attributes to types
+declarations, objects of which we will need in our mapping.Again,
+`device_global` and `specialization_id` are examples here:
+
+```
+template <typename T>
+class
+#if defined(__SYCL_DEVICE_ONLY__) || defined(__INTEL_SYCL_HOST_COMPILER__)
+  [[__sycl_detail__::uniquely_identifiable_object("specialization_id")]]
+#endif
+specialization_id {
+// ...
+};
+```
+
+## Common runtime part
+
+The runtime should implement the following function, which will be called from
+a code generated by the compiler (see the next section):
+
+```
+void __register_uniquely_identifiable_object(
+  void *Address, const char* UniqueID, const char *Kind);
+```
+
+The function accepts the following arguments:
+- `Address` is an address of a variable, which exists in an application on host
+- `UniqueID` is a unique symbolic ID, which corresponds to that variable
+- `Kind` is a string which corresponds to `kind` argument passed to
+  `[[__sycl_detail__::uniquely_identifiable_object(kind)]]` attribute attached
+  to the type of the variable identified by `Address`. It can be used to
+  distinguish different entities like `specialization_id` and `device_global`:
+  for example they could be stored in different maps to speed up certain
+  operations with them.
+
+The compiler guarantees that the function will be called zero or more times
+(depending on the amount of uniquely identifiable objects found in a program)
+_before_ application's `main()` function, i.e. in a global constructor.
+
+That poses some restrictions on those uniquely identifiable object, i.e. that
+they can't be used from another global object due to risk of accessing a
+non-initialized object, but that is an UB anyway because the order of global
+objects initialization is not defined in C++ when those objects are defined in
+different translation unit.
+
+## Compiler driver part
+
+The compiler driver is the component which is responsible for selecting the
+approach we are taking and the decision is made based on whether or not
+3rd-party host compiler is in use.
+
+If `-fsycl-host-compiler` option is present, the compiler driver chooses the
+integration footer approach:
+- it supplies device compilation step with `-fsycl-int-footer` option to
+  instruct device compiler to emit integration footer
+- it appends the integration footer to user-provided translation unit before
+  passing it to a host compiler
+
+Otherwise, if `-fsycl-host-compiler` is not present, then the compiler driver
+chooses another approach by simply doing nothing related to integration footer:
+- `-fsycl-int-footer` is **not** passed to device compiler
+- user-provided translation unit is passes as-is to host compiler
+
+## Integration footer approach
+
+When this approach is used, not only extra file (integration footer) is
+generated, but integration header is also modified: FE compiler generates a
+definition of a namespace scope variable of type
+`__sycl_device_global_registration` whose sole purpose it to run its constructor
+before the application's `main()` function:
+
+```
+namespace sycl::detail {
+namespace {
+
+class __sycl_device_global_registration {
+ public:
+  __sycl_device_global_registration() noexcept;
+};
+__sycl_device_global_registration __sycl_device_global_registrar;
+
+} // namespace (unnamed)
+} // namespace sycl::detail
+```
+
+The integration footer generated by the compiler contains the definition of the
+constructor, which calls a function in the DPC++ runtime, which registers
+needed mappings:
+
+```
+namespace sycl::detail {
+namespace {
+
+__sycl_device_global_registration::__sycl_device_global_registration() noexcept {
+  __register_uniquely_identifiable_object(
+    &::Foo,
+    /* same string returned from __builtin_sycl_unique_stable_id(::Foo) */,
+    "specialization_id");
+  __register_uniquely_identifiable_object(
+    &::inner::Bar,
+    /* same string returned from __builtin_sycl_unique_stable_id(::inner::Bar) */,
+    "device_global");
+}
+
+} // namespace (unnamed)
+} // namespace sycl::detail
+```
+
+## Custom host compiler approach
+
+With this approach, we simply schedule a one more pass in the optimization
+pipeline, which should be executed regardless of the optimization level, because
+it is required for proper functioning of some features.
+
+The pass does similar thing to integration footer: it emits a global constructor
+which in turn calls `__register_uniquely_identifiable_object` to provide the
+runtime with required mapping information.
+
+Unlike with integration footer approach no separate file is being generated,
+which preserves all source files mapping and checksums to be in place and
+correct.
+
+Generated constructor function should have internal linkage to avoid possible
+names clashes and multiple definition errors later at link stage.
+
+Generated constructor contains a call to
+`__register_uniquely_identifiable_object` for each global variable which has
+`sycl-unique-id` and `sycl-uid-kind` attributes, passing values of those
+attributes into the corresponding arguments of the function.