[SYCL] Choose image with inlined default values if default value is s…

…et explicitely
intel · Feb 6, 2024 · 6675b67 · 6675b67
1 parent 2c85e99
commit 6675b67
Show file tree

Hide file tree

Showing 2 changed files with 64 additions and 4 deletions.
diff --git a/sycl/source/detail/device_image_impl.hpp b/sycl/source/detail/device_image_impl.hpp
@@ -149,13 +149,35 @@ class device_image_impl {
     if (MSpecConstSymMap.count(std::string{SpecName}) == 0)
       return;
 
+    const uint8_t *DefVals = nullptr;
+    if (MBinImage) {
+      // get default values for specialization constants
+      const RTDeviceBinaryImage::PropertyRange &SCDefValRange =
+          MBinImage->getSpecConstantsDefaultValues();
+      bool HasDefaultValues = SCDefValRange.begin() != SCDefValRange.end();
+      if (HasDefaultValues) {
+        ByteArray DefValDescriptors =
+            DeviceBinaryProperty(*SCDefValRange.begin()).asByteArray();
+        assert(DefValDescriptors.size() - 8 == MSpecConstsBlob.size() &&
+               "Specialization constant default value blob do not have the "
+               "expected size.");
+        DefVals = &DefValDescriptors[8];
+      }
+    }
+
     std::vector<SpecConstDescT> &Descs =
         MSpecConstSymMap[std::string{SpecName}];
     for (SpecConstDescT &Desc : Descs) {
-      Desc.IsSet = true;
-      std::memcpy(MSpecConstsBlob.data() + Desc.BlobOffset,
-                  static_cast<const char *>(Value) + Desc.CompositeOffset,
-                  Desc.Size);
+      if (DefVals) {
+        if (std::memcmp(DefVals + Desc.BlobOffset,
+                        static_cast<const char *>(Value) + Desc.CompositeOffset,
+                        Desc.Size) != 0) {
+          Desc.IsSet = true;
+          std::memcpy(MSpecConstsBlob.data() + Desc.BlobOffset,
+                      static_cast<const char *>(Value) + Desc.CompositeOffset,
+                      Desc.Size);
+        }
+      }
     }
   }
 

diff --git a/sycl/test-e2e/SpecConstants/2020/image_selection.cpp b/sycl/test-e2e/SpecConstants/2020/image_selection.cpp
@@ -23,12 +23,22 @@
 // RUN: env SYCL_PI_TRACE=-1 %{run} %t3.out | FileCheck --match-full-lines --check-prefix=CHECK-MIX %s
 // clang-format on
 
+// Check the behaviour when -fsycl-add-default-spec-consts-image option is used
+// and default value is overriden with the same value - we are supposed to
+// choose images with inlined values in this case.
+
+// clang-format off
+// RUN: %clangxx  -fsycl-add-default-spec-consts-image -fsycl -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen %gpu_aot_target_opts %s -o %t3.out
+// RUN: env SYCL_PI_TRACE=-1 %{run} %t3.out | FileCheck --match-full-lines --check-prefix=CHECK-DEFAULT-OVERRIDEN %s
+// clang-format on
+
 #include <sycl/sycl.hpp>
 
 constexpr sycl::specialization_id<int> int_id(3);
 
 class Kernel1;
 class Kernel2;
+class Kernel3;
 
 int main() {
   sycl::queue Q;
@@ -189,5 +199,33 @@ int main() {
   else
     std::cout << "Default value of specialization constant was used."
               << std::endl;
+
+  // Test that if user calls set_specialization_constant with the value equal to
+  // default then we choose image with inlined default values of specialization
+  // constants.
+  // CHECK-DEFAULT-OVERRIDEN: Default value overriden
+  // CHECK-DEFAULT-OVERRIDEN: ---> piextKernelSetArgMemObj(
+  // CHECK-DEFAULT-OVERRIDEN-NEXT:	<unknown> : {{.*}}
+  // CHECK-DEFAULT-OVERRIDEN-NEXT:	<unknown> : {{.*}}
+  // CHECK-DEFAULT-OVERRIDEN-NEXT:	<unknown> : {{.*}}
+  // CHECK-DEFAULT-OVERRIDEN-NEXT:	<unknown> : 0
+  // CHECK-DEFAULT-OVERRIDEN-NEXT: ) ---> 	pi_result : PI_SUCCESS
+  // CHECK-DEFAULT-OVERRIDEN: Default value of specialization constant was used.
+  std::cout << "Default value overriden" << std::endl;
+  Q.submit([&](sycl::handler &cgh) {
+     cgh.set_specialization_constant<int_id>(3);
+
+     cgh.single_task<Kernel3>([=](sycl::kernel_handler h) {
+       auto SpecConst = h.get_specialization_constant<int_id>();
+       *Res = SpecConst == 3 ? 0 : 1;
+     });
+   }).wait();
+
+  if (*Res)
+    std::cout << "New specialization constant value was set." << std::endl;
+  else
+    std::cout << "Default value of specialization constant was used."
+              << std::endl;
+
   return 0;
 }