Merge remote-tracking branch 'upstream/SYCL-2020' into steffen/restri…

…ct_bitwise_float_group_reduce

.github/workflows/cts_ci.yml

@@ -59,7 +59,7 @@ jobs:
           # - sycl-impl: computecpp
           # version: 2.11.0
           - sycl-impl: dpcpp
-            version: 0267c1b
+            version: 83f877975c97acdb38d84f94dc146571cd522e0e
           - sycl-impl: hipsycl
             version: 3d8b1cd
     steps:
@@ -119,7 +119,7 @@ jobs:
           - sycl-impl: computecpp
             version: 2.11.0
           - sycl-impl: dpcpp
-            version: 0267c1b
+            version: 83f877975c97acdb38d84f94dc146571cd522e0e
           - sycl-impl: hipsycl
             version: 3d8b1cd
     env:

docker/dpcpp/Dockerfile

@@ -1,24 +1,7 @@
-# DPC++ version (git revision) to install
+# DPC++ build version (git revision) to install
+# Go to https://github.com/intel/llvm/pkgs/container/llvm%2Fsycl_ubuntu2204_nightly to see avilable docker image tags
 ARG IMPL_VERSION
 
-FROM khronosgroup/sycl-cts-ci:common
-
-ARG IMPL_VERSION
-RUN test -n "$IMPL_VERSION" || ( echo "Error: IMPL_VERSION is not set"; exit 1 )
-
-RUN git clone https://github.com/intel/llvm.git \
-      --branch=sycl --single-branch --shallow-since=2021-09-01 \
-      --recurse-submodules /tmp/dpcpp && \
-    cd /tmp/dpcpp && \
-    git checkout $IMPL_VERSION && \
-    python3 /tmp/dpcpp/buildbot/configure.py \
-      --src-dir=/tmp/dpcpp \
-      --obj-dir=/tmp/build \
-      --build-type=Release \
-      --cmake-opt=-DCMAKE_INSTALL_PREFIX=/sycl && \
-    python3 /tmp/dpcpp/buildbot/compile.py \
-      --src-dir=/tmp/dpcpp \
-      --obj-dir=/tmp/build && \
-    rm -rf /tmp/dpcpp /tmp/build
+FROM ghcr.io/intel/llvm/sycl_ubuntu2204_nightly:no-drivers-$IMPL_VERSION
 
 COPY configure.sh /scripts/

docker/dpcpp/configure.sh

@@ -5,8 +5,8 @@ set -o errexit -o pipefail -o noclobber -o nounset
 cmake . -G Ninja -B build \
     -DCMAKE_CXX_COMPILER_LAUNCHER=ccache \
     -DSYCL_IMPLEMENTATION=DPCPP \
-    -DDPCPP_INSTALL_DIR=/sycl \
-    -DCMAKE_CXX_COMPILER=/sycl/bin/clang++ \
+    -DDPCPP_INSTALL_DIR=/opt/sycl \
+    -DCMAKE_CXX_COMPILER=/opt/sycl/bin/clang++ \
     -DCMAKE_BUILD_TYPE=Release \
     -DSYCL_CTS_ENABLE_FULL_CONFORMANCE=0 \
     -DSYCL_CTS_ENABLE_DEPRECATED_FEATURES_TESTS=1 \

test_plans/bfloat16.asciidoc

@@ -0,0 +1,193 @@
+:sectnums:
+:xrefstyle: short
+
+= Test plan for bfloat16
+
+This is a test plan for the APIs described in
+https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/supported/sycl_ext_oneapi_bfloat16.asciidoc[sycl_ext_oneapi_bfloat16]
+
+== Testing scope
+
+=== Device coverage
+
+All of the tests described below are performed only on the default device that
+is selected on the CTS command line.
+
+=== Type coverage
+
+All of the tests described below are performed using each of the following `typename T`:
+
+* `char`
+* `short`
+* `int`
+* `long long`
+* `std::size_t`
+* `bool`
+* `float`
+
+In addition, if the device has aspect `aspect::fp64`:
+
+* `double`
+
+In addition, if the device has `aspect::fp16`:
+
+* `sycl::half`
+
+Note that `T` must be convertible to `float`
+
+== Tests
+
+=== Size
+
+`bfloat16` must occupy 2 bytes of memory, check that `sizeof(bfloat16) == 2`.
+This check is required to calculate special values.
+
+=== Constructors
+
+`bfloat16()` +
+`bfloat16(const bfloat16 &)` +
+`~bfloat16()`
+
+Check that:
+
+* `std::is_default_constructible_v<bfloat16> == true`;
+* `std::is_copy_constructible_v<bfloat16> == true`;
+* `std::is_destructible_v<bfloat16> == true`.
+
+`bfloat16(const float &a)` +
+`bfloat16 &operator=(const float &a)`
+
+* Create a `float` variable `f` equal to 1;
+* Create a `bfloat16` variable `bf1` passing `f` to the constructor;
+* Create a `bfloat16` variable `bf2` by assinging it to `f`;
+* Verify that:
+    ** `bf1 == f`
+    ** `bf2 == f`
+
+`bfloat16(const sycl::half &a)` +
+`bfloat16 &operator=(const sycl::half &a)`
+
+Same as above, but with `sycl::half` instead of `float`
+
+=== Special values
+
+Use `uint*_t` variables representing `bfloat16` and `float` values in bitset format.
+
+==== Minimum positive normal value
+
+[source, c++]
+----
+    uint16_t bfloat16_bits = 0b0000000010000000;
+    uint32_t float_bits    = 0b00000000100000000000000000000000;
+
+    bfloat16 bf_min = sycl::bit_cast<bfloat16>(bfloat16_bits);
+    float float_min = sycl::bit_cast<float>(float_bits);
+----
+
+Verify that the minimum values of `bfloat16` type is equal to the minimum value of `float` by using the `bfloat16` operator `==`.
+
+==== Zero
+
+[source, c++]
+----
+    uint16_t bfloat16_bits = 0b0000000000000000;
+    uint32_t float_bits = 0b00000000000000000000000000000000;
+
+    bfloat16 bf_zero = sycl::bit_cast<bfloat16>(bfloat16_bits);
+    float float_zero = sycl::bit_cast<float>(float_bits);
+----
+    Verify that `bf_zero == float_zero`.
+
+==== NaN
+
+[source, c++]
+----
+    uint16_t bfloat16_bits[4] = {
+        0b0111111111000001, // qNaN
+        0b1111111111000001,
+        0b0111111110000001, // sNaN
+        0b1111111110000001};
+
+    bfloat16 bf_nan[4];
+    for (int i = 0; i < 4; i++) {
+        bf_nan[i] = sycl::bit_cast<bfloat16>(bfloat16_bits[i]);
+    }
+----
+    Verify that `std::isnan(bf_nan)` is `true` for all elements.
+
+==== Infinity
+
+[source, c++]
+----
+    uint16_t bfloat16_bits[2] = {
+        0b0111111110000000,
+        0b1111111110000000};
+
+    bfloat16 bf_inf_0 = sycl::bit_cast<bfloat16>(bfloat16_bits[0]);
+    bfloat16 bf_inf_1 = sycl::bit_cast<bfloat16>(bfloat16_bits[1]);
+----
+     Verify that `std::isinf(bf_inf_0)` and `std::isinf(bf_inf_1)` are `true`.
+
+These tests will fail if the implementation does not use an 8 bit exponent and 7 bit significand for `bfloat16`.
+
+=== Conversion
+
+Check that:
+
+* `std::is_convertible_v<bfloat16, float> == true`
+* `std::is_convertible_v<bfloat16, sycl::half> == true`
+* `std::is_convertible_v<bfloat16, bool> == true`
+* `std::is_convertible_v<float, bfloat16> == true`
+* `std::is_convertible_v<sycl::half, bfloat16> == true`
+
+=== Operators
+
+`operator-(bfloat16 &bf)`
+
+Check that it constructs new instance of `bfloat16` class with negated value.
+Create `neg_bf` using this operator and verify:
+
+* `neg_bf == -bf`
+* `bf == -neg_bf`
+
+(Prefix) +
+`bfloat16 &operator++(bfloat16 &bf)` +
+`bfloat16 &operator--(bfloat16 &bf)`
+
+* Check if it adds/substracts 1 to the value of the object referenced by this `bf`.
+* Check that new value of the referenced object is equal to `(previous value +/- 1)`.
+* Check if it returns the copy of `bf`.
+* Check returned value type.
+
+(Postfix) +
+`bfloat16 operator++(bfloat16 &bf, int)` +
+`bfloat16 operator--(bfloat16 &bf, int)`
+
+Same as above, but check thar it returns value of `bf` before assignment instead of copy.
+
+OP is `+=`, `-=`, `*=`, `/=` +
+`bfloat16 &operatorOP(bfloat16 &lhs, const bfloat16 &rhs)`
+
+* Check results of arithmetic operations returned to initial `bfloat16` object.
+* Check returned value type.
+
+OP is `+`, `-`, `*`, `/` +
+`bfloat16 operatorOP(const bfloat16 &lhs, const bfloat16 &rhs)`
+
+* Check results of arithmetic operations.
+* Check returned value type.
+
+OP is `==`, `!=`, `<`, `>`, `+<=+`, `>=` +
+`bool operatorOP(const bfloat16 &lhs, const bfloat16 &rhs)`
+
+* Check results of equality and inequality between two `bfloat16` objects.
+* Check returned type is `bool`.
+
+OP is `==`, `!=`, `<`, `>`, `+<=+`, `>=` +
+`template <typename T>` +
+`bool operatorOP(const bfloat16 &lhs, const T &rhs)` +
+`template <typename T>` +
+`bool operatorOP(const T &lhs, const bfloat16 &rhs)`
+
+* Check results of equality and inequality between `bfloat16` and `T` objects.
+* Check returned type is `bool`.

test_plans/intel_device_info.asciidoc

@@ -0,0 +1,130 @@
+:sectnums:
+:xrefstyle: short
+
+= Test plan for Intel's Extensions for Device Information
+
+This is a test plan for the APIs described in
+https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/supported/sycl_ext_intel_device_info.md[sycl_ext_intel_device_info]
+
+== Testing scope
+
+=== Device coverage
+
+All of the tests described below are performed only on the default device that
+is selected on the CTS command line.
+
+=== Feature test macro
+
+All of the tests should use `#ifdef SYCL_EXT_INTEL_DEVICE_INFO` so they can be skipped
+if feature is not supported.
+
+== Tests
+
+Run the code in the sections below to check if device has aspect for this `info` and check type of `get_info()` with this `info`.
+
+=== Device ID
+[source, c++]
+----
+if (dev.has(aspect::ext_intel_device_id)) {
+    auto ID = dev.get_info<ext::intel::info::device::device_id>();
+    if (!std::is_same_v<decltype(ID), uint32_t>) { /* test failed */ }
+}
+----
+=== Device UUID
+[source, c++]
+----
+if (dev.has(aspect::ext_intel_device_info_uuid)) {
+    auto UUID = dev.get_info<ext::intel::info::device::uuid>();
+    if (!std::is_same_v<decltype(UUID), std::array<unsigned char, 16>>) { /* test failed */ }
+}
+----
+=== PCI Address
+[source, c++]
+----
+if (dev.has(aspect::ext_intel_pci_address)) {
+    auto BDF = dev.get_info<ext::intel::info::device::pci_address>();
+    if (!std::is_same_v<decltype(BDF), std::string>) { /* test failed */ }
+}
+----
+=== Intel GPU Execution Unit SIMD Width
+[source, c++]
+----
+if (dev.has(aspect::ext_intel_gpu_eu_simd_width)) {
+    auto euSimdWidth = dev.get_info<ext::intel::info::device::gpu_eu_simd_width>();
+    if (!std::is_same_v<decltype(euSimdWidth), uint32_t>) { /* test failed */ }
+}
+----
+=== Intel GPU Execution Unit Count
+[source, c++]
+----
+if (dev.has(aspect::ext_intel_gpu_eu_count)) {
+    auto euCount = dev.get_info<ext::intel::info::device::gpu_eu_count>();
+    if (!std::is_same_v<decltype(euCount), uint32_t>) { /* test failed */ }
+}
+----
+=== Intel GPU Number of Slices
+[source, c++]
+----
+if (dev.has(aspect::ext_intel_gpu_slices)) {
+    auto slices = dev.get_info<ext::intel::info::device::gpu_slices>();
+    if (!std::is_same_v<decltype(slices), uint32_t>) { /* test failed */ }
+}
+----
+=== Intel GPU Number of Subslices per Slice
+[source, c++]
+----
+if (dev.has(aspect::ext_intel_gpu_subslices_per_slice)) {
+    auto subslices = dev.get_info<ext::intel::info::device::gpu_subslices_per_slice>();
+    if (!std::is_same_v<decltype(subslices), uint32_t>) { /* test failed */ }
+}
+----
+=== Intel GPU Number of Execution Units per Subslice
+[source, c++]
+----
+if (dev.has(aspect::ext_intel_gpu_eu_count_per_subslice)) {
+    auto euCount = dev.get_info<ext::intel::info::device::gpu_eu_count_per_subslice>();
+    if (!std::is_same_v<decltype(euCount), uint32_t>) { /* test failed */ }
+}
+----
+=== Intel GPU Number of hardware threads per EU
+[source, c++]
+----
+if (dev.has(aspect::ext_intel_gpu_hw_threads_per_eu)) {
+    auto threadsCount = dev.get_info<ext::intel::info::device::gpu_hw_threads_per_eu>();
+    if (!std::is_same_v<decltype(threadsCount), uint32_t>) { /* test failed */ }
+}
+----
+=== Maximum Memory Bandwidth
+[source, c++]
+----
+if (dev.has(aspect::ext_intel_max_mem_bandwidth)) {
+    auto maxBW = dev.get_info<ext::intel::info::device::max_mem_bandwidth>();
+    if (!std::is_same_v<decltype(maxBW), uint64_t>) { /* test failed */ }
+}
+----
+=== Free Global Memory
+[source, c++]
+----
+if (dev.has(aspect::ext_intel_free_memory)) {
+    auto FreeMemory = dev.get_info<ext::intel::info::device::free_memory>();
+    if (!std::is_same_v<decltype(FreeMemory), uint64_t>) { /* test failed */ }
+}
+----
+=== Memory Clock Rate
+[source, c++]
+----
+if (dev.has(aspect::ext_intel_memory_clock_rate)) {
+    auto MemoryClockRate = dev.get_info<ext::intel::info::device::memory_clock_rate>();
+    if (!std::is_same_v<decltype(MemoryClockRate), uint32_t>) { /* test failed */ }
+}
+----
+=== Memory Bus Width
+[source, c++]
+----
+if (dev.has(aspect::ext_intel_memory_bus_width)) {
+    auto MemoryBusWidth = dev.get_info<ext::intel::info::device::memory_bus_width>();
+    if (!std::is_same_v<decltype(MemoryBusWidth), uint32_t>) { /* test failed */ }
+}
+----
+
+These tests will verify that `ext::intel::info::device` and `aspect` namespaces has these information descriptors.