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">; 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]); +}