[SYCL] Embed fsycl-id-queries-range in device image as property.

clang/lib/Driver/ToolChains/Clang.cpp

@@ -11147,6 +11147,13 @@ static void getNonTripleBasedSYCLPostLinkOpts(const ToolChain &TC,
   if (allowDeviceImageDependencies(TCArgs))
     addArgs(PostLinkArgs, TCArgs, {"-allow-device-image-dependencies"});
 
+  // Forward -fsycl-id-queries-range= to sycl-post-link.
+  if (Arg *A = TCArgs.getLastArg(options::OPT_fsycl_id_queries_range_EQ)) {
+    SmallString<64> IdQueriesRangeOpt("-id-queries-range=");
+    IdQueriesRangeOpt += A->getValue();
+    addArgs(PostLinkArgs, TCArgs, {IdQueriesRangeOpt.str()});
+  }
+
   // For bfloat16 conversions LLVM IR devicelib, we only need to embed it
   // when non-AOT compilation is used.
   if (TC.getTriple().isSPIROrSPIRV() && !TC.getTriple().isSPIRAOT()) {

llvm/include/llvm/SYCLPostLink/ComputeModuleRuntimeInfo.h

@@ -27,6 +27,7 @@ struct GlobalBinImageProps {
   bool EmitExportedSymbols;
   bool EmitImportedSymbols;
   bool EmitDeviceGlobalPropSet;
+  int IdQueriesRange; // 0 = int, 1 = uint, 2 = size_t (default)
 };
 bool isModuleUsingAsan(const Module &M);
 bool isModuleUsingMsan(const Module &M);

llvm/lib/SYCLPostLink/ComputeModuleRuntimeInfo.cpp

@@ -377,6 +377,10 @@ PropSetRegTy computeModuleProperties(const Module &M,
     if (OptLevel != -1)
       PropSet.add(PropSetRegTy::SYCL_MISC_PROP, "optLevel", OptLevel);
   }
+  {
+    PropSet.add(PropSetRegTy::SYCL_MISC_PROP, "idQueriesRange",
+                GlobProps.IdQueriesRange);
+  }
   {
     std::vector<std::pair<StringRef, int>> ArgPos =
         getKernelNamesUsingImplicitLocalMem(M);

llvm/tools/sycl-post-link/sycl-post-link.cpp

@@ -241,6 +241,19 @@ cl::opt<bool> AllowDeviceImageDependencies{
     cl::desc("Allow dependencies between device images"), cl::cat(PostLinkCat),
     cl::init(false)};
 
+enum IdQueriesRangeMode { IDQR_INT = 0, IDQR_UINT = 1, IDQR_SIZE_T = 2 };
+
+cl::opt<IdQueriesRangeMode> IdQueriesRange{
+    "id-queries-range",
+    cl::desc("Specify the assumption about SYCL ID query value ranges"),
+    cl::Optional,
+    cl::init(IDQR_INT),
+    cl::values(
+        clEnumValN(IDQR_INT, "int", "ID query values fit within MAX_INT"),
+        clEnumValN(IDQR_UINT, "uint", "ID query values fit within MAX_UINT"),
+        clEnumValN(IDQR_SIZE_T, "size_t", "No restriction on ID query values")),
+    cl::cat(PostLinkCat)};
+
 struct IrPropSymFilenameTriple {
   std::string Ir;
   std::string Prop;
@@ -311,9 +324,13 @@ Error saveModule(
       continue;
     auto CopyTriple = BaseTriple;
     if (DoPropGen) {
-      GlobalBinImageProps Props = {EmitKernelParamInfo, EmitProgramMetadata,
-                                   EmitKernelNames,     EmitExportedSymbols,
-                                   EmitImportedSymbols, DeviceGlobals};
+      GlobalBinImageProps Props = {EmitKernelParamInfo,
+                                   EmitProgramMetadata,
+                                   EmitKernelNames,
+                                   EmitExportedSymbols,
+                                   EmitImportedSymbols,
+                                   DeviceGlobals,
+                                   static_cast<int>(IdQueriesRange)};
       StringRef Target = OutputFile.Target;
       std::string NewSuff = Suffix.str();
       if (!Target.empty())

sycl/include/sycl/handler.hpp

@@ -15,7 +15,6 @@
 #include <sycl/detail/defines_elementary.hpp>
 #include <sycl/detail/export.hpp>
 #include <sycl/detail/get_device_kernel_info.hpp>
-#include <sycl/detail/id_queries_fit_in_int.hpp>
 #include <sycl/detail/impl_utils.hpp>
 #include <sycl/detail/kernel_desc.hpp>
 #include <sycl/detail/kernel_launch_helper.hpp>
@@ -875,7 +874,6 @@ class __SYCL_EXPORT handler {
       // kernel use items/ids in the user range, which means that
       // ID range assumptions can still be violated. So check the bounds
       // of the user range, instead of the rounded range.
-      detail::checkValueRange<Dims>(UserRange);
       convertToRangeViewAndSetDescriptor(RoundedRange);
       StoreLambda<KName, decltype(Wrapper), Dims, TransformedArgType>(
           std::move(Wrapper));
@@ -903,7 +901,6 @@ class __SYCL_EXPORT handler {
       verifyUsedKernelBundleInternal(Info.Name);
       setKernelLaunchProperties(
           detail::extractKernelProperties<Info.IsESIMD>(Props));
-      detail::checkValueRange<Dims>(UserRange);
       convertToRangeViewAndSetDescriptor(std::move(UserRange));
       StoreLambda<NameT, KernelType, Dims, TransformedArgType>(
           std::move(KernelFunc));
@@ -929,7 +926,6 @@ class __SYCL_EXPORT handler {
 #ifndef __SYCL_DEVICE_ONLY__
     throwIfActionIsCreated();
     setDeviceKernelInfo(std::move(Kernel));
-    detail::checkValueRange<Dims>(NumWorkItems);
     convertToRangeViewAndSetDescriptor(std::move(NumWorkItems));
     setKernelLaunchProperties(detail::extractKernelProperties(Props));
     extractArgsAndReqs();
@@ -952,7 +948,6 @@ class __SYCL_EXPORT handler {
 #ifndef __SYCL_DEVICE_ONLY__
     throwIfActionIsCreated();
     setDeviceKernelInfo(std::move(Kernel));
-    detail::checkValueRange<Dims>(NDRange);
     convertToRangeViewAndSetDescriptor(std::move(NDRange));
     setKernelLaunchProperties(detail::extractKernelProperties(Props));
     extractArgsAndReqs();
@@ -988,7 +983,6 @@ class __SYCL_EXPORT handler {
     throwIfActionIsCreated();
     verifyUsedKernelBundleInternal(Info.Name);
 
-    detail::checkValueRange<Dims>(params...);
     if constexpr (SetNumWorkGroups) {
       convertToRangeViewAndSetDescriptor(std::move(params)...,
                                          /*SetNumWorkGroups=*/true);
@@ -1457,7 +1451,6 @@ class __SYCL_EXPORT handler {
 #ifndef __SYCL_DEVICE_ONLY__
     throwIfActionIsCreated();
     setDeviceKernelInfo(std::move(Kernel));
-    detail::checkValueRange<Dims>(NumWorkItems, WorkItemOffset);
     setNDRangeDescriptor(std::move(NumWorkItems), std::move(WorkItemOffset));
     extractArgsAndReqs();
 #endif

sycl/include/sycl/queue.hpp

@@ -18,7 +18,6 @@
 #include <sycl/detail/common.hpp>             // for code_location
 #include <sycl/detail/defines_elementary.hpp> // for __SYCL2020_DEP...
 #include <sycl/detail/export.hpp>             // for __SYCL_EXPORT
-#include <sycl/detail/id_queries_fit_in_int.hpp> // for checkValueRange
 #include <sycl/detail/info_desc_helpers.hpp>  // for is_queue_info_...
 #include <sycl/detail/kernel_desc.hpp>        // for KernelInfo
 #include <sycl/detail/nd_range_view.hpp>
@@ -3988,10 +3987,6 @@ auto submit_kernel_direct_parallel_for(const queue &Queue, nd_range<Dims> Range,
       "must be either sycl::nd_item or be convertible from sycl::nd_item");
   using TransformedArgType = sycl::nd_item<Dims>;
 
-#ifndef __SYCL_DEVICE_ONLY__
-  detail::checkValueRange<Dims>(Range);
-#endif
-
   return submit_kernel_direct<detail::WrapAs::parallel_for, TransformedArgType,
                               NameT, EventNeeded, PropertiesT,
                               KernelTypeUniversalRef, Dims>(
@@ -4058,14 +4053,6 @@ auto submit_kernel_direct_parallel_for(const queue &Queue, range<Dims> Range,
     using KTypeWrapper = decltype(Wrapper);
     using KName = std::conditional_t<std::is_same<KernelType, NameT>::value,
                                      KTypeWrapper, NameWT>;
-#ifndef __SYCL_DEVICE_ONLY__
-    // We are executing over the rounded range, but there are still
-    // items/ids that are constructed in the range rounded
-    // kernel, use items/ids in the user range, which means that
-    // ID range assumptions can still be violated. So check the bounds
-    // of the user range, instead of the rounded range.
-    detail::checkValueRange<Dims>(Range);
-#endif
     return submit_kernel_direct<detail::WrapAs::parallel_for,
                                 TransformedArgType, KName, EventNeeded,
                                 PropertiesT, KTypeWrapper, Dims>(
@@ -4076,9 +4063,6 @@ auto submit_kernel_direct_parallel_for(const queue &Queue, range<Dims> Range,
        // SYCL_LANGUAGE_VERSION >= 202012L
   {
 #ifndef __SYCL_FORCE_PARALLEL_FOR_RANGE_ROUNDING__
-#ifndef __SYCL_DEVICE_ONLY__
-    detail::checkValueRange<Dims>(Range);
-#endif
     return submit_kernel_direct<detail::WrapAs::parallel_for,
                                 TransformedArgType, NameT, EventNeeded,
                                 PropertiesT, KernelTypeUniversalRef, Dims>(

sycl/source/detail/device_binary_image.cpp

@@ -592,7 +592,9 @@ DynRTDeviceBinaryImage::DynRTDeviceBinaryImage(
       Imgs,
       [](const RTDeviceBinaryImage &Img) { return Img.getMiscProperties(); },
       /*IgnoreDuplicates=*/true, /*DropProperty=*/
-      [](std::string_view PropertyName) { return PropertyName == "optLevel"; });
+      [](std::string_view PropertyName) {
+        return PropertyName == "optLevel" || PropertyName == "idQueriesRange";
+      });
 
   std::array<const std::unordered_map<std::string_view,
                                       const sycl_device_binary_property> *,