Skip to content

Options: add the documented functionControl spelling as an alias #424

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
pvelesko wants to merge 2 commits into
base: master
Choose a base branch
from
+69 −1
Open

Options: add the documented functionControl spelling as an alias #424

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
0e090e0
Tests: the documented -ze-intel-functionControl spelling is ignored
pvelesko Jul 31, 2026
3763c27
Options: add the documented functionControl spelling as an alias
pvelesko Jul 31, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
3 changes: 3 additions & 0 deletions IGC/Options/API-Options-Internal.md
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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` |
Expand Down Expand Up @@ -115,6 +116,7 @@ The `cl-` and `ze-` prefixes correspond to OpenCL and Level Zero specific option
|`-<cl-\|ze->intel-force-enable-a64WA`| | `KIND_FLAG` |
|`-<cl-\|ze->intel-force-global-mem-allocation`| | `KIND_FLAG` |
|`-<cl-\|ze->intel-fp64-gen-emu`| Enable full FP64 emulation. | `KIND_FLAG` |
|`-<cl-\|ze->intel-functionControl`| | `KIND_SEPARATE` |
|`-<cl-\|ze->intel-functonControl`| | `KIND_SEPARATE` |
|`-<cl-\|ze->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` |
|`-<cl-\|ze->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` |
Expand Down Expand Up @@ -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` |
Expand Down
5 changes: 4 additions & 1 deletion IGC/Options/include/igc/Options/IGCInternalOptions.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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 [<n>] -ze-intel-functionControl [<n>]
// -cl-intel-functonControl [<n>] -ze-intel-functonControl [<n>]
defm functonControl : CommonSeparate<"functonControl">;
// Correctly spelled alias, not a rename: the misspelled form keeps working.
// -cl-intel-functionControl [<n>] -ze-intel-functionControl [<n>]
defm : CommonSeparate<"functionControl">, Alias<functonControl_common>;

defm fail_on_spill : CommonFlag<"fail-on-spill">;

Expand Down
62 changes: 62 additions & 0 deletions IGC/ocloc_tests/features/function_control/function_control_option_spelling.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -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 [<n>] -ze-intel-functionControl [<n>]
// 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]);
}