[SYCL] Add normalized -fsycl-id-queries-range= option

clang/include/clang/Basic/LangOptions.def

@@ -277,12 +277,13 @@ LANGOPT(DeclareSPIRVBuiltins, 1, 0, NotCompatible, "Declare SPIR-V builtin funct
 LANGOPT(SYCLExplicitSIMD  , 1, 0, NotCompatible, "SYCL compilation with explicit SIMD extension")
 LANGOPT(EnableDAEInSpirKernels , 1, 0, NotCompatible, "Enable Dead Argument Elimination in SPIR kernels")
 LANGOPT(SYCLDecomposeStruct, 1, 1, NotCompatible, "Force top level decomposition of SYCL functor")
-LANGOPT(
-    SYCLValueFitInMaxInt, 1, 1,
+ENUM_LANGOPT(
+    SYCLIdQueriesRange, SYCLIdQueriesRangeKind, 2,
+    SYCLIdQueriesRangeKind::Int,
     NotCompatible,
-    "SYCL compiler assumes value fits within MAX_INT for member function of "
-    "get/operator[], get_id/operator[] and get_global_id/get_global_linear_id "
-    "in SYCL class id, iterm and nd_iterm")
+    "SYCL compiler assumption about the range of values returned by ID query "
+    "member functions (get/operator[], get_id/operator[], "
+    "get_global_id/get_global_linear_id) in SYCL class id, item and nd_item")
 LANGOPT(SYCLCUDACompat, 1, 0,
         NotCompatible,
         "Enable CUDA definitions and implicit includes when building for the "

clang/include/clang/Basic/LangOptions.h

@@ -203,6 +203,12 @@ class LangOptionsBase {
     Force,
   };
 
+  enum class SYCLIdQueriesRangeKind {
+    Int,
+    UInt,
+    SizeT,
+  };
+
   enum HLSLLangStd {
     HLSL_Unset = 0,
     HLSL_2015 = 2015,

clang/include/clang/Options/Options.td

@@ -7701,12 +7701,24 @@ def flink_huge_device_code : Flag<["-"], "flink-huge-device-code">,
 def fno_link_huge_device_code : Flag<["-"], "fno-link-huge-device-code">,
   HelpText<"Do not generate or use a custom linker script for huge device "
            "code sections (default)">;
-defm sycl_id_queries_fit_in_int: BoolFOption<"sycl-id-queries-fit-in-int",
-  LangOpts<"SYCLValueFitInMaxInt">, DefaultTrue,
-  PosFlag<SetTrue, [], [ClangOption, CLOption], "Assume">,
-  NegFlag<SetFalse, [], [ClangOption, CLOption], "Do not assume">,
-  BothFlags<[], [ClangOption, CLOption, CC1Option], " that SYCL ID queries fit "
-            "within MAX_INT.">>;
+def fsycl_id_queries_range_EQ : Joined<["-"], "fsycl-id-queries-range=">,
+  Visibility<[ClangOption, CLOption, CC1Option]>,
+  Values<"int,uint,size_t">,
+  NormalizedValuesScope<"LangOptions::SYCLIdQueriesRangeKind">,
+  NormalizedValues<["Int", "UInt", "SizeT"]>,
+  MarshallingInfoEnum<LangOpts<"SYCLIdQueriesRange">, "Int">,
+  HelpText<"Specify the assumption about SYCL ID query value ranges. "
+           "Valid values: int (queries fit in signed int, default), "
+           "uint (queries fit in unsigned int), "
+           "size_t (queries fit in size_t, no additional assumption)">;
+def fsycl_id_queries_fit_in_int : Flag<["-"], "fsycl-id-queries-fit-in-int">,
+  Alias<fsycl_id_queries_range_EQ>, AliasArgs<["int"]>,
+  HelpText<"Assume that SYCL ID queries fit within MAX_INT "
+           "(compatibility alias for -fsycl-id-queries-range=int)">;
+def fno_sycl_id_queries_fit_in_int : Flag<["-"], "fno-sycl-id-queries-fit-in-int">,
+  Alias<fsycl_id_queries_range_EQ>, AliasArgs<["size_t"]>,
+  HelpText<"Do not assume SYCL ID queries fit within MAX_INT "
+           "(compatibility alias for -fsycl-id-queries-range=size_t)">;
 def fsycl_device_obj_EQ : Joined<["-"], "fsycl-device-obj=">,
   Values<"spirv,llvmir,asm">, HelpText<"Specify format of device code stored "
   "in the resulting object. Valid values are: spirv, asm, llvmir (default)">;

clang/lib/Driver/ToolChains/Clang.cpp

@@ -5714,8 +5714,10 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
       CmdArgs.push_back("-sycl-std=2020");
     }
 
-    if (Arg *A = Args.getLastArg(options::OPT_fsycl_id_queries_fit_in_int,
-                                 options::OPT_fno_sycl_id_queries_fit_in_int))
+    // Pass SYCL ID queries range option to cc1. The flags
+    // -fsycl-id-queries-fit-in-int and -fno-sycl-id-queries-fit-in-int are
+    // aliases that get rewritten to -fsycl-id-queries-range=.
+    if (Arg *A = Args.getLastArg(options::OPT_fsycl_id_queries_range_EQ))
       A->render(Args, CmdArgs);
 
     if (Args.hasArg(options::OPT_fpreview_breaking_changes))

clang/lib/Driver/ToolChains/Cuda.cpp

@@ -973,10 +973,17 @@ void CudaToolChain::addClangTargetOptions(
 
     CC1Args.append({"-mllvm", "-enable-memcpyopt-without-libcalls"});
 
-    if (DriverArgs.hasFlag(options::OPT_fsycl_id_queries_fit_in_int,
-                           options::OPT_fno_sycl_id_queries_fit_in_int, false))
-      CC1Args.append(
-          {"-mllvm", "-nvvm-reflect-add=__CUDA_ID_QUERIES_FIT_IN_INT=1"});
+    // Add NVVM reflect flags based on SYCL ID queries range assumption.
+    if (DriverArgs.hasArg(options::OPT_fsycl_id_queries_range_EQ)) {
+      StringRef RangeValue =
+          DriverArgs.getLastArgValue(options::OPT_fsycl_id_queries_range_EQ);
+      if (RangeValue == "int")
+        CC1Args.append(
+            {"-mllvm", "-nvvm-reflect-add=__CUDA_ID_QUERIES_FIT_IN_INT=1"});
+      else if (RangeValue == "uint")
+        CC1Args.append(
+            {"-mllvm", "-nvvm-reflect-add=__CUDA_ID_QUERIES_FIT_IN_UINT=1"});
+    }
 
     SYCLInstallation.addLibspirvLinkArgs(getEffectiveTriple(), DriverArgs,
                                          HostTC.getTriple(), CC1Args);

clang/lib/Frontend/InitPreprocessor.cpp

@@ -568,8 +568,18 @@ static void InitializeStandardPredefinedMacros(const TargetInfo &TI,
          getSYCLVersionMacros(LangOpts))
       Builder.defineMacro(Macro.first, Macro.second);
 
-    if (LangOpts.SYCLValueFitInMaxInt)
+    // Define macros based on SYCL ID queries range assumption
+    switch (LangOpts.getSYCLIdQueriesRange()) {
+    case LangOptions::SYCLIdQueriesRangeKind::Int:
       Builder.defineMacro("__SYCL_ID_QUERIES_FIT_IN_INT__");
+      break;
+    case LangOptions::SYCLIdQueriesRangeKind::UInt:
+      Builder.defineMacro("__SYCL_ID_QUERIES_FIT_IN_UINT__");
+      break;
+    case LangOptions::SYCLIdQueriesRangeKind::SizeT:
+      // No macro defined - queries fit in size_t per SYCL spec
+      break;
+    }
 
     // Set __SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__ macro for
     // both host and device compilations if -fsycl-disable-range-rounding

clang/test/Driver/sycl-nvptx-id-queries-fit-in-int.cpp

@@ -1,16 +1,44 @@
 // REQUIRES: nvptx-registered-target
 
+// Test default behavior (no flag specified - should not add NVVM flag)
 // RUN: %clang -### -fno-sycl-libspirv -nocudalib \
 // RUN:   -fsycl -fsycl-targets=nvptx64-nvidia-cuda %s 2>&1 \
 // RUN: | FileCheck --check-prefix=CHECK-DEFAULT %s
 
+// Test new normalized option with different values
 // RUN: %clang -### -fno-sycl-libspirv -nocudalib \
-// RUN:   -fsycl -fsycl-targets=nvptx64-nvidia-cuda -fno-sycl-id-queries-fit-in-int %s 2>&1 \
+// RUN:   -fsycl -fsycl-targets=nvptx64-nvidia-cuda -fsycl-id-queries-range=int %s 2>&1 \
+// RUN: | FileCheck --check-prefix=CHECK-INT %s
+
+// RUN: %clang -### -fno-sycl-libspirv -nocudalib \
+// RUN:   -fsycl -fsycl-targets=nvptx64-nvidia-cuda -fsycl-id-queries-range=uint %s 2>&1 \
+// RUN: | FileCheck --check-prefix=CHECK-UINT %s
+
+// RUN: %clang -### -fno-sycl-libspirv -nocudalib \
+// RUN:   -fsycl -fsycl-targets=nvptx64-nvidia-cuda -fsycl-id-queries-range=size_t %s 2>&1 \
 // RUN: | FileCheck --check-prefix=CHECK-DEFAULT %s
 
+// Test legacy compatibility flags
 // RUN: %clang -### -fno-sycl-libspirv -nocudalib \
 // RUN:   -fsycl -fsycl-targets=nvptx64-nvidia-cuda -fsycl-id-queries-fit-in-int %s 2>&1 \
 // RUN: | FileCheck --check-prefix=CHECK-INT %s
 
+// RUN: %clang -### -fno-sycl-libspirv -nocudalib \
+// RUN:   -fsycl -fsycl-targets=nvptx64-nvidia-cuda -fno-sycl-id-queries-fit-in-int %s 2>&1 \
+// RUN: | FileCheck --check-prefix=CHECK-DEFAULT %s
+
+// Test precedence: last one wins
+// RUN: %clang -### -fno-sycl-libspirv -nocudalib \
+// RUN:   -fsycl -fsycl-targets=nvptx64-nvidia-cuda -fsycl-id-queries-range=uint -fsycl-id-queries-fit-in-int %s 2>&1 \
+// RUN: | FileCheck --check-prefix=CHECK-INT %s
+
+// RUN: %clang -### -fno-sycl-libspirv -nocudalib \
+// RUN:   -fsycl -fsycl-targets=nvptx64-nvidia-cuda -fsycl-id-queries-fit-in-int -fsycl-id-queries-range=size_t %s 2>&1 \
+// RUN: | FileCheck --check-prefix=CHECK-DEFAULT %s
+
 // CHECK-INT: "-mllvm" "-nvvm-reflect-add=__CUDA_ID_QUERIES_FIT_IN_INT=1"
+// CHECK-INT-NOT: "-nvvm-reflect-add=__CUDA_ID_QUERIES_FIT_IN_UINT=1"
+// CHECK-UINT: "-mllvm" "-nvvm-reflect-add=__CUDA_ID_QUERIES_FIT_IN_UINT=1"
+// CHECK-UINT-NOT: "-nvvm-reflect-add=__CUDA_ID_QUERIES_FIT_IN_INT=1"
 // CHECK-DEFAULT-NOT: "-nvvm-reflect-add=__CUDA_ID_QUERIES_FIT_IN_INT=1"
+// CHECK-DEFAULT-NOT: "-nvvm-reflect-add=__CUDA_ID_QUERIES_FIT_IN_UINT=1"

clang/test/Driver/sycl-specific-args-diagnostics.cpp

@@ -87,6 +87,10 @@
 // RUN:   %clang_cl -### -fsycl-id-queries-fit-in-int  %s 2>&1 \
 // RUN:   | FileCheck -check-prefix=WARNING-UNUSED-ARG -DOPT=-fsycl-id-queries-fit-in-int %s
 
+// Warning should be emitted when using -fsycl-id-queries-range= without -fsycl
+// RUN:   %clang -### -fsycl-id-queries-range=int  %s 2>&1 \
+// RUN:   | FileCheck -check-prefix=WARNING-UNUSED-ARG -DOPT=-fsycl-id-queries-range=int %s
+
 // Warning should be emitted when using -fsycl-instrument-device-code without -fsycl
 // RUN:   %clang -### -fsycl-instrument-device-code  %s 2>&1 \
 // RUN:   | FileCheck -check-prefix=WARNING-UNUSED-ARG -DOPT=-fsycl-instrument-device-code %s

clang/test/Driver/sycl.cpp

@@ -118,17 +118,27 @@
 // SYCL_HELP_ORDER: Emitting help information for opencl-aot
 // SYCL_HELP_ORDER: opencl-aot{{(\.exe)?}}" "--help"
 
-// -fsycl-id-queries-fit-in-int
-// RUN: %clang -### -fsycl -fsycl-id-queries-fit-in-int  %s 2>&1 | FileCheck %s --check-prefix=ID_QUERIES
-// RUN: %clang_cl -### -fsycl -fsycl-id-queries-fit-in-int  %s 2>&1 | FileCheck %s --check-prefix=ID_QUERIES
-// RUN: %clang -### -fsycl-device-only -fsycl-id-queries-fit-in-int  %s 2>&1 | FileCheck %s --check-prefix=ID_QUERIES
-// RUN: %clang_cl -### -fsycl-device-only -fsycl-id-queries-fit-in-int  %s 2>&1 | FileCheck %s --check-prefix=ID_QUERIES
-// RUN: %clang -### -fsycl -fno-sycl-id-queries-fit-in-int  %s 2>&1 | FileCheck %s --check-prefix=NO_ID_QUERIES
-// RUN: %clang_cl -### -fsycl -fno-sycl-id-queries-fit-in-int  %s 2>&1 | FileCheck %s --check-prefix=NO_ID_QUERIES
-// RUN: %clang -### -fsycl-device-only -fno-sycl-id-queries-fit-in-int  %s 2>&1 | FileCheck %s --check-prefix=NO_ID_QUERIES
-// RUN: %clang_cl -### -fsycl-device-only -fno-sycl-id-queries-fit-in-int  %s 2>&1 | FileCheck %s --check-prefix=NO_ID_QUERIES
-// ID_QUERIES: "-fsycl-id-queries-fit-in-int"
-// NO_ID_QUERIES: "-fno-sycl-id-queries-fit-in-int"
+// -fsycl-id-queries-range and compatibility flags
+// Test new normalized option
+// RUN: %clang -### -fsycl -fsycl-id-queries-range=int  %s 2>&1 | FileCheck %s --check-prefix=ID_QUERIES_INT
+// RUN: %clang -### -fsycl -fsycl-id-queries-range=uint  %s 2>&1 | FileCheck %s --check-prefix=ID_QUERIES_UINT
+// RUN: %clang -### -fsycl -fsycl-id-queries-range=size_t  %s 2>&1 | FileCheck %s --check-prefix=ID_QUERIES_SIZET
+// Test legacy compatibility flags (should be rewritten to new option)
+// RUN: %clang -### -fsycl -fsycl-id-queries-fit-in-int  %s 2>&1 | FileCheck %s --check-prefix=ID_QUERIES_INT
+// RUN: %clang_cl -### -fsycl -fsycl-id-queries-fit-in-int  %s 2>&1 | FileCheck %s --check-prefix=ID_QUERIES_INT
+// RUN: %clang -### -fsycl-device-only -fsycl-id-queries-fit-in-int  %s 2>&1 | FileCheck %s --check-prefix=ID_QUERIES_INT
+// RUN: %clang_cl -### -fsycl-device-only -fsycl-id-queries-fit-in-int  %s 2>&1 | FileCheck %s --check-prefix=ID_QUERIES_INT
+// RUN: %clang -### -fsycl -fno-sycl-id-queries-fit-in-int  %s 2>&1 | FileCheck %s --check-prefix=ID_QUERIES_SIZET
+// RUN: %clang_cl -### -fsycl -fno-sycl-id-queries-fit-in-int  %s 2>&1 | FileCheck %s --check-prefix=ID_QUERIES_SIZET
+// RUN: %clang -### -fsycl-device-only -fno-sycl-id-queries-fit-in-int  %s 2>&1 | FileCheck %s --check-prefix=ID_QUERIES_SIZET
+// RUN: %clang_cl -### -fsycl-device-only -fno-sycl-id-queries-fit-in-int  %s 2>&1 | FileCheck %s --check-prefix=ID_QUERIES_SIZET
+// Test precedence: last one wins
+// RUN: %clang -### -fsycl -fsycl-id-queries-fit-in-int -fsycl-id-queries-range=uint  %s 2>&1 | FileCheck %s --check-prefix=ID_QUERIES_UINT
+// RUN: %clang -### -fsycl -fsycl-id-queries-range=uint -fno-sycl-id-queries-fit-in-int  %s 2>&1 | FileCheck %s --check-prefix=ID_QUERIES_SIZET
+// RUN: %clang -### -fsycl -fsycl-id-queries-range=size_t -fsycl-id-queries-fit-in-int  %s 2>&1 | FileCheck %s --check-prefix=ID_QUERIES_INT
+// ID_QUERIES_INT: "-fsycl-id-queries-range=int"
+// ID_QUERIES_UINT: "-fsycl-id-queries-range=uint"
+// ID_QUERIES_SIZET: "-fsycl-id-queries-range=size_t"
 
 // RUN: %clang -### -fsycl  %s 2>&1 | FileCheck %s --check-prefix=DEFAULT_STD
 // RUN: %clangxx -### -fsycl %s 2>&1 | FileCheck %s --check-prefix=DEFAULT_STD

clang/test/Preprocessor/sycl-macro.cpp

@@ -1,10 +1,13 @@
 // RUN: %clang_cc1 %s -E -dM | FileCheck %s
 // RUN: %clang_cc1 %s -fsycl-is-device -E -dM | FileCheck --check-prefix=CHECK-SYCL-ID %s
-// RUN: %clang_cc1 %s -fsycl-id-queries-fit-in-int -fsycl-is-host -E -dM | FileCheck --check-prefix=CHECK-SYCL-STD-2020 %s
-// RUN: %clang_cc1 %s -fsycl-id-queries-fit-in-int -fsycl-is-device -E -dM | FileCheck --check-prefix=CHECK-SYCL-STD-2020 %s
-// RUN: %clang_cc1 %s -fsycl-id-queries-fit-in-int -fsycl-is-device -E -dM | FileCheck --check-prefix=CHECK-SYCL-STD-DEVICE %s
-// RUNx: %clang_cc1 %s -fsycl-id-queries-fit-in-int -fsycl-is-device -E -dM -fms-compatibility | FileCheck --check-prefix=CHECK-MSVC %s
-// RUN: %clang_cc1 -fno-sycl-id-queries-fit-in-int %s -E -dM | FileCheck \
+// Test new option syntax
+// RUN: %clang_cc1 %s -fsycl-id-queries-range=int -fsycl-is-host -E -dM | FileCheck --check-prefix=CHECK-SYCL-STD-2020 %s
+// RUN: %clang_cc1 %s -fsycl-id-queries-range=int -fsycl-is-device -E -dM | FileCheck --check-prefix=CHECK-SYCL-STD-2020 %s
+// RUN: %clang_cc1 %s -fsycl-id-queries-range=int -fsycl-is-device -E -dM | FileCheck --check-prefix=CHECK-SYCL-STD-DEVICE %s
+// RUN: %clang_cc1 %s -fsycl-id-queries-range=uint -fsycl-is-device -E -dM | FileCheck --check-prefix=CHECK-SYCL-UINT %s
+// RUN: %clang_cc1 %s -fsycl-id-queries-range=size_t -fsycl-is-device -E -dM | FileCheck --check-prefix=CHECK-SYCL-SIZET %s
+// RUNx: %clang_cc1 %s -fsycl-id-queries-range=int -fsycl-is-device -E -dM -fms-compatibility | FileCheck --check-prefix=CHECK-MSVC %s
+// RUN: %clang_cc1 -fsycl-id-queries-range=size_t %s -E -dM | FileCheck \
 // RUN: --check-prefix=CHECK-NO-SYCL_FIT_IN_INT %s
 // RUN: %clang_cc1 %s  -triple nvptx64-nvidia-cuda -target-cpu sm_80 -fsycl-is-device -E -dM | FileCheck \
 // RUN: --check-prefix=CHECK-CUDA %s -DARCH_CODE=800
@@ -34,6 +37,12 @@
 // CHECK-NO-SYCL_FIT_IN_INT-NOT:#define __SYCL_ID_QUERIES_FIT_IN_INT__ 1
 // CHECK-SYCL-ID:#define __SYCL_ID_QUERIES_FIT_IN_INT__ 1
 
+// CHECK-SYCL-UINT:#define __SYCL_ID_QUERIES_FIT_IN_UINT__ 1
+// CHECK-SYCL-UINT-NOT:#define __SYCL_ID_QUERIES_FIT_IN_INT__
+
+// CHECK-SYCL-SIZET-NOT:#define __SYCL_ID_QUERIES_FIT_IN_INT__
+// CHECK-SYCL-SIZET-NOT:#define __SYCL_ID_QUERIES_FIT_IN_UINT__
+
 // CHECK-CUDA:#define __SYCL_CUDA_ARCH__ [[ARCH_CODE]]
 // CHECK-CUDA-NOT:#define __CUDA_ARCH__ {{[0-9]+}}
 // CHECK-SYCL-CUDA-COMPAT:#define __CUDA_ARCH__ {{[0-9][0-9]+}}

libclc/libspirv/lib/ptx-nvidiacl/workitem/get_global_id.cl

@@ -11,7 +11,8 @@
 extern int __nvvm_reflect_ocl(constant char *);
 
 _CLC_DEF _CLC_OVERLOAD size_t __spirv_BuiltInGlobalInvocationId(int dim) {
-  if (__nvvm_reflect_ocl("__CUDA_ID_QUERIES_FIT_IN_INT")) {
+  if (__nvvm_reflect_ocl("__CUDA_ID_QUERIES_FIT_IN_INT") ||
+      __nvvm_reflect_ocl("__CUDA_ID_QUERIES_FIT_IN_UINT")) {
     return (uint)__spirv_BuiltInWorkgroupId(dim) *
                (uint)__spirv_BuiltInWorkgroupSize(dim) +
            (uint)__spirv_BuiltInLocalInvocationId(dim) +

sycl/doc/UsersManual.md

@@ -181,6 +181,23 @@ and not recommended to use in production environment.
     Currently has effect only on spir64\* targets.
     Enabled by default.
 
+**`-fsycl-id-queries-range={int,uint,size_t}`**
+
+    Specify the assumption about SYCL ID query value ranges. This affects the
+    following member functions:
+    * id class get() member function and operator[]
+    * item class get_id() member function and operator[]
+    * nd_item class get_global_id()/get_global_linear_id() member functions
+
+    Valid values:
+    * int - assume that ID queries fit within INT_MAX (default)
+    * uint - assume that ID queries fit within UINT_MAX
+    * size_t - queries fit within SIZE_MAX per SYCL spec (no additional assumption)
+
+    The compiler uses these assumptions to optimize code generation. When an
+    assumption is specified, the runtime validates kernel launch parameters and
+    throws an exception if they would violate the assumption.
+
 **`-f[no-]sycl-id-queries-fit-in-int`**
 
     Assume/Do not assume that SYCL ID queries fit within MAX_INT. It assumes
@@ -190,6 +207,11 @@ and not recommended to use in production environment.
     * nd_item class get_global_id()/get_global_linear_id() member functions
     Enabled by default.
 
+    Note: This option is equivalent to `-fsycl-id-queries-range=int` (when
+    enabled) or `-fsycl-id-queries-range=size_t` (when disabled). The normalized
+    `-fsycl-id-queries-range=` option provides additional control with support
+    for unsigned int range.
+
 **`-f[no-]sycl-force-inline-kernel-lambda`**
 
   Enables/Disables inlining of the kernel lambda operator into the compiler

sycl/doc/extensions/proposed/sycl_ext_oneapi_range_type.asciidoc

@@ -168,19 +168,30 @@ struct SmallKernel
 q.parallel_for(2147483648, SmallKernel());
 ```
 
-== Interaction with the {dpcpp} "-fsycl-id-queries-fit-in-int" option
+== Interaction with the {dpcpp} "-fsycl-id-queries-range" option
 
-The `-fsycl-id-queries-fit-in-int` option is specific to the {dpcpp}
+The `-fsycl-id-queries-range=` option is specific to the {dpcpp}
 implementation.
 Therefore, this section that describes the interaction between this extension
 and that option is non-normative and does not apply to other SYCL
 implementations that may support this extension.
 
-If a translation unit is compiled with the `-fsycl-id-queries-fit-in-int`
-option, all kernels and `SYCL_EXTERNAL` functions without an explicitly
-specified `range_type` property are compiled as-if `range_type<int>` was
+The `-fsycl-id-queries-range=` option accepts three values: `int`, `uint`, or
+`size_t`.
+If a translation unit is compiled with `-fsycl-id-queries-range=int` (the
+default) or `-fsycl-id-queries-range=uint`, all kernels and `SYCL_EXTERNAL`
+functions without an explicitly specified `range_type` property are compiled
+as-if `range_type<int>` or `range_type<unsigned int>` (respectively) was
 specified as a property of that kernel or function.
 
+If a translation unit is compiled with `-fsycl-id-queries-range=size_t`, no
+default `range_type` is applied, adhering to the SYCL spec requirement that
+the total number of work-items must be representable as a `size_t`.
+
+The `-fsycl-id-queries-fit-in-int` flag is equivalent to
+`-fsycl-id-queries-range=int`, and `-fno-sycl-id-queries-fit-in-int` is
+equivalent to `-fsycl-id-queries-range=size_t`.
+
 
 == Implementation notes
 

sycl/include/sycl/detail/defines.hpp

@@ -10,12 +10,17 @@
 
 #include <sycl/detail/defines_elementary.hpp> // for __SYCL_ID_QUERIES_FIT_...
 
+// Unified macro for ID query range assumptions
 #if __SYCL_ID_QUERIES_FIT_IN_INT__ && __has_builtin(__builtin_assume)
 #include <climits>
-#define __SYCL_ASSUME_INT(x) __builtin_assume((x) <= INT_MAX)
+#define __SYCL_ASSUME_ID_RANGE(x) __builtin_assume((x) <= INT_MAX)
+#elif __SYCL_ID_QUERIES_FIT_IN_UINT__ && __has_builtin(__builtin_assume)
+#include <climits>
+#define __SYCL_ASSUME_ID_RANGE(x) __builtin_assume((x) <= UINT_MAX)
 #else
-#define __SYCL_ASSUME_INT(x)
-#if __SYCL_ID_QUERIES_FIT_IN_INT__ && !__has_builtin(__builtin_assume)
+#define __SYCL_ASSUME_ID_RANGE(x)
+#if (__SYCL_ID_QUERIES_FIT_IN_INT__ || __SYCL_ID_QUERIES_FIT_IN_UINT__) &&     \
+    !__has_builtin(__builtin_assume)
 #warning "No assumptions will be emitted due to no __builtin_assume available"
 #endif
 #endif

sycl/include/sycl/detail/defines_elementary.hpp

@@ -57,6 +57,10 @@
 #define __SYCL_ID_QUERIES_FIT_IN_INT__ 0
 #endif
 
+#ifndef __SYCL_ID_QUERIES_FIT_IN_UINT__
+#define __SYCL_ID_QUERIES_FIT_IN_UINT__ 0
+#endif
+
 #ifndef __SYCL_DEPRECATED
 #if !defined(__SYCL_DISABLE_DEPRECATION_WARNINGS)
 #define __SYCL_DEPRECATED(message) [[deprecated(message)]]