Skip to content

Commit 1a98d9f

Browse files
amielczaigcbot
authored andcommitted
Provide IGC API option descriptions
Provide IGC API option descriptions
1 parent 0e2022d commit 1a98d9f

File tree

3 files changed

+103
-55
lines changed

3 files changed

+103
-55
lines changed

IGC/Options/include/igc/Options/CommonApiOptions.td

Lines changed: 8 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -26,13 +26,15 @@ defm gtpin_scratch_area_size : CommonSeparate<"gtpin-scratch-area-size">,
2626
defm : CommonJoined<"gtpin-scratch-area-size=">, Alias<gtpin_scratch_area_size_common>,
2727
HelpText<"Alias for -ze-gtpin-scratch-area-size">;
2828

29-
defm gtpin_indir_ref : CommonFlag<"gtpin-indir-ref">;
29+
defm gtpin_indir_ref : CommonFlag<"gtpin-indir-ref">,
30+
HelpText<"Ask finalizer to provide list of registers used by indirect operand per %ip">;
3031

3132
defm skip_fde : CommonFlag<"skip-fde">;
3233

3334
defm no_fusedCallWA : CommonFlag<"no-fusedCallWA">;
3435

35-
defm disable_compaction : CommonFlag<"disable-compaction">;
36+
defm disable_compaction : CommonFlag<"disable-compaction">,
37+
HelpText<"Disable compaction pass in finalizer. This pass is used to decide whether to use a compacted, i.e. shorter encoding of machine instructions wherever ISA allows.">;
3638

3739
def emit_debug : PlainFlag<"g">,
3840
HelpText<"Enable generation of debug information and enables kernel debug">;
@@ -44,8 +46,10 @@ defm opt_disable : CommonFlag<"opt-disable">,
4446
defm library_compilation : CommonFlag<"library-compilation">;
4547

4648
// -library-compile-simd=[8|16|32]
47-
defm library_compile_simd : CommonSeparate<"library-compile-simd">;
48-
defm : CommonJoined<"library-compile-simd=">, Alias<library_compile_simd_common>;
49+
defm library_compile_simd : CommonSeparate<"library-compile-simd">,
50+
HelpText<"Select SIMD size for library compilations [8|16|32]">;
51+
defm : CommonJoined<"library-compile-simd=">, Alias<library_compile_simd_common>,
52+
HelpText<"Select SIMD size for library compilations [8|16|32]">;
4953

5054
defm exp_register_file_size : CommonSeparate<"exp-register-file-size">,
5155
HelpText<"Set amount of registers used by regalloc">;

IGC/Options/include/igc/Options/IGCApiOptions.td

Lines changed: 56 additions & 30 deletions
Original file line numberDiff line numberDiff line change
@@ -14,13 +14,16 @@ let Flags = [IGCApiOption] in {
1414
// Backend API options {{
1515

1616
// -cl-fp32-correctly-rounded-divide-sqrt, -ze-fp32-correctly-rounded-divide-sqrt
17-
defm fp32_correctly_rounded_divide_sqrt : CommonFlag<"fp32-correctly-rounded-divide-sqrt">;
17+
defm fp32_correctly_rounded_divide_sqrt : CommonFlag<"fp32-correctly-rounded-divide-sqrt">,
18+
HelpText<"Allows an application to specify that single precision floating-point divide (x/y and 1/x) and sqrt used in the program source are correctly rounded.">;
1819

1920
// -cl-no-subgroup-ifp, -ze-no-subgroup-ifp
20-
defm no_subgroup_ifp : CommonFlag<"no-subgroup-ifp">;
21+
defm no_subgroup_ifp : CommonFlag<"no-subgroup-ifp">,
22+
HelpText<"This indicates that kernels in this program do not require sub-groups to make independent forward progress.">;
2123

2224
// -cl-uniform-work-group-size, -ze-uniform-work-group-size
23-
defm uniform_work_group_size : CommonFlag<"uniform-work-group-size">;
25+
defm uniform_work_group_size : CommonFlag<"uniform-work-group-size">,
26+
HelpText<"This requires that the global work-size be a multiple of the work-group size specified to clEnqueueNDRangeKernel.">;
2427

2528
// -cl-take-global-address, -ze-take-global-address
2629
defm take_global_address : CommonFlag<"take-global-address">;
@@ -32,10 +35,12 @@ defm required_thread_count : CommonSeparate<"reqd-eu-thread-count">;
3235
defm : CommonJoined<"reqd-eu-thread-count=">, Alias<required_thread_count_common>;
3336

3437
// -ze-opt-large-grf-kernel
35-
defm large_grf_kernel : CommonSeparate<"large-grf-kernel">;
38+
defm large_grf_kernel : CommonSeparate<"large-grf-kernel">,
39+
HelpText<"-ze-opt-large-grf-kernel <string> tells IGC to use large GRF size if kernel name contains <string> regardless of module-level options.">;
3640

3741
// -ze-opt-regular-grf-kernel
38-
defm regular_grf_kernel : CommonSeparate<"regular-grf-kernel">;
42+
defm regular_grf_kernel : CommonSeparate<"regular-grf-kernel">,
43+
HelpText<"-ze-opt-regular-grf-kernel <string> tells IGC to use regular GRF size if kernel name contains <string >regardless of module-level options.">;
3944

4045
// -cl-intel-no-local-to-generic, -ze-opt-no-local-to-generic
4146
defm no_local_to_generic : CommonFlag<"no-local-to-generic">;
@@ -47,14 +52,15 @@ defm 128_grf_per_thread : CommonFlag<"128-GRF-per-thread">;
4752
defm 256_grf_per_thread : CommonFlag<"256-GRF-per-thread">;
4853

4954
// -cl-intel-greater-than-4GB-buffer-required, -ze-opt-greater-than-4GB-buffer-required
50-
defm greater_than_4GB_buffer_required : CommonFlag<"greater-than-4GB-buffer-required">;
55+
defm greater_than_4GB_buffer_required : CommonFlag<"greater-than-4GB-buffer-required">,
56+
HelpText<"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.">;
5157

5258
// -cl-poison-unsupported-fp64-kernels -ze-poison-unsupported-fp64-kernels
5359
defm poison_unsupported_fp64_kernels : CommonFlag<"poison-unsupported-fp64-kernels">;
5460

5561
// -cl-intel-enable-ieee-float-exception-trap
56-
// This flags enables the IEEE exception trap bit in Control Register in the prolog of the kernel.
57-
defm enable_ieee_float_exception_trap : CommonFlag<"enable-ieee-float-exception-trap">;
62+
defm enable_ieee_float_exception_trap : CommonFlag<"enable-ieee-float-exception-trap">,
63+
HelpText<"This flag enables the IEEE exception trap bit in Control Register in the prolog of the kernel.">;
5864

5965
// -cl-fp64-gen-emu -ze-fp64-gen-emu
6066
defm fp64_gen_emu : CommonFlag<"fp64-gen-emu">;
@@ -63,46 +69,64 @@ defm fp64_gen_emu : CommonFlag<"fp64-gen-emu">;
6369
defm fp64_gen_conv_emu : CommonFlag<"fp64-gen-conv-emu">;
6470

6571
// -cl-intel-profile-guided-trimming, -ze-opt-profile-guided-trimming
66-
defm static_profile_guided_trimming : CommonFlag<"static-profile-guided-trimming">;
72+
defm static_profile_guided_trimming : CommonFlag<"static-profile-guided-trimming">,
73+
HelpText<"Enable static analysis in the kernel trimming.">;
6774

6875
// }} Backend API options
6976

7077
// API options from source translation {{
71-
def cl_std : ClJoined<"std=">;
78+
def cl_std : ClJoined<"std=">,
79+
HelpText<"Determine the language version to be accepted by the compiler.">;
7280

73-
def single_precision_constant : ClFlag<"single-precision-constant">;
81+
def single_precision_constant : ClFlag<"single-precision-constant">,
82+
HelpText<"Forces implicit conversions of double-precision floating-point literals to single precision.">;
7483

75-
def denorms_are_zero : ClFlag<"denorms-are-zero">;
84+
def denorms_are_zero : ClFlag<"denorms-are-zero">,
85+
HelpText<"Controls how single precision and double precision denormalized numbers are handled. If specified as a build option, the single precision denormalized numbers may be flushed to zero; double precision denormalized numbers may also be flushed to zero if the optional extension for double precision is supported">;
7686

77-
def strict_aliasing : ClFlag<"strict-aliasing">;
87+
def strict_aliasing : ClFlag<"strict-aliasing">,
88+
HelpText<"Allows the compiler to assume the strictest aliasing rules.">;
7889

79-
def mad_enable : ClFlag<"mad-enable">;
90+
def mad_enable : ClFlag<"mad-enable">,
91+
HelpText<"Allow a * b + c to be replaced by a mad instruction.">;
8092

81-
def no_signed_zeros : ClFlag<"no-signed-zeros">;
93+
def no_signed_zeros : ClFlag<"no-signed-zeros">,
94+
HelpText<"Allow optimizations for floating-point arithmetic that ignore the signedness of zero.">;
8295

83-
def unsafe_math_optimizations : ClFlag<"unsafe-math-optimizations">;
96+
def unsafe_math_optimizations : ClFlag<"unsafe-math-optimizations">,
97+
HelpText<"Allow optimizations for floating-point arithmetic that (a) assume that arguments and results are valid, (b) may violate the IEEE 754 standard, (c) assume relaxed OpenCL numerical compliance requirements as defined in the unsafe math optimization section of the OpenCL C or OpenCL SPIR-V Environment specifications, and (d) may violate edge case behavior in the OpenCL C or OpenCL SPIR-V Environment specifications.">;
8498

85-
def finite_math_only : ClFlag<"finite-math-only">;
99+
def finite_math_only : ClFlag<"finite-math-only">,
100+
HelpText<"Allow optimizations for floating-point arithmetic that assume that arguments and results are not NaNs, +Inf, -Inf. This option may violate the OpenCL numerical compliance requirements for single precision and double precision floating-point, as well as edge case behavior.">;
86101

87-
def fast_relaxed_math : ClFlag<"fast-relaxed-math">;
102+
def fast_relaxed_math : ClFlag<"fast-relaxed-math">,
103+
HelpText<"Sets the optimization options -cl-finite-math-only and -cl-unsafe-math-optimizations. This option causes the preprocessor macro __FAST_RELAXED_MATH__ to be defined in the OpenCL program.">;
88104

89105
def match_sincospi : ClFlag<"match-sincospi">;
90106

91-
def oclfe_w : PlainFlag<"w">;
107+
def oclfe_w : PlainFlag<"w">,
108+
HelpText<"Remove all warning messages.">;
92109

93-
def oclfe_werror : PlainFlag<"Werror">;
110+
def oclfe_werror : PlainFlag<"Werror">,
111+
HelpText<"Treat every warning as an error.">;
94112

95-
def kernel_arg_info : ClFlag<"kernel-arg-info">;
113+
def kernel_arg_info : ClFlag<"kernel-arg-info">,
114+
HelpText<"Allow the compiler to store information about the arguments of a kernel(s) in the program executable. The argument information stored includes the argument name, its type, the address space and access qualifiers used.">;
96115

97-
def oclfe_x : PlainSeparate<"x">;
116+
def oclfe_x : PlainSeparate<"x">,
117+
HelpText<"Manualy provide type of file. Takes only spir or spir64 as argument.">;
98118

99-
def oclfe_D : PlainJoinedOrSeparate<"D">;
119+
def oclfe_D : PlainJoinedOrSeparate<"D">,
120+
HelpText<"Manually define macros.">;
100121

101-
def oclfe_I : PlainJoinedOrSeparate<"I">;
122+
def oclfe_I : PlainJoinedOrSeparate<"I">,
123+
HelpText<"Add directory to the list of directories which will be searched for header files."> ;
102124

103-
def oclfe_spir_std : PlainJoined<"spir-std=">;
125+
def oclfe_spir_std : PlainJoined<"spir-std=">,
126+
HelpText<"Specify the SPIR version.">;
104127

105-
def oclfe_gline_tables_only : PlainFlag<"gline-tables-only">;
128+
def oclfe_gline_tables_only : PlainFlag<"gline-tables-only">,
129+
HelpText<"Generate only line table debug information.">;
106130

107131
def oclfe_triple : PlainSeparate<"triple">;
108132

@@ -112,15 +136,17 @@ def oclfe_dwarf_column_info : PlainFlag<"dwarf-column-info">;
112136
def debug_info : ClFlag<"intel-debug-info">;
113137

114138
def disable_a64wa : ClFlag<"intel-disable-a64WA">;
115-
116139
def oclfe_profiler : PlainFlag<"profiler">;
117140

118-
def oclfe_s : PlainSeparate<"s">;
141+
def oclfe_s : PlainSeparate<"s">,
142+
HelpText<"Strip all symbol table and debug informaton from the output binary.">;
119143

120144
// Additional debug options.
121-
def oclfe_igc_opts : PlainJoinedOrSeparate<"igc_opts">;
145+
def oclfe_igc_opts : PlainJoinedOrSeparate<"igc_opts">,
146+
HelpText<"Pass IGC options delimited by ',' or ' '.">;
122147

123-
def oclfe_dump_opt_llvm : PlainJoined<"dump-opt-llvm=">;
148+
def oclfe_dump_opt_llvm : PlainJoined<"dump-opt-llvm=">,
149+
HelpText<"Dump the llvm output to the specified file.">;
124150
// }} API options from source translation
125151

126152
}

IGC/Options/include/igc/Options/IGCInternalOptions.td

Lines changed: 39 additions & 21 deletions
Original file line numberDiff line numberDiff line change
@@ -14,7 +14,8 @@ let Flags = [IGCInternalOption] in {
1414
// Backend internal options {{
1515

1616
// -cl-replace-global-offsets-by-zero, -ze-replace-global-offsets-by-zero
17-
defm replace_global_offsets_by_zero : CommonFlag<"replace-global-offsets-by-zero">;
17+
defm replace_global_offsets_by_zero : CommonFlag<"replace-global-offsets-by-zero">,
18+
HelpText<"OpenCL's global IDs are assumed to start from the origin at global offsets (offset_x, offset_y, offset_z) When this flag is present, it indicates that the global offsets are (0,0,0).">;
1819

1920
defm kernel_debug_enable : CommonFlag<"kernel-debug-enable">;
2021

@@ -28,16 +29,20 @@ defm include_sip_kernel_local_debug : CommonFlag<"include-sip-kernel-local-debug
2829
defm use_32_bit_ptr_arith : CommonFlag<"use-32bit-ptr-arith">;
2930

3031
// -cl-intel-greater-than-4GB-buffer-required, -ze-opt-greater-than-4GB-buffer-required
31-
defm greater_than_4GB_buffer_required : CommonFlag<"greater-than-4GB-buffer-required">;
32+
defm greater_than_4GB_buffer_required : CommonFlag<"greater-than-4GB-buffer-required">,
33+
HelpText<"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.">;
3234

3335
// -cl-intel-has-buffer-offset-arg, -ze-opt-has-buffer-offset-arg, -ze-intel-has-buffer-offset-arg
34-
defm has_buffer_offset_arg : CommonFlag<"has-buffer-offset-arg">;
36+
defm has_buffer_offset_arg : CommonFlag<"has-buffer-offset-arg">,
37+
HelpText<"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.">;
3538

3639
// -cl-intel-buffer-offset-arg-required, -ze-opt-buffer-offset-arg-required
37-
defm buffer_offset_arg_required : CommonFlag<"buffer-offset-arg-required">;
40+
defm buffer_offset_arg_required : CommonFlag<"buffer-offset-arg-required">,
41+
HelpText<"Tell IGC to always use buffer offset. It is valid only if -intel-has-buffer-offset-arg is present.">;
3842

3943
// -cl-intel-has-positive-pointer-offset, -ze-opt-has-positive-pointer-offset
40-
defm has_positive_pointer_offset : CommonFlag<"has-positive-pointer-offset">;
44+
defm has_positive_pointer_offset : CommonFlag<"has-positive-pointer-offset">,
45+
HelpText<"For any load and store (aka message) whose address = *ptrArg + offset*, where *ptrArg* is a kernel pointer argument, offset is assumed to be non-negative if this flag is present.">;
4146

4247
// -cl-intel-disable-a64WA
4348
defm disable_a64wa : CommonFlag<"disable-a64WA">;
@@ -49,12 +54,13 @@ defm force_enable_a64wa : CommonFlag<"force-enable-a64WA">;
4954
defm no_prera_scheduling : CommonFlag<"no-prera-scheduling">;
5055

5156
// (All start with -cl-intel or -ze-opt)
52-
defm num_thread_per_eu : CommonSeparate<"num-thread-per-eu">;
57+
defm num_thread_per_eu : CommonSeparate<"num-thread-per-eu">,
58+
HelpText<"Overrides the current number of threads value defined by the user's command line option for the entire module or the compiler choice by heuristics.">;
5359

5460
defm exp_register_file_size : CommonSeparate<"exp-register-file-size">,
55-
HelpText<"Set amount of registers used by regalloc">;
61+
HelpText<"Set amount of registers used by regalloc.">;
5662
defm : CommonJoined<"exp-register-file-size=">, Alias<exp_register_file_size_common>,
57-
HelpText<"Alias for -ze-exp-register-file-size">;
63+
HelpText<"Alias for -ze-exp-register-file-size.">;
5864

5965
// -cl-intel-128-GRF-per-thread
6066
defm 128_grf_per_thread : CommonFlag<"128-GRF-per-thread">;
@@ -64,7 +70,7 @@ defm 256_grf_per_thread : CommonFlag<"256-GRF-per-thread">;
6470

6571
// -cl-intel-enable-auto-large-GRF-mode, -ze-opt-enable-auto-large-GRF-mode
6672
defm enable_auto_large_GRF_mode : CommonFlag<"enable-auto-large-GRF-mode">,
67-
HelpText<"Use compiler heuristics to determine number of GRF">;
73+
HelpText<"Use compiler heuristics to determine number of GRF.">;
6874

6975
// -cl-intel-force-global-mem-allocation, -cl-force-global-mem-allocation, -ze-force-global-mem-allocation
7076
defm force_global_mem_allocation : CommonFlag<"force-global-mem-allocation">;
@@ -76,10 +82,12 @@ defm force_global_mem_allocation : CommonFlag<"force-global-mem-allocation">;
7682
defm disable_recompilation : CommonFlag<"disable-recompilation">;
7783

7884
// -cl-intel-force-emu-int32divrem
79-
defm force_emu_int32divrem : CommonFlag<"force-emu-int32divrem">;
85+
defm force_emu_int32divrem : CommonFlag<"force-emu-int32divrem">,
86+
HelpText<"Use emulation fp64-based emulation functions if fp64 is supported natively.">;
8087

8188
// -cl-intel-force-emu-sp-int32divrem
82-
defm force_emu_sp_int32divrem : CommonFlag<"force-emu-sp-int32divrem">;
89+
defm force_emu_sp_int32divrem : CommonFlag<"force-emu-sp-int32divrem">,
90+
HelpText<"Force the fp64-based emulation regardless of native support.">;
8391

8492
// -cl-intel-force-disable-4GB-buffer
8593
defm force_disable_4GB_buffer : CommonFlag<"force-disable-4GB-buffer">;
@@ -101,7 +109,8 @@ defm vector_coalescing : CommonSeparate<"vector-coalesing">;
101109
defm : CommonJoined<"vector-coalesing=">, Alias<vector_coalescing_common>;
102110

103111
// -cl-intel-exclude-ir-from-zebin
104-
defm exclude_ir_from_zebin : CommonFlag<"exclude-ir-from-zebin">;
112+
defm exclude_ir_from_zebin : CommonFlag<"exclude-ir-from-zebin">,
113+
HelpText<"Exclude SPIR-V section from files generated in ZEBIN format.">;
105114

106115
// -cl-intel-no-spill
107116
defm no_spill : CommonFlag<"no-spill">;
@@ -110,10 +119,12 @@ defm no_spill : CommonFlag<"no-spill">;
110119
defm disable_noMaskWA : CommonFlag<"disable-noMaskWA">;
111120

112121
// -cl-intel-ignoreBFRounding, -ze-intel-ignoreBFRounding
113-
defm ignoreBFRounding : CommonFlag<"ignoreBFRounding">;
122+
defm ignoreBFRounding : CommonFlag<"ignoreBFRounding">,
123+
HelpText<"Folds BF operands into mul/add/cmp operations.">;
114124

115125
// -cl-compile-one-at-time
116-
defm compile_one_at_time : CommonFlag<"compile-one-at-time">;
126+
defm compile_one_at_time : CommonFlag<"compile-one-at-time">,
127+
HelpText<"Enables llvm::module splitting to compile only one kernel at a time.">;
117128

118129
// -cl-skip-reloc-add
119130
defm skip_reloc_add : CommonFlag<"skip-reloc-add">;
@@ -135,18 +146,23 @@ defm store_cache_default : CommonSeparate<"store-cache-default">;
135146
defm : CommonJoined<"store-cache-default=">, Alias<store_cache_default_common>;
136147

137148
// -cl-fp64-gen-emu -ze-fp64-gen-emu
138-
defm fp64_gen_emu : CommonFlag<"fp64-gen-emu">;
149+
defm fp64_gen_emu : CommonFlag<"fp64-gen-emu">,
150+
HelpText<"Enable full FP64 emulation.">;
139151

140152
// *-private-memory-minimal-size-per-thread <SIZE>
141-
defm private_memory_minimal_size_per_thread : CommonSeparate<"private-memory-minimal-size-per-thread">;
153+
defm private_memory_minimal_size_per_thread : CommonSeparate<"private-memory-minimal-size-per-thread">,
154+
HelpText<"When this flag is present, it guarantees that size of private memory allocated per thread can not be less then the given value. Constraint: <SIZE> >= 0.">;
142155

143156
// *-scratch-space-private-memory-minimal-size-per-thread <SIZE>
144-
defm scratch_space_private_memory_minimal_size_per_thread : CommonSeparate<"scratch-space-private-memory-minimal-size-per-thread">;
157+
defm scratch_space_private_memory_minimal_size_per_thread : CommonSeparate<"scratch-space-private-memory-minimal-size-per-thread">,
158+
HelpText<"When this flag is present, it guarantees that size of scratch space private memory allocated per thread can not be less then the given value. Constraint: <SIZE> >= 0.">;
145159

146-
defm enable_divergent_barrier_handling : CommonFlag<"enable-divergent-barrier-handling">;
160+
defm enable_divergent_barrier_handling : CommonFlag<"enable-divergent-barrier-handling">,
161+
HelpText<"Enable the divergent barrier pass.">;
147162

148163
// -cl-intel-high-accuracy-nolut-math
149-
defm high_accuracy_nolut_math : CommonFlag<"high-accuracy-nolut-math">;
164+
defm high_accuracy_nolut_math : CommonFlag<"high-accuracy-nolut-math">,
165+
HelpText<"Enbales experimental high accuracy implementations of transcendentals.">;
150166

151167
// -[cl-intel|ze-opt]-ldstcombine=[0|1]
152168
// -[cl-intel|ze-opt]-ldstcombine-max-storebytes=[4|8|16|32]
@@ -163,13 +179,15 @@ defm : CommonJoined<"ldstcombine-max-loadbytes=">, Alias<ldstcombine_max_loadbyt
163179
// Internal options from source translation {{
164180
def oclfe_ocl_version : PlainJoined<"ocl-version=">;
165181

166-
def oclfe_force_cl_std : PlainFlag<"force-cl-std">;
182+
def oclfe_force_cl_std : PlainFlag<"force-cl-std">,
183+
HelpText<"Force a specific OpenCL C version.">;
167184

168185
def oclfe_32bit : PlainFlag<"m32">;
169186

170187
def oclfe_64bit : PlainFlag<"m64">;
171188

172-
def oclfe_D : PlainJoinedOrSeparate<"D">;
189+
def oclfe_D : PlainJoinedOrSeparate<"D">,
190+
HelpText<"Manually define macros.">;
173191

174192
def oclfe_cl_ext : ClJoined<"ext=">;
175193
// }} Internal options from source translation

0 commit comments

Comments
 (0)