From 0e090e0969b75f67167d793077fc1a87f512f99f Mon Sep 17 00:00:00 2001 From: Paulius Velesko Date: Fri, 31 Jul 2026 17:49:46 +0300 Subject: [PATCH 1/2] Tests: the documented -ze-intel-functionControl spelling is ignored IGCInternalOptions.td advertises "-cl-intel-functonControl [] -ze-intel-functionControl []" above defm functonControl : CommonSeparate<"functonControl">; CommonSeparate builds every accepted spelling out of that single base string, so the typo is carried by all six prefixes and the documented name exists nowhere. An unrecognized internal option is dropped without a diagnostic, so passing it is silently a no-op. The new ocloc test observes the option's effect rather than just its parsing: FunctionControl=3 is FLAG_FCALL_FORCE_STACKCALL, which makes ProcessFuncAttributes mark the callee noinline + visaStackCall and PrivateMemoryResolution emit "Stack call has been detected". The misspelled spelling produces that warning, the documented spelling does not. Value 1 (FLAG_FCALL_FORCE_INLINE) is also checked so that the integer operand, not merely the option name, is exercised. Only -internal_options is used, so the test does not depend on regkeys. Signed-off-by: Paulius Velesko --- .../function_control_option_spelling.cl | 62 +++++++++++++++++++ 1 file changed, 62 insertions(+) create mode 100644 IGC/ocloc_tests/features/function_control/function_control_option_spelling.cl diff --git a/IGC/ocloc_tests/features/function_control/function_control_option_spelling.cl b/IGC/ocloc_tests/features/function_control/function_control_option_spelling.cl new file mode 100644 index 000000000000..3455d9f16cc0 --- /dev/null +++ b/IGC/ocloc_tests/features/function_control/function_control_option_spelling.cl @@ -0,0 +1,62 @@ +/*========================== begin_copyright_notice ============================ + +Copyright (C) 2026 Intel Corporation + +SPDX-License-Identifier: MIT + +============================= end_copyright_notice ===========================*/ + +// Checks that the FunctionControl internal option is accepted under the name its +// own definition documents: -cl-intel-functionControl / -ze-intel-functionControl. +// +// Background: +// IGCInternalOptions.td defines the option as +// // -cl-intel-functonControl [] -ze-intel-functionControl [] +// defm functonControl : CommonSeparate<"functonControl">; +// CommonSeparate derives every accepted spelling (-cl-, -cl-intel-, -ze-, +// -ze-intel-, -ze-opt- and bare) from that one base string, so the typo in the +// base string is carried by all of them and the spelling promised by the comment +// exists nowhere. llvm::opt drops an unrecognized internal option without a +// diagnostic, so -ze-intel-functionControl was silently ignored. +// +// The option is observed through its effect rather than through parsing alone: +// FunctionControl=3 is FLAG_FCALL_FORCE_STACKCALL, so ProcessFuncAttributes marks +// every non-kernel function noinline + visaStackCall and PrivateMemoryResolution +// then emits the "Stack call has been detected" warning. FunctionControl=1 is +// FLAG_FCALL_FORCE_INLINE and, like the default, leaves no stack call behind. +// Only -internal_options is used, so the test does not depend on regkeys. + +// REQUIRES: dg2-supported + +// Control: with no option at all the callee is inlined, so there is no stack call. +// RUN: ocloc compile -file %s -device dg2 2>&1 | FileCheck %s --check-prefix=CHECK-INLINED + +// The long-standing misspelled spelling has to keep working. +// RUN: ocloc compile -file %s -device dg2 -internal_options "-cl-intel-functonControl 3" 2>&1 | FileCheck %s --check-prefix=CHECK-STACKCALL +// RUN: ocloc compile -file %s -device dg2 -internal_options "-ze-intel-functonControl 3" 2>&1 | FileCheck %s --check-prefix=CHECK-STACKCALL + +// The documented spelling has to work as well. +// RUN: ocloc compile -file %s -device dg2 -internal_options "-cl-intel-functionControl 3" 2>&1 | FileCheck %s --check-prefix=CHECK-STACKCALL +// RUN: ocloc compile -file %s -device dg2 -internal_options "-ze-intel-functionControl 3" 2>&1 | FileCheck %s --check-prefix=CHECK-STACKCALL + +// ... and the value has to be carried through, not just the option name: 1 is +// FLAG_FCALL_FORCE_INLINE and must not produce a stack call. +// RUN: ocloc compile -file %s -device dg2 -internal_options "-ze-intel-functionControl 1" 2>&1 | FileCheck %s --check-prefix=CHECK-INLINED + +// CHECK-INLINED-NOT: Stack call has been detected +// CHECK-INLINED: Build succeeded. + +// CHECK-STACKCALL: warning: in kernel 'fctl': Stack call has been detected +// CHECK-STACKCALL: Build succeeded. + +int helper(int a, int b) { + int r = 0; + for (int i = 0; i < 8; i++) + r += (a ^ (b + i)) * (i + 1); + return r; +} + +__kernel void fctl(__global int *in, __global int *out) { + int i = get_global_id(0); + out[i] = helper(in[i], in[i + 1]); +} From 3763c27510e184f5ce532b330e7ec1eb75e5e0ef Mon Sep 17 00:00:00 2001 From: Paulius Velesko Date: Fri, 31 Jul 2026 17:49:47 +0300 Subject: [PATCH 2/2] Options: add the documented functionControl spelling as an alias CommonSeparate<"functonControl"> derives all six accepted prefixes from one base string, so the typo in that string was carried by every spelling and -ze-intel-functionControl, which the comment right above the defm advertises, matched nothing and was dropped without a diagnostic. Define the correctly spelled name as an alias of functonControl_common rather than renaming the base string, so the misspelled spelling that users depend on today keeps working. This follows the anonymous "defm : Common*<...>, Alias<..._common>" idiom already used in this file for exp-register-file-size=, vector-coalesing=, load-cache-default=, store-cache-default= and the ldstcombine options. The single consumer, OpenCLOptions.cpp, keeps reading OPT_functonControl_common and needs no change; llvm::opt resolves the alias before matching. The comment on the original defm is corrected to name what that line actually defines. API-Options-Internal.md is generated from the .td and lists alias records, so the three new rows it would gain are added alongside. Signed-off-by: Paulius Velesko --- IGC/Options/API-Options-Internal.md | 3 +++ IGC/Options/include/igc/Options/IGCInternalOptions.td | 5 ++++- 2 files changed, 7 insertions(+), 1 deletion(-) diff --git a/IGC/Options/API-Options-Internal.md b/IGC/Options/API-Options-Internal.md index faa1a878f815..ea59e8c9d07c 100644 --- a/IGC/Options/API-Options-Internal.md +++ b/IGC/Options/API-Options-Internal.md @@ -49,6 +49,7 @@ The `cl-` and `ze-` prefixes correspond to OpenCL and Level Zero specific option |`-[cl-\|ze-]force-enable-a64WA`| | `KIND_FLAG` | |`-[cl-\|ze-]force-global-mem-allocation`| | `KIND_FLAG` | |`-[cl-\|ze-]fp64-gen-emu`| Enable full FP64 emulation. | `KIND_FLAG` | +|`-[cl-\|ze-]functionControl`| | `KIND_SEPARATE` | |`-[cl-\|ze-]functonControl`| | `KIND_SEPARATE` | |`-[cl-\|ze-]greater-than-4GB-buffer-required`| When this flag is present, it indicates that any OpenCL buffers can be more than 4GB in size. If it is absent, all buffers are not more than 4GB in size. | `KIND_FLAG` | |`-[cl-\|ze-]has-buffer-offset-arg`| This flag, together with *[-cl-intel\|-ze-opt]-greater-than-4GB-buffer-required* is used to convert stateless memory accesses, called messages or load/store, into stateful ones. The OpenCL runtime can create a surface whose base is either *buffer_base* or *buffer_base + buffer_offset*, based on whether *buffer_offset* is used. | `KIND_FLAG` | @@ -115,6 +116,7 @@ The `cl-` and `ze-` prefixes correspond to OpenCL and Level Zero specific option |`-intel-force-enable-a64WA`| | `KIND_FLAG` | |`-intel-force-global-mem-allocation`| | `KIND_FLAG` | |`-intel-fp64-gen-emu`| Enable full FP64 emulation. | `KIND_FLAG` | +|`-intel-functionControl`| | `KIND_SEPARATE` | |`-intel-functonControl`| | `KIND_SEPARATE` | |`-intel-greater-than-4GB-buffer-required`| When this flag is present, it indicates that any OpenCL buffers can be more than 4GB in size. If it is absent, all buffers are not more than 4GB in size. | `KIND_FLAG` | |`-intel-has-buffer-offset-arg`| This flag, together with *[-cl-intel\|-ze-opt]-greater-than-4GB-buffer-required* is used to convert stateless memory accesses, called messages or load/store, into stateful ones. The OpenCL runtime can create a surface whose base is either *buffer_base* or *buffer_base + buffer_offset*, based on whether *buffer_offset* is used. | `KIND_FLAG` | @@ -180,6 +182,7 @@ The `cl-` and `ze-` prefixes correspond to OpenCL and Level Zero specific option |`-ze-opt-force-enable-a64WA`| | `KIND_FLAG` | |`-ze-opt-force-global-mem-allocation`| | `KIND_FLAG` | |`-ze-opt-fp64-gen-emu`| Enable full FP64 emulation. | `KIND_FLAG` | +|`-ze-opt-functionControl`| | `KIND_SEPARATE` | |`-ze-opt-functonControl`| | `KIND_SEPARATE` | |`-ze-opt-greater-than-4GB-buffer-required`| When this flag is present, it indicates that any OpenCL buffers can be more than 4GB in size. If it is absent, all buffers are not more than 4GB in size. | `KIND_FLAG` | |`-ze-opt-has-buffer-offset-arg`| This flag, together with *[-cl-intel\|-ze-opt]-greater-than-4GB-buffer-required* is used to convert stateless memory accesses, called messages or load/store, into stateful ones. The OpenCL runtime can create a surface whose base is either *buffer_base* or *buffer_base + buffer_offset*, based on whether *buffer_offset* is used. | `KIND_FLAG` | diff --git a/IGC/Options/include/igc/Options/IGCInternalOptions.td b/IGC/Options/include/igc/Options/IGCInternalOptions.td index f95be4035772..94a37235f44a 100644 --- a/IGC/Options/include/igc/Options/IGCInternalOptions.td +++ b/IGC/Options/include/igc/Options/IGCInternalOptions.td @@ -135,8 +135,11 @@ defm skip_reloc_add : CommonFlag<"skip-reloc-add">; // -cl-intel-disableEUFusion -ze-intel-disableEUFusion defm disableEUFusion : CommonFlag<"disableEUFusion">; -// -cl-intel-functonControl [] -ze-intel-functionControl [] +// -cl-intel-functonControl [] -ze-intel-functonControl [] defm functonControl : CommonSeparate<"functonControl">; +// Correctly spelled alias, not a rename: the misspelled form keeps working. +// -cl-intel-functionControl [] -ze-intel-functionControl [] +defm : CommonSeparate<"functionControl">, Alias; defm fail_on_spill : CommonFlag<"fail-on-spill">;