2 # Copyright (C) 2018 Red Hat
3 # Copyright (C) 2014 Intel Corporation
5 # Permission is hereby granted, free of charge, to any person obtaining a
6 # copy of this software and associated documentation files (the "Software"),
7 # to deal in the Software without restriction, including without limitation
8 # the rights to use, copy, modify, merge, publish, distribute, sublicense,
9 # and/or sell copies of the Software, and to permit persons to whom the
10 # Software is furnished to do so, subject to the following conditions:
12 # The above copyright notice and this permission notice (including the next
13 # paragraph) shall be included in all copies or substantial portions of the
16 # THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
17 # IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
18 # FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL
19 # THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
20 # LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
21 # FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
25 # This file defines all the available intrinsics in one place.
27 # The Intrinsic class corresponds one-to-one with nir_intrinsic_info
37 def __init__(self, c_data_type, name):
38 self.c_data_type = c_data_type
41 class Intrinsic(object):
42 """Class that represents all the information about an intrinsic opcode.
43 NOTE: this must be kept in sync with nir_intrinsic_info.
45 def __init__(self, name, src_components, dest_components,
46 indices, flags, sysval, bit_sizes):
49 - name: the intrinsic name
50 - src_components: list of the number of components per src, 0 means
51 vectorized instruction with number of components given in the
52 num_components field in nir_intrinsic_instr.
53 - dest_components: number of destination components, -1 means no
54 dest, 0 means number of components given in num_components field
55 in nir_intrinsic_instr.
56 - indices: list of constant indicies
57 - flags: list of semantic flags
58 - sysval: is this a system-value intrinsic
59 - bit_sizes: allowed dest bit_sizes or the source it must match
61 assert isinstance(name, str)
62 assert isinstance(src_components, list)
64 assert isinstance(src_components[0], int)
65 assert isinstance(dest_components, int)
66 assert isinstance(indices, list)
68 assert isinstance(indices[0], Index)
69 assert isinstance(flags, list)
71 assert isinstance(flags[0], str)
72 assert isinstance(sysval, bool)
73 if isinstance(bit_sizes, list):
74 assert not bit_sizes or isinstance(bit_sizes[0], int)
76 assert isinstance(bit_sizes, tuple)
77 assert bit_sizes[0] == 'src'
78 assert isinstance(bit_sizes[1], int)
81 self.num_srcs = len(src_components)
82 self.src_components = src_components
83 self.has_dest = (dest_components >= 0)
84 self.dest_components = dest_components
85 self.num_indices = len(indices)
86 self.indices = indices
89 self.bit_sizes = bit_sizes if isinstance(bit_sizes, list) else []
90 self.bit_size_src = bit_sizes[1] if isinstance(bit_sizes, tuple) else -1
96 CAN_ELIMINATE = "NIR_INTRINSIC_CAN_ELIMINATE"
97 CAN_REORDER = "NIR_INTRINSIC_CAN_REORDER"
102 def index(c_data_type, name):
103 idx = Index(c_data_type, name)
104 INTR_INDICES.append(idx)
105 globals()[name.upper()] = idx
107 # Defines a new NIR intrinsic. By default, the intrinsic will have no sources
108 # and no destination.
110 # You can set dest_comp=n to enable a destination for the intrinsic, in which
111 # case it will have that many components, or =0 for "as many components as the
112 # NIR destination value."
114 # Set src_comp=n to enable sources for the intruction. It can be an array of
115 # component counts, or (for convenience) a scalar component count if there's
116 # only one source. If a component count is 0, it will be as many components as
117 # the intrinsic has based on the dest_comp.
118 def intrinsic(name, src_comp=[], dest_comp=-1, indices=[],
119 flags=[], sysval=False, bit_sizes=[]):
120 assert name not in INTR_OPCODES
121 INTR_OPCODES[name] = Intrinsic(name, src_comp, dest_comp,
122 indices, flags, sysval, bit_sizes)
128 # Generally instructions that take a offset src argument, can encode
129 # a constant 'base' value which is added to the offset.
132 # For store instructions, a writemask for the store.
133 index("unsigned", "write_mask")
135 # The stream-id for GS emit_vertex/end_primitive intrinsics.
136 index("unsigned", "stream_id")
138 # The clip-plane id for load_user_clip_plane intrinsic.
139 index("unsigned", "ucp_id")
141 # The offset to the start of the NIR_INTRINSIC_RANGE. This is an alternative
142 # to NIR_INTRINSIC_BASE for describing the valid range in intrinsics that don't
143 # have the implicit addition of a base to the offset.
145 # If the [range_base, range] is [0, ~0], then we don't know the possible
146 # range of the access.
147 index("unsigned", "range_base")
149 # The amount of data, starting from BASE or RANGE_BASE, that this
150 # instruction may access. This is used to provide bounds if the offset is
152 index("unsigned", "range")
154 # The Vulkan descriptor set for vulkan_resource_index intrinsic.
155 index("unsigned", "desc_set")
157 # The Vulkan descriptor set binding for vulkan_resource_index intrinsic.
158 index("unsigned", "binding")
161 index("unsigned", "component")
163 # Column index for matrix system values
164 index("unsigned", "column")
166 # Interpolation mode (only meaningful for FS inputs)
167 index("unsigned", "interp_mode")
169 # A binary nir_op to use when performing a reduction or scan operation
170 index("unsigned", "reduction_op")
172 # Cluster size for reduction operations
173 index("unsigned", "cluster_size")
175 # Parameter index for a load_param intrinsic
176 index("unsigned", "param_idx")
178 # Image dimensionality for image intrinsics
179 index("enum glsl_sampler_dim", "image_dim")
181 # Non-zero if we are accessing an array image
182 index("bool", "image_array")
184 # Image format for image intrinsics
185 # Vertex buffer format for load_typed_buffer_amd
186 index("enum pipe_format", "format")
188 # Access qualifiers for image and memory access intrinsics. ACCESS_RESTRICT is
189 # not set at the intrinsic if the NIR was created from SPIR-V.
190 index("enum gl_access_qualifier", "access")
192 # call index for split raytracing shaders
193 index("unsigned", "call_idx")
195 # The stack size increment/decrement for split raytracing shaders
196 index("unsigned", "stack_size")
198 # Alignment for offsets and addresses
200 # These two parameters, specify an alignment in terms of a multiplier and
201 # an offset. The multiplier is always a power of two. The offset or
202 # address parameter X of the intrinsic is guaranteed to satisfy the
205 # (X - align_offset) % align_mul == 0
207 # For constant offset values, align_mul will be NIR_ALIGN_MUL_MAX and the
208 # align_offset will be modulo that.
209 index("unsigned", "align_mul")
210 index("unsigned", "align_offset")
212 # The Vulkan descriptor type for a vulkan_resource_[re]index intrinsic.
213 index("unsigned", "desc_type")
215 # The nir_alu_type of input data to a store or conversion
216 index("nir_alu_type", "src_type")
218 # The nir_alu_type of the data output from a load or conversion
219 index("nir_alu_type", "dest_type")
221 # The swizzle mask for quad_swizzle_amd & masked_swizzle_amd
222 index("unsigned", "swizzle_mask")
224 # Offsets for load_shared2_amd/store_shared2_amd
225 index("uint8_t", "offset0")
226 index("uint8_t", "offset1")
228 # If true, both offsets have an additional stride of 64 dwords (ie. they are multiplied by 256 bytes
229 # in hardware, instead of 4).
230 index("bool", "st64")
232 # When set, range analysis will use it for nir_unsigned_upper_bound
233 index("unsigned", "arg_upper_bound_u32_amd")
235 # Separate source/dest access flags for copies
236 index("enum gl_access_qualifier", "dst_access")
237 index("enum gl_access_qualifier", "src_access")
239 # Driver location of attribute
240 index("unsigned", "driver_location")
242 # Ordering and visibility of a memory operation
243 index("nir_memory_semantics", "memory_semantics")
245 # Modes affected by a memory operation
246 index("nir_variable_mode", "memory_modes")
248 # Scope of a memory operation
249 index("nir_scope", "memory_scope")
251 # Scope of a control barrier
252 index("nir_scope", "execution_scope")
254 # Semantics of an IO instruction
255 index("struct nir_io_semantics", "io_semantics")
257 # Transform feedback info
258 index("struct nir_io_xfb", "io_xfb")
259 index("struct nir_io_xfb", "io_xfb2")
261 # Ray query values accessible from the RayQueryKHR object
262 index("nir_ray_query_value", "ray_query_value")
264 # Rounding mode for conversions
265 index("nir_rounding_mode", "rounding_mode")
267 # Whether or not to saturate in conversions
268 index("unsigned", "saturate")
270 # Whether or not trace_ray_intel is synchronous
271 index("bool", "synchronous")
273 # Value ID to identify SSA value loaded/stored on the stack
274 index("unsigned", "value_id")
276 # Whether to sign-extend offsets in address arithmatic (else zero extend)
277 index("bool", "sign_extend")
279 # Instruction specific flags
280 index("unsigned", "flags")
282 intrinsic("nop", flags=[CAN_ELIMINATE])
284 intrinsic("convert_alu_types", dest_comp=0, src_comp=[0],
285 indices=[SRC_TYPE, DEST_TYPE, ROUNDING_MODE, SATURATE],
286 flags=[CAN_ELIMINATE, CAN_REORDER])
288 intrinsic("load_param", dest_comp=0, indices=[PARAM_IDX], flags=[CAN_ELIMINATE])
290 intrinsic("load_deref", dest_comp=0, src_comp=[-1],
291 indices=[ACCESS], flags=[CAN_ELIMINATE])
292 intrinsic("store_deref", src_comp=[-1, 0], indices=[WRITE_MASK, ACCESS])
293 intrinsic("copy_deref", src_comp=[-1, -1], indices=[DST_ACCESS, SRC_ACCESS])
294 intrinsic("memcpy_deref", src_comp=[-1, -1, 1], indices=[DST_ACCESS, SRC_ACCESS])
296 # Interpolation of input. The interp_deref_at* intrinsics are similar to the
297 # load_var intrinsic acting on a shader input except that they interpolate the
298 # input differently. The at_sample, at_offset and at_vertex intrinsics take an
299 # additional source that is an integer sample id, a vec2 position offset, or a
300 # vertex ID respectively.
302 intrinsic("interp_deref_at_centroid", dest_comp=0, src_comp=[1],
303 flags=[ CAN_ELIMINATE, CAN_REORDER])
304 intrinsic("interp_deref_at_sample", src_comp=[1, 1], dest_comp=0,
305 flags=[CAN_ELIMINATE, CAN_REORDER])
306 intrinsic("interp_deref_at_offset", src_comp=[1, 2], dest_comp=0,
307 flags=[CAN_ELIMINATE, CAN_REORDER])
308 intrinsic("interp_deref_at_vertex", src_comp=[1, 1], dest_comp=0,
309 flags=[CAN_ELIMINATE, CAN_REORDER])
311 # Gets the length of an unsized array at the end of a buffer
312 intrinsic("deref_buffer_array_length", src_comp=[-1], dest_comp=1,
313 indices=[ACCESS], flags=[CAN_ELIMINATE, CAN_REORDER])
315 # Ask the driver for the size of a given SSBO. It takes the buffer index
317 intrinsic("get_ssbo_size", src_comp=[-1], dest_comp=1, bit_sizes=[32],
318 indices=[ACCESS], flags=[CAN_ELIMINATE, CAN_REORDER])
319 intrinsic("get_ubo_size", src_comp=[-1], dest_comp=1,
320 flags=[CAN_ELIMINATE, CAN_REORDER])
322 # Intrinsics which provide a run-time mode-check. Unlike the compile-time
323 # mode checks, a pointer can only have exactly one mode at runtime.
324 intrinsic("deref_mode_is", src_comp=[-1], dest_comp=1,
325 indices=[MEMORY_MODES], flags=[CAN_ELIMINATE, CAN_REORDER])
326 intrinsic("addr_mode_is", src_comp=[-1], dest_comp=1,
327 indices=[MEMORY_MODES], flags=[CAN_ELIMINATE, CAN_REORDER])
329 intrinsic("is_sparse_texels_resident", dest_comp=1, src_comp=[1], bit_sizes=[1,32],
330 flags=[CAN_ELIMINATE, CAN_REORDER])
331 # result code is resident only if both inputs are resident
332 intrinsic("sparse_residency_code_and", dest_comp=1, src_comp=[1, 1], bit_sizes=[32],
333 flags=[CAN_ELIMINATE, CAN_REORDER])
335 # a barrier is an intrinsic with no inputs/outputs but which can't be moved
336 # around/optimized in general
342 # Demote fragment shader invocation to a helper invocation. Any stores to
343 # memory after this instruction are suppressed and the fragment does not write
344 # outputs to the framebuffer. Unlike discard, demote needs to ensure that
345 # derivatives will still work for invocations that were not demoted.
347 # As specified by SPV_EXT_demote_to_helper_invocation.
349 intrinsic("is_helper_invocation", dest_comp=1, flags=[CAN_ELIMINATE])
351 # SpvOpTerminateInvocation from SPIR-V. Essentially a discard "for real".
354 # A workgroup-level control barrier. Any thread which hits this barrier will
355 # pause until all threads within the current workgroup have also hit the
356 # barrier. For compute shaders, the workgroup is defined as the local group.
357 # For tessellation control shaders, the workgroup is defined as the current
358 # patch. This intrinsic does not imply any sort of memory barrier.
359 barrier("control_barrier")
361 # Memory barrier with semantics analogous to the memoryBarrier() GLSL
363 barrier("memory_barrier")
365 # Control/Memory barrier with explicit scope. Follows the semantics of SPIR-V
366 # OpMemoryBarrier and OpControlBarrier, used to implement Vulkan Memory Model.
367 # Storage that the barrier applies is represented using NIR variable modes.
368 # For an OpMemoryBarrier, set EXECUTION_SCOPE to NIR_SCOPE_NONE.
369 intrinsic("scoped_barrier",
370 indices=[EXECUTION_SCOPE, MEMORY_SCOPE, MEMORY_SEMANTICS, MEMORY_MODES])
372 # Shader clock intrinsic with semantics analogous to the clock2x32ARB()
374 # The latter can be used as code motion barrier, which is currently not
376 intrinsic("shader_clock", dest_comp=2, bit_sizes=[32], flags=[CAN_ELIMINATE],
377 indices=[MEMORY_SCOPE])
379 # Shader ballot intrinsics with semantics analogous to the
382 # readInvocationARB()
383 # readFirstInvocationARB()
385 # GLSL functions from ARB_shader_ballot.
386 intrinsic("ballot", src_comp=[1], dest_comp=0, flags=[CAN_ELIMINATE])
387 intrinsic("read_invocation", src_comp=[0, 1], dest_comp=0, bit_sizes=src0, flags=[CAN_ELIMINATE])
388 intrinsic("read_first_invocation", src_comp=[0], dest_comp=0, bit_sizes=src0, flags=[CAN_ELIMINATE])
390 # Returns the value of the first source for the lane where the second source is
391 # true. The second source must be true for exactly one lane.
392 intrinsic("read_invocation_cond_ir3", src_comp=[0, 1], dest_comp=0, flags=[CAN_ELIMINATE])
394 # Additional SPIR-V ballot intrinsics
396 # These correspond to the SPIR-V opcodes
398 # OpGroupNonUniformElect
399 # OpSubgroupFirstInvocationKHR
400 intrinsic("elect", dest_comp=1, flags=[CAN_ELIMINATE])
401 intrinsic("first_invocation", dest_comp=1, bit_sizes=[32], flags=[CAN_ELIMINATE])
402 intrinsic("last_invocation", dest_comp=1, bit_sizes=[32], flags=[CAN_ELIMINATE])
404 # Memory barrier with semantics analogous to the compute shader
405 # groupMemoryBarrier(), memoryBarrierAtomicCounter(), memoryBarrierBuffer(),
406 # memoryBarrierImage() and memoryBarrierShared() GLSL intrinsics.
407 barrier("group_memory_barrier")
408 barrier("memory_barrier_atomic_counter")
409 barrier("memory_barrier_buffer")
410 barrier("memory_barrier_image")
411 barrier("memory_barrier_shared")
412 barrier("begin_invocation_interlock")
413 barrier("end_invocation_interlock")
415 # Memory barrier for synchronizing TCS patch outputs
416 barrier("memory_barrier_tcs_patch")
418 # A conditional discard/demote/terminate, with a single boolean source.
419 intrinsic("discard_if", src_comp=[1])
420 intrinsic("demote_if", src_comp=[1])
421 intrinsic("terminate_if", src_comp=[1])
423 # ARB_shader_group_vote intrinsics
424 intrinsic("vote_any", src_comp=[1], dest_comp=1, flags=[CAN_ELIMINATE])
425 intrinsic("vote_all", src_comp=[1], dest_comp=1, flags=[CAN_ELIMINATE])
426 intrinsic("vote_feq", src_comp=[0], dest_comp=1, flags=[CAN_ELIMINATE])
427 intrinsic("vote_ieq", src_comp=[0], dest_comp=1, flags=[CAN_ELIMINATE])
429 # Ballot ALU operations from SPIR-V.
431 # These operations work like their ALU counterparts except that the operate
432 # on a uvec4 which is treated as a 128bit integer. Also, they are, in
433 # general, free to ignore any bits which are above the subgroup size.
434 intrinsic("ballot_bitfield_extract", src_comp=[4, 1], dest_comp=1, flags=[CAN_ELIMINATE])
435 intrinsic("ballot_bit_count_reduce", src_comp=[4], dest_comp=1, flags=[CAN_ELIMINATE])
436 intrinsic("ballot_bit_count_inclusive", src_comp=[4], dest_comp=1, flags=[CAN_ELIMINATE])
437 intrinsic("ballot_bit_count_exclusive", src_comp=[4], dest_comp=1, flags=[CAN_ELIMINATE])
438 intrinsic("ballot_find_lsb", src_comp=[4], dest_comp=1, flags=[CAN_ELIMINATE])
439 intrinsic("ballot_find_msb", src_comp=[4], dest_comp=1, flags=[CAN_ELIMINATE])
441 # Shuffle operations from SPIR-V.
442 intrinsic("shuffle", src_comp=[0, 1], dest_comp=0, bit_sizes=src0, flags=[CAN_ELIMINATE])
443 intrinsic("shuffle_xor", src_comp=[0, 1], dest_comp=0, bit_sizes=src0, flags=[CAN_ELIMINATE])
444 intrinsic("shuffle_up", src_comp=[0, 1], dest_comp=0, bit_sizes=src0, flags=[CAN_ELIMINATE])
445 intrinsic("shuffle_down", src_comp=[0, 1], dest_comp=0, bit_sizes=src0, flags=[CAN_ELIMINATE])
447 # Quad operations from SPIR-V.
448 intrinsic("quad_broadcast", src_comp=[0, 1], dest_comp=0, flags=[CAN_ELIMINATE])
449 intrinsic("quad_swap_horizontal", src_comp=[0], dest_comp=0, flags=[CAN_ELIMINATE])
450 intrinsic("quad_swap_vertical", src_comp=[0], dest_comp=0, flags=[CAN_ELIMINATE])
451 intrinsic("quad_swap_diagonal", src_comp=[0], dest_comp=0, flags=[CAN_ELIMINATE])
453 # Rotate operation from SPIR-V: SpvOpGroupNonUniformRotateKHR.
454 intrinsic("rotate", src_comp=[0, 1], dest_comp=0, bit_sizes=src0,
455 indices=[EXECUTION_SCOPE, CLUSTER_SIZE], flags=[CAN_ELIMINATE]);
457 intrinsic("reduce", src_comp=[0], dest_comp=0, bit_sizes=src0,
458 indices=[REDUCTION_OP, CLUSTER_SIZE], flags=[CAN_ELIMINATE])
459 intrinsic("inclusive_scan", src_comp=[0], dest_comp=0, bit_sizes=src0,
460 indices=[REDUCTION_OP], flags=[CAN_ELIMINATE])
461 intrinsic("exclusive_scan", src_comp=[0], dest_comp=0, bit_sizes=src0,
462 indices=[REDUCTION_OP], flags=[CAN_ELIMINATE])
464 # AMD shader ballot operations
465 intrinsic("quad_swizzle_amd", src_comp=[0], dest_comp=0, bit_sizes=src0,
466 indices=[SWIZZLE_MASK], flags=[CAN_ELIMINATE])
467 intrinsic("masked_swizzle_amd", src_comp=[0], dest_comp=0, bit_sizes=src0,
468 indices=[SWIZZLE_MASK], flags=[CAN_ELIMINATE])
469 intrinsic("write_invocation_amd", src_comp=[0, 0, 1], dest_comp=0, bit_sizes=src0,
470 flags=[CAN_ELIMINATE])
471 # src = [ mask, addition ]
472 intrinsic("mbcnt_amd", src_comp=[1, 1], dest_comp=1, bit_sizes=[32], flags=[CAN_ELIMINATE])
473 # Compiled to v_permlane16_b32. src = [ value, lanesel_lo, lanesel_hi ]
474 intrinsic("lane_permute_16_amd", src_comp=[1, 1, 1], dest_comp=1, bit_sizes=[32], flags=[CAN_ELIMINATE])
476 # Basic Geometry Shader intrinsics.
478 # emit_vertex implements GLSL's EmitStreamVertex() built-in. It takes a single
479 # index, which is the stream ID to write to.
481 # end_primitive implements GLSL's EndPrimitive() built-in.
482 intrinsic("emit_vertex", indices=[STREAM_ID])
483 intrinsic("end_primitive", indices=[STREAM_ID])
485 # Geometry Shader intrinsics with a vertex count.
487 # Alternatively, drivers may implement these intrinsics, and use
488 # nir_lower_gs_intrinsics() to convert from the basic intrinsics.
490 # These contain two additional unsigned integer sources:
491 # 1. The total number of vertices emitted so far.
492 # 2. The number of vertices emitted for the current primitive
493 # so far if we're counting, otherwise undef.
494 intrinsic("emit_vertex_with_counter", src_comp=[1, 1], indices=[STREAM_ID])
495 intrinsic("end_primitive_with_counter", src_comp=[1, 1], indices=[STREAM_ID])
496 # Contains the final total vertex and primitive counts in the current GS thread.
497 intrinsic("set_vertex_and_primitive_count", src_comp=[1, 1], indices=[STREAM_ID])
499 # Launches mesh shader workgroups from a task shader, with explicit task_payload.
501 # - This is a terminating instruction.
502 # - May only occur in workgroup-uniform control flow.
503 # - Dispatch sizes may be divergent (in which case the values
504 # from the first invocation are used).
505 # Meaning of indices:
506 # - BASE: address of the task_payload variable used.
507 # - RANGE: size of the task_payload variable used.
509 # src[] = {vec(x, y, z)}
510 intrinsic("launch_mesh_workgroups", src_comp=[3], indices=[BASE, RANGE])
512 # Launches mesh shader workgroups from a task shader, with task_payload variable deref.
513 # Same rules as launch_mesh_workgroups apply here as well.
514 # src[] = {vec(x, y, z), payload pointer}
515 intrinsic("launch_mesh_workgroups_with_payload_deref", src_comp=[3, -1], indices=[])
517 # Trace a ray through an acceleration structure
519 # This instruction has a lot of parameters:
520 # 0. Acceleration Structure
525 # 5. Miss shader index
531 intrinsic("trace_ray", src_comp=[-1, 1, 1, 1, 1, 1, 3, 1, 3, 1, -1])
532 # src[] = { hit_t, hit_kind }
533 intrinsic("report_ray_intersection", src_comp=[1, 1], dest_comp=1)
534 intrinsic("ignore_ray_intersection")
535 intrinsic("accept_ray_intersection") # Not in SPIR-V; useful for lowering
536 intrinsic("terminate_ray")
537 # src[] = { sbt_index, payload }
538 intrinsic("execute_callable", src_comp=[1, -1])
540 # Initialize a ray query
543 # 1. Acceleration Structure
550 intrinsic("rq_initialize", src_comp=[-1, -1, 1, 1, 3, 1, 3, 1])
552 intrinsic("rq_terminate", src_comp=[-1])
554 intrinsic("rq_proceed", src_comp=[-1], dest_comp=1)
555 # src[] = { query, hit }
556 intrinsic("rq_generate_intersection", src_comp=[-1, 1])
558 intrinsic("rq_confirm_intersection", src_comp=[-1])
559 # src[] = { query, committed }
560 intrinsic("rq_load", src_comp=[-1, 1], dest_comp=0, indices=[RAY_QUERY_VALUE,COLUMN])
562 # Driver independent raytracing helpers
564 # rt_resume is a helper that that be the first instruction accesing the
565 # stack/scratch in a resume shader for a raytracing pipeline. It includes the
566 # resume index (for nir_lower_shader_calls_internal reasons) and the stack size
567 # of the variables spilled during the call. The stack size can be use to e.g.
568 # adjust a stack pointer.
569 intrinsic("rt_resume", indices=[CALL_IDX, STACK_SIZE])
571 # Lowered version of execute_callabe that includes the index of the resume
572 # shader, and the amount of scratch space needed for this call (.ie. how much
573 # to increase a stack pointer by).
574 # src[] = { sbt_index, payload }
575 intrinsic("rt_execute_callable", src_comp=[1, -1], indices=[CALL_IDX,STACK_SIZE])
577 # Lowered version of trace_ray in a similar vein to rt_execute_callable.
578 # src same as trace_ray
579 intrinsic("rt_trace_ray", src_comp=[-1, 1, 1, 1, 1, 1, 3, 1, 3, 1, -1],
580 indices=[CALL_IDX, STACK_SIZE])
585 # The *_deref variants take an atomic_uint nir_variable, while the other,
586 # lowered, variants take a buffer index and register offset. The buffer index
587 # is always constant, as there's no way to declare an array of atomic counter
590 # The register offset may be non-constant but must by dynamically uniform
591 # ("Atomic counters aggregated into arrays within a shader can only be indexed
592 # with dynamically uniform integral expressions, otherwise results are
594 def atomic(name, flags=[]):
595 intrinsic(name + "_deref", src_comp=[-1], dest_comp=1, flags=flags)
596 intrinsic(name, src_comp=[1], dest_comp=1, indices=[BASE, RANGE_BASE], flags=flags)
599 intrinsic(name + "_deref", src_comp=[-1, 1], dest_comp=1)
600 intrinsic(name, src_comp=[1, 1], dest_comp=1, indices=[BASE, RANGE_BASE])
603 intrinsic(name + "_deref", src_comp=[-1, 1, 1], dest_comp=1)
604 intrinsic(name, src_comp=[1, 1, 1], dest_comp=1, indices=[BASE, RANGE_BASE])
606 atomic("atomic_counter_inc")
607 atomic("atomic_counter_pre_dec")
608 atomic("atomic_counter_post_dec")
609 atomic("atomic_counter_read", flags=[CAN_ELIMINATE])
610 atomic2("atomic_counter_add")
611 atomic2("atomic_counter_min")
612 atomic2("atomic_counter_max")
613 atomic2("atomic_counter_and")
614 atomic2("atomic_counter_or")
615 atomic2("atomic_counter_xor")
616 atomic2("atomic_counter_exchange")
617 atomic3("atomic_counter_comp_swap")
619 # Image load, store and atomic intrinsics.
621 # All image intrinsics come in three versions. One which take an image target
622 # passed as a deref chain as the first source, one which takes an index as the
623 # first source, and one which takes a bindless handle as the first source.
624 # In the first version, the image variable contains the memory and layout
625 # qualifiers that influence the semantics of the intrinsic. In the second and
626 # third, the image format and access qualifiers are provided as constant
627 # indices. Up through GLSL ES 3.10, the image index source may only be a
628 # constant array access. GLSL ES 3.20 and GLSL 4.00 allow dynamically uniform
631 # All image intrinsics take a four-coordinate vector and a sample index as
632 # 2nd and 3rd sources, determining the location within the image that will be
633 # accessed by the intrinsic. Components not applicable to the image target
634 # in use are undefined. Image store takes an additional four-component
635 # argument with the value to be written, and image atomic operations take
636 # either one or two additional scalar arguments with the same meaning as in
637 # the ARB_shader_image_load_store specification.
638 def image(name, src_comp=[], extra_indices=[], **kwargs):
639 intrinsic("image_deref_" + name, src_comp=[-1] + src_comp,
640 indices=[IMAGE_DIM, IMAGE_ARRAY, FORMAT, ACCESS] + extra_indices, **kwargs)
641 intrinsic("image_" + name, src_comp=[1] + src_comp,
642 indices=[IMAGE_DIM, IMAGE_ARRAY, FORMAT, ACCESS, RANGE_BASE] + extra_indices, **kwargs)
643 intrinsic("bindless_image_" + name, src_comp=[-1] + src_comp,
644 indices=[IMAGE_DIM, IMAGE_ARRAY, FORMAT, ACCESS] + extra_indices, **kwargs)
646 image("load", src_comp=[4, 1, 1], extra_indices=[DEST_TYPE], dest_comp=0, flags=[CAN_ELIMINATE])
647 image("sparse_load", src_comp=[4, 1, 1], extra_indices=[DEST_TYPE], dest_comp=0, flags=[CAN_ELIMINATE])
648 image("store", src_comp=[4, 1, 0, 1], extra_indices=[SRC_TYPE])
649 image("atomic_add", src_comp=[4, 1, 1], dest_comp=1)
650 image("atomic_imin", src_comp=[4, 1, 1], dest_comp=1)
651 image("atomic_umin", src_comp=[4, 1, 1], dest_comp=1)
652 image("atomic_imax", src_comp=[4, 1, 1], dest_comp=1)
653 image("atomic_umax", src_comp=[4, 1, 1], dest_comp=1)
654 image("atomic_and", src_comp=[4, 1, 1], dest_comp=1)
655 image("atomic_or", src_comp=[4, 1, 1], dest_comp=1)
656 image("atomic_xor", src_comp=[4, 1, 1], dest_comp=1)
657 image("atomic_exchange", src_comp=[4, 1, 1], dest_comp=1)
658 image("atomic_comp_swap", src_comp=[4, 1, 1, 1], dest_comp=1)
659 image("atomic_fadd", src_comp=[4, 1, 1], dest_comp=1)
660 image("atomic_fmin", src_comp=[4, 1, 1], dest_comp=1)
661 image("atomic_fmax", src_comp=[4, 1, 1], dest_comp=1)
662 image("size", dest_comp=0, src_comp=[1], flags=[CAN_ELIMINATE, CAN_REORDER])
663 image("samples", dest_comp=1, flags=[CAN_ELIMINATE, CAN_REORDER])
664 image("atomic_inc_wrap", src_comp=[4, 1, 1], dest_comp=1)
665 image("atomic_dec_wrap", src_comp=[4, 1, 1], dest_comp=1)
666 # This returns true if all samples within the pixel have equal color values.
667 image("samples_identical", dest_comp=1, src_comp=[4], flags=[CAN_ELIMINATE])
668 # Non-uniform access is not lowered for image_descriptor_amd.
669 # dest_comp can be either 4 (buffer) or 8 (image).
670 image("descriptor_amd", dest_comp=0, src_comp=[], flags=[CAN_ELIMINATE, CAN_REORDER])
671 # CL-specific format queries
672 image("format", dest_comp=1, flags=[CAN_ELIMINATE, CAN_REORDER])
673 image("order", dest_comp=1, flags=[CAN_ELIMINATE, CAN_REORDER])
674 # Multisample fragment mask load
675 # src_comp[0] is same as image load src_comp[0]
676 image("fragment_mask_load_amd", src_comp=[4], dest_comp=1, bit_sizes=[32], flags=[CAN_ELIMINATE, CAN_REORDER])
678 # Vulkan descriptor set intrinsics
680 # The Vulkan API uses a different binding model from GL. In the Vulkan
681 # API, all external resources are represented by a tuple:
683 # (descriptor set, binding, array index)
685 # where the array index is the only thing allowed to be indirect. The
686 # vulkan_surface_index intrinsic takes the descriptor set and binding as
687 # its first two indices and the array index as its source. The third
688 # index is a nir_variable_mode in case that's useful to the backend.
690 # The intended usage is that the shader will call vulkan_surface_index to
691 # get an index and then pass that as the buffer index ubo/ssbo calls.
693 # The vulkan_resource_reindex intrinsic takes a resource index in src0
694 # (the result of a vulkan_resource_index or vulkan_resource_reindex) which
695 # corresponds to the tuple (set, binding, index) and computes an index
696 # corresponding to tuple (set, binding, idx + src1).
697 intrinsic("vulkan_resource_index", src_comp=[1], dest_comp=0,
698 indices=[DESC_SET, BINDING, DESC_TYPE],
699 flags=[CAN_ELIMINATE, CAN_REORDER])
700 intrinsic("vulkan_resource_reindex", src_comp=[0, 1], dest_comp=0,
701 indices=[DESC_TYPE], flags=[CAN_ELIMINATE, CAN_REORDER])
702 intrinsic("load_vulkan_descriptor", src_comp=[-1], dest_comp=0,
703 indices=[DESC_TYPE], flags=[CAN_ELIMINATE, CAN_REORDER])
707 # All of these atomic memory operations read a value from memory, compute a new
708 # value using one of the operations below, write the new value to memory, and
709 # return the original value read.
711 # All variable operations take 2 sources except CompSwap that takes 3. These
714 # 0: A deref to the memory on which to perform the atomic
715 # 1: The data parameter to the atomic function (i.e. the value to add
716 # in shared_atomic_add, etc).
717 # 2: For CompSwap only: the second data parameter.
719 # All SSBO operations take 3 sources except CompSwap that takes 4. These
722 # 0: The SSBO buffer index (dynamically uniform in GLSL, possibly non-uniform
723 # with VK_EXT_descriptor_indexing).
724 # 1: The offset into the SSBO buffer of the variable that the atomic
725 # operation will operate on.
726 # 2: The data parameter to the atomic function (i.e. the value to add
727 # in ssbo_atomic_add, etc).
728 # 3: For CompSwap only: the second data parameter.
730 # All shared (and task payload) variable operations take 2 sources
731 # except CompSwap that takes 3.
732 # These sources represent:
734 # 0: The offset into the shared variable storage region that the atomic
735 # operation will operate on.
736 # 1: The data parameter to the atomic function (i.e. the value to add
737 # in shared_atomic_add, etc).
738 # 2: For CompSwap only: the second data parameter.
740 # All global operations take 2 sources except CompSwap that takes 3. These
743 # 0: The memory address that the atomic operation will operate on.
744 # 1: The data parameter to the atomic function (i.e. the value to add
745 # in shared_atomic_add, etc).
746 # 2: For CompSwap only: the second data parameter.
748 # The 2x32 global variants use a vec2 for the memory address where component X
749 # has the low 32-bit and component Y has the high 32-bit.
751 # IR3 global operations take 32b vec2 as memory address. IR3 doesn't support
754 def memory_atomic_data1(name):
755 intrinsic("deref_atomic_" + name, src_comp=[-1, 1], dest_comp=1, indices=[ACCESS])
756 intrinsic("ssbo_atomic_" + name, src_comp=[-1, 1, 1], dest_comp=1, indices=[ACCESS])
757 intrinsic("shared_atomic_" + name, src_comp=[1, 1], dest_comp=1, indices=[BASE])
758 intrinsic("task_payload_atomic_" + name, src_comp=[1, 1], dest_comp=1, indices=[BASE])
759 intrinsic("global_atomic_" + name, src_comp=[1, 1], dest_comp=1, indices=[])
760 intrinsic("global_atomic_" + name + "_2x32", src_comp=[2, 1], dest_comp=1, indices=[])
761 intrinsic("global_atomic_" + name + "_amd", src_comp=[1, 1, 1], dest_comp=1, indices=[BASE])
762 if not name.startswith('f'):
763 intrinsic("global_atomic_" + name + "_ir3", src_comp=[2, 1], dest_comp=1, indices=[BASE])
765 def memory_atomic_data2(name):
766 intrinsic("deref_atomic_" + name, src_comp=[-1, 1, 1], dest_comp=1, indices=[ACCESS])
767 intrinsic("ssbo_atomic_" + name, src_comp=[-1, 1, 1, 1], dest_comp=1, indices=[ACCESS])
768 intrinsic("shared_atomic_" + name, src_comp=[1, 1, 1], dest_comp=1, indices=[BASE])
769 intrinsic("task_payload_atomic_" + name, src_comp=[1, 1, 1], dest_comp=1, indices=[BASE])
770 intrinsic("global_atomic_" + name, src_comp=[1, 1, 1], dest_comp=1, indices=[])
771 intrinsic("global_atomic_" + name + "_2x32", src_comp=[2, 1, 1], dest_comp=1, indices=[])
772 intrinsic("global_atomic_" + name + "_amd", src_comp=[1, 1, 1, 1], dest_comp=1, indices=[BASE])
773 if not name.startswith('f'):
774 intrinsic("global_atomic_" + name + "_ir3", src_comp=[2, 1, 1], dest_comp=1, indices=[BASE])
776 memory_atomic_data1("add")
777 memory_atomic_data1("imin")
778 memory_atomic_data1("umin")
779 memory_atomic_data1("imax")
780 memory_atomic_data1("umax")
781 memory_atomic_data1("and")
782 memory_atomic_data1("or")
783 memory_atomic_data1("xor")
784 memory_atomic_data1("exchange")
785 memory_atomic_data1("fadd")
786 memory_atomic_data1("fmin")
787 memory_atomic_data1("fmax")
788 memory_atomic_data2("comp_swap")
789 memory_atomic_data2("fcomp_swap")
791 def system_value(name, dest_comp, indices=[], bit_sizes=[32]):
792 intrinsic("load_" + name, [], dest_comp, indices,
793 flags=[CAN_ELIMINATE, CAN_REORDER], sysval=True,
796 system_value("frag_coord", 4)
797 system_value("point_coord", 2)
798 system_value("line_coord", 1)
799 system_value("front_face", 1, bit_sizes=[1, 32])
800 system_value("vertex_id", 1)
801 system_value("vertex_id_zero_base", 1)
802 system_value("first_vertex", 1)
803 system_value("is_indexed_draw", 1)
804 system_value("base_vertex", 1)
805 system_value("instance_id", 1)
806 system_value("base_instance", 1)
807 system_value("draw_id", 1)
808 system_value("sample_id", 1)
809 # sample_id_no_per_sample is like sample_id but does not imply per-
810 # sample shading. See the lower_helper_invocation option.
811 system_value("sample_id_no_per_sample", 1)
812 system_value("sample_pos", 2)
813 # sample_pos_or_center is like sample_pos but does not imply per-sample
814 # shading. When per-sample dispatch is not enabled, it returns (0.5, 0.5).
815 system_value("sample_pos_or_center", 2)
816 system_value("sample_mask_in", 1)
817 system_value("primitive_id", 1)
818 system_value("invocation_id", 1)
819 system_value("tess_coord", 3)
820 system_value("tess_level_outer", 4)
821 system_value("tess_level_inner", 2)
822 system_value("tess_level_outer_default", 4)
823 system_value("tess_level_inner_default", 2)
824 system_value("patch_vertices_in", 1)
825 system_value("local_invocation_id", 3)
826 system_value("local_invocation_index", 1)
827 # zero_base indicates it starts from 0 for the current dispatch
828 # non-zero_base indicates the base is included
829 system_value("workgroup_id", 3, bit_sizes=[32, 64])
830 system_value("workgroup_id_zero_base", 3)
831 # The workgroup_index is intended for situations when a 3 dimensional
832 # workgroup_id is not available on the HW, but a 1 dimensional index is.
833 system_value("workgroup_index", 1)
834 system_value("base_workgroup_id", 3, bit_sizes=[32, 64])
835 system_value("user_clip_plane", 4, indices=[UCP_ID])
836 system_value("num_workgroups", 3, bit_sizes=[32, 64])
837 system_value("num_vertices", 1)
838 system_value("helper_invocation", 1, bit_sizes=[1, 32])
839 system_value("layer_id", 1)
840 system_value("view_index", 1)
841 system_value("subgroup_size", 1)
842 system_value("subgroup_invocation", 1)
843 system_value("subgroup_eq_mask", 0, bit_sizes=[32, 64])
844 system_value("subgroup_ge_mask", 0, bit_sizes=[32, 64])
845 system_value("subgroup_gt_mask", 0, bit_sizes=[32, 64])
846 system_value("subgroup_le_mask", 0, bit_sizes=[32, 64])
847 system_value("subgroup_lt_mask", 0, bit_sizes=[32, 64])
848 system_value("num_subgroups", 1)
849 system_value("subgroup_id", 1)
850 system_value("workgroup_size", 3)
851 # note: the definition of global_invocation_id_zero_base is based on
852 # (workgroup_id * workgroup_size) + local_invocation_id.
853 # it is *not* based on workgroup_id_zero_base, meaning the work group
854 # base is already accounted for, and the global base is additive on top of that
855 system_value("global_invocation_id", 3, bit_sizes=[32, 64])
856 system_value("global_invocation_id_zero_base", 3, bit_sizes=[32, 64])
857 system_value("base_global_invocation_id", 3, bit_sizes=[32, 64])
858 system_value("global_invocation_index", 1, bit_sizes=[32, 64])
859 system_value("work_dim", 1)
860 system_value("line_width", 1)
861 system_value("aa_line_width", 1)
862 # BASE=0 for global/shader, BASE=1 for local/function
863 system_value("scratch_base_ptr", 0, bit_sizes=[32,64], indices=[BASE])
864 system_value("constant_base_ptr", 0, bit_sizes=[32,64])
865 system_value("shared_base_ptr", 0, bit_sizes=[32,64])
866 system_value("global_base_ptr", 0, bit_sizes=[32,64])
867 # Address and size of a transform feedback buffer, indexed by BASE
868 system_value("xfb_address", 1, bit_sizes=[32,64], indices=[BASE])
869 system_value("xfb_size", 1, bit_sizes=[32], indices=[BASE])
871 # Address of the associated index buffer in a transform feedback program for an
872 # indexed draw. This will be used so transform feedback can pull the gl_VertexID
873 # from the index buffer.
874 system_value("xfb_index_buffer", 1, bit_sizes=[32,64])
876 system_value("frag_size", 2)
877 system_value("frag_invocation_count", 1)
879 # System values for ray tracing.
880 system_value("ray_launch_id", 3)
881 system_value("ray_launch_size", 3)
882 system_value("ray_world_origin", 3)
883 system_value("ray_world_direction", 3)
884 system_value("ray_object_origin", 3)
885 system_value("ray_object_direction", 3)
886 system_value("ray_t_min", 1)
887 system_value("ray_t_max", 1)
888 system_value("ray_object_to_world", 3, indices=[COLUMN])
889 system_value("ray_world_to_object", 3, indices=[COLUMN])
890 system_value("ray_hit_kind", 1)
891 system_value("ray_flags", 1)
892 system_value("ray_geometry_index", 1)
893 system_value("ray_instance_custom_index", 1)
894 system_value("shader_record_ptr", 1, bit_sizes=[64])
895 system_value("cull_mask", 1)
897 # Driver-specific viewport scale/offset parameters.
899 # VC4 and V3D need to emit a scaled version of the position in the vertex
900 # shaders for binning, and having system values lets us move the math for that
903 # Panfrost needs to implement all coordinate transformation in the
904 # vertex shader; system values allow us to share this routine in NIR.
905 system_value("viewport_x_scale", 1)
906 system_value("viewport_y_scale", 1)
907 system_value("viewport_z_scale", 1)
908 system_value("viewport_x_offset", 1)
909 system_value("viewport_y_offset", 1)
910 system_value("viewport_z_offset", 1)
911 system_value("viewport_scale", 3)
912 system_value("viewport_offset", 3)
913 # Pack xy scale and offset into a vec4 load (used by AMD NGG primitive culling)
914 system_value("viewport_xy_scale_and_offset", 4)
916 # Blend constant color values. Float values are clamped. Vectored versions are
917 # provided as well for driver convenience
919 system_value("blend_const_color_r_float", 1)
920 system_value("blend_const_color_g_float", 1)
921 system_value("blend_const_color_b_float", 1)
922 system_value("blend_const_color_a_float", 1)
923 system_value("blend_const_color_rgba", 4)
924 system_value("blend_const_color_rgba8888_unorm", 1)
925 system_value("blend_const_color_aaaa8888_unorm", 1)
927 # System values for gl_Color, for radeonsi which interpolates these in the
928 # shader prolog to handle two-sided color without recompiles and therefore
929 # doesn't handle these in the main shader part like normal varyings.
930 system_value("color0", 4)
931 system_value("color1", 4)
933 # System value for internal compute shaders in radeonsi.
934 system_value("user_data_amd", 4)
936 # Barycentric coordinate intrinsics.
938 # These set up the barycentric coordinates for a particular interpolation.
939 # The first four are for the simple cases: pixel, centroid, per-sample
940 # (at gl_SampleID), or pull model (1/W, 1/I, 1/J) at the pixel center. The next
941 # two handle interpolating at a specified sample location, or interpolating
942 # with a vec2 offset,
944 # The interp_mode index should be either the INTERP_MODE_SMOOTH or
945 # INTERP_MODE_NOPERSPECTIVE enum values.
947 # The vec2 value produced by these intrinsics is intended for use as the
948 # barycoord source of a load_interpolated_input intrinsic.
950 def barycentric(name, dst_comp, src_comp=[]):
951 intrinsic("load_barycentric_" + name, src_comp=src_comp, dest_comp=dst_comp,
952 indices=[INTERP_MODE], flags=[CAN_ELIMINATE, CAN_REORDER])
955 barycentric("pixel", 2)
956 barycentric("centroid", 2)
957 barycentric("sample", 2)
958 barycentric("model", 3)
959 # src[] = { sample_id }.
960 barycentric("at_sample", 2, [1])
961 # src[] = { offset.xy }.
962 barycentric("at_offset", 2, [2])
964 # Load sample position:
966 # Takes a sample # and returns a sample position. Used for lowering
967 # interpolateAtSample() to interpolateAtOffset()
968 intrinsic("load_sample_pos_from_id", src_comp=[1], dest_comp=2,
969 flags=[CAN_ELIMINATE, CAN_REORDER])
971 intrinsic("load_persp_center_rhw_ir3", dest_comp=1,
972 flags=[CAN_ELIMINATE, CAN_REORDER])
974 # Load texture scaling values:
976 # Takes a sampler # and returns 1/size values for multiplying to normalize
977 # texture coordinates. Used for lowering rect textures.
978 intrinsic("load_texture_rect_scaling", src_comp=[1], dest_comp=2,
979 flags=[CAN_ELIMINATE, CAN_REORDER])
981 # Fragment shader input interpolation delta intrinsic.
983 # For hw where fragment shader input interpolation is handled in shader, the
984 # load_fs_input_interp deltas intrinsics can be used to load the input deltas
985 # used for interpolation as follows:
987 # vec3 iid = load_fs_input_interp_deltas(varying_slot)
988 # vec2 bary = load_barycentric_*(...)
989 # float result = iid.x + iid.y * bary.y + iid.z * bary.x
991 intrinsic("load_fs_input_interp_deltas", src_comp=[1], dest_comp=3,
992 indices=[BASE, COMPONENT, IO_SEMANTICS], flags=[CAN_ELIMINATE, CAN_REORDER])
994 # Load operations pull data from some piece of GPU memory. All load
995 # operations operate in terms of offsets into some piece of theoretical
996 # memory. Loads from externally visible memory (UBO and SSBO) simply take a
997 # byte offset as a source. Loads from opaque memory (uniforms, inputs, etc.)
998 # take a base+offset pair where the nir_intrinsic_base() gives the location
999 # of the start of the variable being loaded and and the offset source is a
1000 # offset into that variable.
1002 # Uniform load operations have a nir_intrinsic_range() index that specifies the
1003 # range (starting at base) of the data from which we are loading. If
1004 # range == 0, then the range is unknown.
1006 # UBO load operations have a nir_intrinsic_range_base() and
1007 # nir_intrinsic_range() that specify the byte range [range_base,
1008 # range_base+range] of the UBO that the src offset access must lie within.
1010 # Some load operations such as UBO/SSBO load and per_vertex loads take an
1011 # additional source to specify which UBO/SSBO/vertex to load from.
1013 # The exact address type depends on the lowering pass that generates the
1014 # load/store intrinsics. Typically, this is vec4 units for things such as
1015 # varying slots and float units for fragment shader inputs. UBO and SSBO
1016 # offsets are always in bytes.
1018 def load(name, src_comp, indices=[], flags=[]):
1019 intrinsic("load_" + name, src_comp, dest_comp=0, indices=indices,
1022 # src[] = { offset }.
1023 load("uniform", [1], [BASE, RANGE, DEST_TYPE], [CAN_ELIMINATE, CAN_REORDER])
1024 # src[] = { buffer_index, offset }.
1025 load("ubo", [-1, 1], [ACCESS, ALIGN_MUL, ALIGN_OFFSET, RANGE_BASE, RANGE], flags=[CAN_ELIMINATE, CAN_REORDER])
1026 # src[] = { buffer_index, offset in vec4 units }. base is also in vec4 units.
1027 load("ubo_vec4", [-1, 1], [ACCESS, BASE, COMPONENT], flags=[CAN_ELIMINATE, CAN_REORDER])
1028 # src[] = { offset }.
1029 load("input", [1], [BASE, COMPONENT, DEST_TYPE, IO_SEMANTICS], [CAN_ELIMINATE, CAN_REORDER])
1030 # src[] = { vertex_id, offset }.
1031 load("input_vertex", [1, 1], [BASE, COMPONENT, DEST_TYPE, IO_SEMANTICS], [CAN_ELIMINATE, CAN_REORDER])
1032 # src[] = { vertex, offset }.
1033 load("per_vertex_input", [1, 1], [BASE, COMPONENT, DEST_TYPE, IO_SEMANTICS], [CAN_ELIMINATE, CAN_REORDER])
1034 # src[] = { barycoord, offset }.
1035 load("interpolated_input", [2, 1], [BASE, COMPONENT, DEST_TYPE, IO_SEMANTICS], [CAN_ELIMINATE, CAN_REORDER])
1037 # src[] = { buffer_index, offset }.
1038 load("ssbo", [-1, 1], [ACCESS, ALIGN_MUL, ALIGN_OFFSET], [CAN_ELIMINATE])
1039 # src[] = { buffer_index }
1040 load("ssbo_address", [1], [], [CAN_ELIMINATE, CAN_REORDER])
1041 # src[] = { offset }.
1042 load("output", [1], [BASE, COMPONENT, DEST_TYPE, IO_SEMANTICS], flags=[CAN_ELIMINATE])
1043 # src[] = { vertex, offset }.
1044 load("per_vertex_output", [1, 1], [BASE, COMPONENT, DEST_TYPE, IO_SEMANTICS], [CAN_ELIMINATE])
1045 # src[] = { primitive, offset }.
1046 load("per_primitive_output", [1, 1], [BASE, COMPONENT, DEST_TYPE, IO_SEMANTICS], [CAN_ELIMINATE])
1047 # src[] = { offset }.
1048 load("shared", [1], [BASE, ALIGN_MUL, ALIGN_OFFSET], [CAN_ELIMINATE])
1049 # src[] = { offset }.
1050 load("task_payload", [1], [BASE, ALIGN_MUL, ALIGN_OFFSET], [CAN_ELIMINATE])
1051 # src[] = { offset }.
1052 load("push_constant", [1], [BASE, RANGE, ALIGN_MUL, ALIGN_OFFSET], [CAN_ELIMINATE, CAN_REORDER])
1053 # src[] = { offset }.
1054 load("constant", [1], [BASE, RANGE, ALIGN_MUL, ALIGN_OFFSET],
1055 [CAN_ELIMINATE, CAN_REORDER])
1056 # src[] = { address }.
1057 load("global", [1], [ACCESS, ALIGN_MUL, ALIGN_OFFSET], [CAN_ELIMINATE])
1058 # src[] = { address }.
1059 load("global_2x32", [2], [ACCESS, ALIGN_MUL, ALIGN_OFFSET], [CAN_ELIMINATE])
1060 # src[] = { address }.
1061 load("global_constant", [1], [ACCESS, ALIGN_MUL, ALIGN_OFFSET],
1062 [CAN_ELIMINATE, CAN_REORDER])
1063 # src[] = { base_address, offset }.
1064 load("global_constant_offset", [1, 1], [ACCESS, ALIGN_MUL, ALIGN_OFFSET],
1065 [CAN_ELIMINATE, CAN_REORDER])
1066 # src[] = { base_address, offset, bound }.
1067 load("global_constant_bounded", [1, 1, 1], [ACCESS, ALIGN_MUL, ALIGN_OFFSET],
1068 [CAN_ELIMINATE, CAN_REORDER])
1069 # src[] = { address }.
1070 load("kernel_input", [1], [BASE, RANGE, ALIGN_MUL, ALIGN_OFFSET], [CAN_ELIMINATE, CAN_REORDER])
1071 # src[] = { offset }.
1072 load("scratch", [1], [ALIGN_MUL, ALIGN_OFFSET], [CAN_ELIMINATE])
1074 # Stores work the same way as loads, except now the first source is the value
1075 # to store and the second (and possibly third) source specify where to store
1076 # the value. SSBO and shared memory stores also have a
1077 # nir_intrinsic_write_mask()
1079 def store(name, srcs, indices=[], flags=[]):
1080 intrinsic("store_" + name, [0] + srcs, indices=indices, flags=flags)
1082 # src[] = { value, offset }.
1083 store("output", [1], [BASE, WRITE_MASK, COMPONENT, SRC_TYPE, IO_SEMANTICS, IO_XFB, IO_XFB2])
1084 # src[] = { value, vertex, offset }.
1085 store("per_vertex_output", [1, 1], [BASE, WRITE_MASK, COMPONENT, SRC_TYPE, IO_SEMANTICS])
1086 # src[] = { value, primitive, offset }.
1087 store("per_primitive_output", [1, 1], [BASE, WRITE_MASK, COMPONENT, SRC_TYPE, IO_SEMANTICS])
1088 # src[] = { value, block_index, offset }
1089 store("ssbo", [-1, 1], [WRITE_MASK, ACCESS, ALIGN_MUL, ALIGN_OFFSET])
1090 # src[] = { value, offset }.
1091 store("shared", [1], [BASE, WRITE_MASK, ALIGN_MUL, ALIGN_OFFSET])
1092 # src[] = { value, offset }.
1093 store("task_payload", [1], [BASE, WRITE_MASK, ALIGN_MUL, ALIGN_OFFSET])
1094 # src[] = { value, address }.
1095 store("global", [1], [WRITE_MASK, ACCESS, ALIGN_MUL, ALIGN_OFFSET])
1096 # src[] = { value, address }.
1097 store("global_2x32", [2], [WRITE_MASK, ACCESS, ALIGN_MUL, ALIGN_OFFSET])
1098 # src[] = { value, offset }.
1099 store("scratch", [1], [ALIGN_MUL, ALIGN_OFFSET, WRITE_MASK])
1101 # Intrinsic to load/store from the call stack.
1102 # BASE is the offset relative to the current position of the stack
1104 intrinsic("load_stack", [], dest_comp=0,
1105 indices=[BASE, ALIGN_MUL, ALIGN_OFFSET, CALL_IDX, VALUE_ID],
1106 flags=[CAN_ELIMINATE])
1107 # src[] = { value }.
1108 intrinsic("store_stack", [0],
1109 indices=[BASE, ALIGN_MUL, ALIGN_OFFSET, WRITE_MASK, CALL_IDX, VALUE_ID])
1112 # A bit field to implement SPIRV FragmentShadingRateKHR
1113 # bit | name | description
1114 # 0 | Vertical2Pixels | Fragment invocation covers 2 pixels vertically
1115 # 1 | Vertical4Pixels | Fragment invocation covers 4 pixels vertically
1116 # 2 | Horizontal2Pixels | Fragment invocation covers 2 pixels horizontally
1117 # 3 | Horizontal4Pixels | Fragment invocation covers 4 pixels horizontally
1118 intrinsic("load_frag_shading_rate", dest_comp=1, bit_sizes=[32],
1119 flags=[CAN_ELIMINATE, CAN_REORDER])
1121 # Whether the rasterized fragment is fully covered by the generating primitive.
1122 system_value("fully_covered", dest_comp=1, bit_sizes=[1])
1124 # OpenCL printf instruction
1125 # First source is a deref to the format string
1126 # Second source is a deref to a struct containing the args
1127 # Dest is success or failure
1128 intrinsic("printf", src_comp=[1, 1], dest_comp=1, bit_sizes=[32])
1129 # Since most drivers will want to lower to just dumping args
1130 # in a buffer, nir_lower_printf will do that, but requires
1131 # the driver to at least provide a base location
1132 system_value("printf_buffer_address", 1, bit_sizes=[32,64])
1134 # Mesh shading MultiView intrinsics
1135 system_value("mesh_view_count", 1)
1136 load("mesh_view_indices", [1], [BASE, RANGE], [CAN_ELIMINATE, CAN_REORDER])
1138 # Used to pass values from the preamble to the main shader.
1139 # This should use something similar to Vulkan push constants and load_preamble
1140 # should be relatively cheap.
1141 # For now we only support accesses with a constant offset.
1142 load("preamble", [], indices=[BASE], flags=[CAN_ELIMINATE, CAN_REORDER])
1143 store("preamble", [], indices=[BASE])
1145 # A 32 bits bitfield storing 1 in bits corresponding to varyings
1146 # that have the flat interpolation specifier in the fragment shader
1148 system_value("flat_mask", 1)
1150 # Whether provoking vertex mode is last
1151 system_value("provoking_last", 1)
1153 # IR3-specific version of most SSBO intrinsics. The only different
1154 # compare to the originals is that they add an extra source to hold
1155 # the dword-offset, which is needed by the backend code apart from
1156 # the byte-offset already provided by NIR in one of the sources.
1158 # NIR lowering pass 'ir3_nir_lower_io_offset' will replace the
1159 # original SSBO intrinsics by these, placing the computed
1160 # dword-offset always in the last source.
1162 # The float versions are not handled because those are not supported
1164 store("ssbo_ir3", [1, 1, 1],
1165 indices=[WRITE_MASK, ACCESS, ALIGN_MUL, ALIGN_OFFSET])
1166 load("ssbo_ir3", [1, 1, 1],
1167 indices=[ACCESS, ALIGN_MUL, ALIGN_OFFSET], flags=[CAN_ELIMINATE])
1168 intrinsic("ssbo_atomic_add_ir3", src_comp=[1, 1, 1, 1], dest_comp=1, indices=[ACCESS])
1169 intrinsic("ssbo_atomic_imin_ir3", src_comp=[1, 1, 1, 1], dest_comp=1, indices=[ACCESS])
1170 intrinsic("ssbo_atomic_umin_ir3", src_comp=[1, 1, 1, 1], dest_comp=1, indices=[ACCESS])
1171 intrinsic("ssbo_atomic_imax_ir3", src_comp=[1, 1, 1, 1], dest_comp=1, indices=[ACCESS])
1172 intrinsic("ssbo_atomic_umax_ir3", src_comp=[1, 1, 1, 1], dest_comp=1, indices=[ACCESS])
1173 intrinsic("ssbo_atomic_and_ir3", src_comp=[1, 1, 1, 1], dest_comp=1, indices=[ACCESS])
1174 intrinsic("ssbo_atomic_or_ir3", src_comp=[1, 1, 1, 1], dest_comp=1, indices=[ACCESS])
1175 intrinsic("ssbo_atomic_xor_ir3", src_comp=[1, 1, 1, 1], dest_comp=1, indices=[ACCESS])
1176 intrinsic("ssbo_atomic_exchange_ir3", src_comp=[1, 1, 1, 1], dest_comp=1, indices=[ACCESS])
1177 intrinsic("ssbo_atomic_comp_swap_ir3", src_comp=[1, 1, 1, 1, 1], dest_comp=1, indices=[ACCESS])
1179 # System values for freedreno geometry shaders.
1180 system_value("vs_primitive_stride_ir3", 1)
1181 system_value("vs_vertex_stride_ir3", 1)
1182 system_value("gs_header_ir3", 1)
1183 system_value("primitive_location_ir3", 1, indices=[DRIVER_LOCATION])
1185 # System values for freedreno tessellation shaders.
1186 system_value("hs_patch_stride_ir3", 1)
1187 system_value("tess_factor_base_ir3", 2)
1188 system_value("tess_param_base_ir3", 2)
1189 system_value("tcs_header_ir3", 1)
1190 system_value("rel_patch_id_ir3", 1)
1192 # System values for freedreno compute shaders.
1193 system_value("subgroup_id_shift_ir3", 1)
1195 # IR3-specific intrinsics for tessellation control shaders. cond_end_ir3 end
1196 # the shader when src0 is false and is used to narrow down the TCS shader to
1197 # just thread 0 before writing out tessellation levels.
1198 intrinsic("cond_end_ir3", src_comp=[1])
1199 # end_patch_ir3 is used just before thread 0 exist the TCS and presumably
1200 # signals the TE that the patch is complete and can be tessellated.
1201 intrinsic("end_patch_ir3")
1203 # IR3-specific load/store intrinsics. These access a buffer used to pass data
1204 # between geometry stages - perhaps it's explicit access to the vertex cache.
1206 # src[] = { value, offset }.
1207 store("shared_ir3", [1], [BASE, ALIGN_MUL, ALIGN_OFFSET])
1208 # src[] = { offset }.
1209 load("shared_ir3", [1], [BASE, ALIGN_MUL, ALIGN_OFFSET], [CAN_ELIMINATE])
1211 # IR3-specific load/store global intrinsics. They take a 64-bit base address
1212 # and a 32-bit offset. The hardware will add the base and the offset, which
1213 # saves us from doing 64-bit math on the base address.
1215 # src[] = { value, address(vec2 of hi+lo uint32_t), offset }.
1216 # const_index[] = { write_mask, align_mul, align_offset }
1217 store("global_ir3", [2, 1], indices=[ACCESS, ALIGN_MUL, ALIGN_OFFSET])
1218 # src[] = { address(vec2 of hi+lo uint32_t), offset }.
1219 # const_index[] = { access, align_mul, align_offset }
1220 load("global_ir3", [2, 1], indices=[ACCESS, ALIGN_MUL, ALIGN_OFFSET], flags=[CAN_ELIMINATE])
1222 # IR3-specific bindless handle specifier. Similar to vulkan_resource_index, but
1223 # without the binding because the hardware expects a single flattened index
1224 # rather than a (binding, index) pair. We may also want to use this with GL.
1225 # Note that this doesn't actually turn into a HW instruction.
1226 intrinsic("bindless_resource_ir3", [1], dest_comp=1, indices=[DESC_SET], flags=[CAN_ELIMINATE, CAN_REORDER])
1228 # IR3-specific intrinsics for shader preamble. These are meant to be used like
1231 # if (preamble_start()) {
1232 # if (subgroupElect()) {
1241 intrinsic("preamble_start_ir3", [], dest_comp=1, flags=[CAN_ELIMINATE, CAN_REORDER])
1243 barrier("preamble_end_ir3")
1245 # IR3-specific intrinsic for stc. Should be used in the shader preamble.
1246 store("uniform_ir3", [], indices=[BASE])
1248 # IR3-specific intrinsic for ldc.k. Copies UBO to constant file.
1249 # base is the const file base in components, range is the amount to copy in
1251 intrinsic("copy_ubo_to_uniform_ir3", [1, 1], indices=[BASE, RANGE])
1253 # DXIL specific intrinsics
1254 # src[] = { value, mask, index, offset }.
1255 intrinsic("store_ssbo_masked_dxil", [1, 1, 1, 1])
1256 # src[] = { value, index }.
1257 intrinsic("store_shared_dxil", [1, 1])
1258 # src[] = { value, mask, index }.
1259 intrinsic("store_shared_masked_dxil", [1, 1, 1])
1260 # src[] = { value, index }.
1261 intrinsic("store_scratch_dxil", [1, 1])
1262 # src[] = { index }.
1263 load("shared_dxil", [1], [], [CAN_ELIMINATE])
1264 # src[] = { index }.
1265 load("scratch_dxil", [1], [], [CAN_ELIMINATE])
1266 # src[] = { deref_var, offset }
1267 load("ptr_dxil", [1, 1], [], [])
1268 # src[] = { index, 16-byte-based-offset }
1269 load("ubo_dxil", [1, 1], [], [CAN_ELIMINATE, CAN_REORDER])
1271 # DXIL Shared atomic intrinsics
1273 # All of the shared variable atomic memory operations read a value from
1274 # memory, compute a new value using one of the operations below, write the
1275 # new value to memory, and return the original value read.
1277 # All operations take 2 sources:
1279 # 0: The index in the i32 array for by the shared memory region
1280 # 1: The data parameter to the atomic function (i.e. the value to add
1281 # in shared_atomic_add, etc).
1282 intrinsic("shared_atomic_add_dxil", src_comp=[1, 1], dest_comp=1)
1283 intrinsic("shared_atomic_imin_dxil", src_comp=[1, 1], dest_comp=1)
1284 intrinsic("shared_atomic_umin_dxil", src_comp=[1, 1], dest_comp=1)
1285 intrinsic("shared_atomic_imax_dxil", src_comp=[1, 1], dest_comp=1)
1286 intrinsic("shared_atomic_umax_dxil", src_comp=[1, 1], dest_comp=1)
1287 intrinsic("shared_atomic_and_dxil", src_comp=[1, 1], dest_comp=1)
1288 intrinsic("shared_atomic_or_dxil", src_comp=[1, 1], dest_comp=1)
1289 intrinsic("shared_atomic_xor_dxil", src_comp=[1, 1], dest_comp=1)
1290 intrinsic("shared_atomic_exchange_dxil", src_comp=[1, 1], dest_comp=1)
1291 intrinsic("shared_atomic_comp_swap_dxil", src_comp=[1, 1, 1], dest_comp=1)
1293 # Intrinsics used by the Midgard/Bifrost blend pipeline. These are defined
1294 # within a blend shader to read/write the raw value from the tile buffer,
1295 # without applying any format conversion in the process. If the shader needs
1296 # usable pixel values, it must apply format conversions itself.
1298 # These definitions are generic, but they are explicitly vendored to prevent
1299 # other drivers from using them, as their semantics is defined in terms of the
1300 # Midgard/Bifrost hardware tile buffer and may not line up with anything sane.
1301 # One notable divergence is sRGB, which is asymmetric: raw_input_pan requires
1302 # an sRGB->linear conversion, but linear values should be written to
1303 # raw_output_pan and the hardware handles linear->sRGB.
1305 # store_raw_output_pan is used only for blend shaders, and writes out only a
1306 # single 128-bit chunk. To support multisampling, the BASE index specifies the
1307 # bas sample index written out.
1310 store("raw_output_pan", [], [IO_SEMANTICS, BASE])
1311 store("combined_output_pan", [1, 1, 1, 4], [IO_SEMANTICS, COMPONENT, SRC_TYPE, DEST_TYPE])
1312 load("raw_output_pan", [1], [IO_SEMANTICS], [CAN_ELIMINATE, CAN_REORDER])
1314 # Loads the sampler paramaters <min_lod, max_lod, lod_bias>
1315 # src[] = { sampler_index }
1316 load("sampler_lod_parameters_pan", [1], flags=[CAN_ELIMINATE, CAN_REORDER])
1318 # Like load_output but using a specified render target conversion descriptor
1319 load("converted_output_pan", [1], indices=[DEST_TYPE, IO_SEMANTICS], flags=[CAN_ELIMINATE])
1321 # Load the render target conversion descriptor for a given render target given
1322 # in the BASE index. Converts to a type with size given by the source type.
1323 # Valid in fragment and blend stages.
1324 system_value("rt_conversion_pan", 1, indices=[BASE, SRC_TYPE], bit_sizes=[32])
1326 # Loads the sample position array on Bifrost, in a packed Arm-specific format
1327 system_value("sample_positions_pan", 1, bit_sizes=[64])
1329 # In a fragment shader, is the framebuffer single-sampled? 0/~0 bool
1330 system_value("multisampled_pan", 1, bit_sizes=[32])
1332 # In a fragment shader, the current coverage mask. Affected by writes.
1333 intrinsic("load_coverage_mask_pan", [], 1, [], flags=[CAN_ELIMINATE],
1334 sysval=True, bit_sizes=[32])
1336 # R600 specific instrincs
1338 # location where the tesselation data is stored in LDS
1339 system_value("tcs_in_param_base_r600", 4)
1340 system_value("tcs_out_param_base_r600", 4)
1341 system_value("tcs_rel_patch_id_r600", 1)
1342 system_value("tcs_tess_factor_base_r600", 1)
1344 # the tess coords come as xy only, z has to be calculated
1345 system_value("tess_coord_r600", 2)
1347 # load as many components as needed giving per-component addresses
1348 intrinsic("load_local_shared_r600", src_comp=[0], dest_comp=0, indices = [], flags = [CAN_ELIMINATE])
1350 store("local_shared_r600", [1], [WRITE_MASK])
1351 store("tf_r600", [])
1353 # AMD GCN/RDNA specific intrinsics
1355 # This barrier is a hint that prevents moving the instruction that computes
1356 # src after this barrier. It's a constraint for the instruction scheduler.
1357 # Otherwise it's identical to a move instruction.
1358 # On AMD, it also forces the src value to be stored in a VGPR.
1359 intrinsic("optimization_barrier_vgpr_amd", dest_comp=0, src_comp=[0],
1360 flags=[CAN_ELIMINATE])
1362 # Untyped buffer load/store instructions of arbitrary length.
1363 # src[] = { descriptor, vector byte offset, scalar byte offset, index offset }
1364 # The index offset is multiplied by the stride in the descriptor.
1365 # The vector/scalar offsets are in bytes, BASE is a constant byte offset.
1366 intrinsic("load_buffer_amd", src_comp=[4, 1, 1, 1], dest_comp=0, indices=[BASE, MEMORY_MODES, ACCESS], flags=[CAN_ELIMINATE])
1367 # src[] = { store value, descriptor, vector byte offset, scalar byte offset, index offset }
1368 intrinsic("store_buffer_amd", src_comp=[0, 4, 1, 1, 1], indices=[BASE, WRITE_MASK, MEMORY_MODES, ACCESS])
1370 # Typed buffer load of arbitrary length, using a specified format.
1371 # src[] = { descriptor, vector byte offset, scalar byte offset, index offset }
1373 # The compiler backend is responsible for emitting correct HW instructions according to alignment, range etc.
1374 # Users of this intrinsic must ensure that the first component being loaded is really the first component
1375 # of the specified format, because range analysis assumes this.
1376 # The size of the specified format also determines the memory range that this instruction is allowed to access.
1378 # The index offset is multiplied by the stride in the descriptor, if any.
1379 # The vector/scalar offsets are in bytes, BASE is a constant byte offset.
1380 intrinsic("load_typed_buffer_amd", src_comp=[4, 1, 1, 1], dest_comp=0, indices=[BASE, MEMORY_MODES, ACCESS, FORMAT, ALIGN_MUL, ALIGN_OFFSET], flags=[CAN_ELIMINATE])
1382 # src[] = { address, unsigned 32-bit offset }.
1383 load("global_amd", [1, 1], indices=[BASE, ACCESS, ALIGN_MUL, ALIGN_OFFSET], flags=[CAN_ELIMINATE])
1384 # src[] = { value, address, unsigned 32-bit offset }.
1385 store("global_amd", [1, 1], indices=[BASE, ACCESS, ALIGN_MUL, ALIGN_OFFSET, WRITE_MASK])
1387 # Same as shared_atomic_add, but with GDS. src[] = {store_val, gds_addr, m0}
1388 intrinsic("gds_atomic_add_amd", src_comp=[1, 1, 1], dest_comp=1, indices=[BASE])
1390 # src[] = { descriptor, add_value }
1391 intrinsic("buffer_atomic_add_amd", src_comp=[4, 1], dest_comp=1, indices=[BASE])
1393 # src[] = { sample_id, num_samples }
1394 intrinsic("load_sample_positions_amd", src_comp=[1, 1], dest_comp=2, flags=[CAN_ELIMINATE, CAN_REORDER])
1396 # Descriptor where TCS outputs are stored for TES
1397 system_value("ring_tess_offchip_amd", 4)
1398 system_value("ring_tess_offchip_offset_amd", 1)
1399 # Descriptor where TCS outputs are stored for the HW tessellator
1400 system_value("ring_tess_factors_amd", 4)
1401 system_value("ring_tess_factors_offset_amd", 1)
1402 # Descriptor where ES outputs are stored for GS to read on GFX6-8
1403 system_value("ring_esgs_amd", 4)
1404 system_value("ring_es2gs_offset_amd", 1)
1405 # Address of the task shader draw ring (used for VARYING_SLOT_TASK_COUNT)
1406 system_value("ring_task_draw_amd", 4)
1407 # Address of the task shader payload ring (used for all other outputs)
1408 system_value("ring_task_payload_amd", 4)
1409 # Address of the mesh shader scratch ring (used for excess mesh shader outputs)
1410 system_value("ring_mesh_scratch_amd", 4)
1411 system_value("ring_mesh_scratch_offset_amd", 1)
1412 # Pointer into the draw and payload rings
1413 system_value("task_ring_entry_amd", 1)
1414 # Descriptor where NGG attributes are stored on GFX11.
1415 system_value("ring_attr_amd", 4)
1416 system_value("ring_attr_offset_amd", 1)
1418 # Number of patches processed by each TCS workgroup
1419 system_value("tcs_num_patches_amd", 1)
1420 # Relative tessellation patch ID within the current workgroup
1421 system_value("tess_rel_patch_id_amd", 1)
1422 # Vertex offsets used for GS per-vertex inputs
1423 system_value("gs_vertex_offset_amd", 1, [BASE])
1424 # Number of rasterization samples
1425 system_value("rasterization_samples_amd", 1)
1427 # Descriptor where GS outputs are stored for GS copy shader to read on GFX6-9
1428 system_value("ring_gsvs_amd", 4, indices=[STREAM_ID])
1429 # Write offset in gsvs ring for legacy GS shader
1430 system_value("ring_gs2vs_offset_amd", 1)
1432 # Streamout configuration
1433 system_value("streamout_config_amd", 1)
1434 # Position to write within the streamout buffers
1435 system_value("streamout_write_index_amd", 1)
1436 # Offset to write within a streamout buffer
1437 system_value("streamout_offset_amd", 1, indices=[BASE])
1439 # AMD merged shader intrinsics
1441 # Whether the current invocation index in the subgroup is less than the source. The source must be
1442 # subgroup uniform and bits 0-7 must be less than or equal to the wave size.
1443 intrinsic("is_subgroup_invocation_lt_amd", src_comp=[1], dest_comp=1, bit_sizes=[1], flags=[CAN_ELIMINATE])
1445 # AMD NGG intrinsics
1447 # Number of initial input vertices in the current workgroup.
1448 system_value("workgroup_num_input_vertices_amd", 1)
1449 # Number of initial input primitives in the current workgroup.
1450 system_value("workgroup_num_input_primitives_amd", 1)
1451 # For NGG passthrough mode only. Pre-packed argument for export_primitive_amd.
1452 system_value("packed_passthrough_primitive_amd", 1)
1453 # Whether NGG should execute shader query for pipeline statistics.
1454 system_value("pipeline_stat_query_enabled_amd", dest_comp=1, bit_sizes=[1])
1455 # Whether NGG should execute shader query for primitive generated.
1456 system_value("prim_gen_query_enabled_amd", dest_comp=1, bit_sizes=[1])
1457 # Whether NGG should execute shader query for primitive streamouted.
1458 system_value("prim_xfb_query_enabled_amd", dest_comp=1, bit_sizes=[1])
1459 # Merged wave info. Bits 0-7 are the ES thread count, 8-15 are the GS thread count, 16-24 is the
1460 # GS Wave ID, 24-27 is the wave index in the workgroup, and 28-31 is the workgroup size in waves.
1461 system_value("merged_wave_info_amd", dest_comp=1)
1462 # Whether the shader should clamp vertex color outputs to [0, 1].
1463 system_value("clamp_vertex_color_amd", dest_comp=1, bit_sizes=[1])
1464 # Whether the shader should cull front facing triangles.
1465 intrinsic("load_cull_front_face_enabled_amd", dest_comp=1, bit_sizes=[1], flags=[CAN_ELIMINATE])
1466 # Whether the shader should cull back facing triangles.
1467 intrinsic("load_cull_back_face_enabled_amd", dest_comp=1, bit_sizes=[1], flags=[CAN_ELIMINATE])
1468 # True if face culling should use CCW (false if CW).
1469 intrinsic("load_cull_ccw_amd", dest_comp=1, bit_sizes=[1], flags=[CAN_ELIMINATE])
1470 # Whether the shader should cull small primitives that are not visible in a pixel.
1471 intrinsic("load_cull_small_primitives_enabled_amd", dest_comp=1, bit_sizes=[1], flags=[CAN_ELIMINATE])
1472 # Whether any culling setting is enabled in the shader.
1473 intrinsic("load_cull_any_enabled_amd", dest_comp=1, bit_sizes=[1], flags=[CAN_ELIMINATE])
1474 # Small primitive culling precision
1475 intrinsic("load_cull_small_prim_precision_amd", dest_comp=1, bit_sizes=[32], flags=[CAN_ELIMINATE, CAN_REORDER])
1476 # Initial edge flags in a Vertex Shader, packed into the format the HW needs for primitive export.
1477 intrinsic("load_initial_edgeflags_amd", src_comp=[], dest_comp=1, bit_sizes=[32], indices=[])
1478 # Allocates export space for vertices and primitives. src[] = {num_vertices, num_primitives}.
1479 intrinsic("alloc_vertices_and_primitives_amd", src_comp=[1, 1], indices=[])
1480 # Overwrites VS input registers, for use with vertex compaction after culling. src = {vertex_id, instance_id}.
1481 intrinsic("overwrite_vs_arguments_amd", src_comp=[1, 1], indices=[])
1482 # Overwrites TES input registers, for use with vertex compaction after culling. src = {tes_u, tes_v, rel_patch_id, patch_id}.
1483 intrinsic("overwrite_tes_arguments_amd", src_comp=[1, 1, 1, 1], indices=[])
1485 # The address of the sbt descriptors.
1486 system_value("sbt_base_amd", 1, bit_sizes=[64])
1489 # 2. BVH node(64-bit pointer as 2x32 ...)
1493 # 6. inverse ray direction (componentwise 1.0/ray direction)
1494 intrinsic("bvh64_intersect_ray_amd", [4, 2, 1, 3, 3, 3], 4, flags=[CAN_ELIMINATE, CAN_REORDER])
1496 # Return of a callable in raytracing pipelines
1497 intrinsic("rt_return_amd")
1499 # offset into scratch for the input callable data in a raytracing pipeline.
1500 system_value("rt_arg_scratch_offset_amd", 1)
1502 # Whether to call the anyhit shader for an intersection in an intersection shader.
1503 system_value("intersection_opaque_amd", 1, bit_sizes=[1])
1505 # Used for indirect ray tracing.
1506 system_value("ray_launch_size_addr_amd", 1, bit_sizes=[64])
1508 # Scratch base of callable stack for ray tracing.
1509 system_value("rt_dynamic_callable_stack_base_amd", 1)
1511 # Ray Tracing Traversal inputs
1512 system_value("sbt_offset_amd", 1)
1513 system_value("sbt_stride_amd", 1)
1514 system_value("accel_struct_amd", 1, bit_sizes=[64])
1515 system_value("cull_mask_and_flags_amd", 1)
1521 # 4. Geometry Id and Flags
1523 intrinsic("execute_closest_hit_amd", src_comp=[1, 1, 1, 1, 1, 1])
1526 intrinsic("execute_miss_amd", src_comp=[1])
1528 # Used for saving and restoring hit attribute variables.
1530 intrinsic("load_hit_attrib_amd", dest_comp=1, bit_sizes=[32], indices=[BASE])
1531 intrinsic("store_hit_attrib_amd", src_comp=[1], indices=[BASE])
1533 # Load forced VRS rates.
1534 intrinsic("load_force_vrs_rates_amd", dest_comp=1, bit_sizes=[32], flags=[CAN_ELIMINATE, CAN_REORDER])
1536 intrinsic("load_scalar_arg_amd", dest_comp=0, bit_sizes=[32], indices=[BASE, ARG_UPPER_BOUND_U32_AMD], flags=[CAN_ELIMINATE, CAN_REORDER])
1537 intrinsic("load_vector_arg_amd", dest_comp=0, bit_sizes=[32], indices=[BASE, ARG_UPPER_BOUND_U32_AMD], flags=[CAN_ELIMINATE, CAN_REORDER])
1539 # src[] = { 32/64-bit base address, 32-bit offset }.
1540 intrinsic("load_smem_amd", src_comp=[1, 1], dest_comp=0, bit_sizes=[32],
1541 indices=[ALIGN_MUL, ALIGN_OFFSET],
1542 flags=[CAN_ELIMINATE, CAN_REORDER])
1544 # src[] = { offset }.
1545 intrinsic("load_shared2_amd", [1], dest_comp=2, indices=[OFFSET0, OFFSET1, ST64], flags=[CAN_ELIMINATE])
1547 # src[] = { value, offset }.
1548 intrinsic("store_shared2_amd", [2, 1], indices=[OFFSET0, OFFSET1, ST64])
1550 # Vertex stride in LS-HS buffer
1551 system_value("lshs_vertex_stride_amd", 1)
1553 # Vertex stride in ES-GS buffer
1554 system_value("esgs_vertex_stride_amd", 1)
1556 # Per patch data offset in HS VRAM output buffer
1557 system_value("hs_out_patch_data_offset_amd", 1)
1559 # line_width * 0.5 / abs(viewport_scale[2])
1560 system_value("clip_half_line_width_amd", 2)
1562 # Number of vertices in a primitive
1563 system_value("num_vertices_per_primitive_amd", 1)
1565 # Load streamout buffer desc
1566 # BASE = buffer index
1567 intrinsic("load_streamout_buffer_amd", dest_comp=4, indices=[BASE], bit_sizes=[32], flags=[CAN_ELIMINATE, CAN_REORDER])
1569 # An ID for each workgroup ordered by primitve sequence
1570 system_value("ordered_id_amd", 1)
1572 # Add src1 to global streamout buffer offsets in the specified order
1573 # src[] = { ordered_id, counter }
1574 # WRITE_MASK = mask for counter channel to update
1575 intrinsic("ordered_xfb_counter_add_amd", dest_comp=0, src_comp=[1, 0], indices=[WRITE_MASK], bit_sizes=[32])
1576 # Subtract from global streamout buffer offsets. Used to fix up the offsets
1577 # when we overflow streamout buffers.
1578 # src[] = { offsets }
1579 # WRITE_MASK = mask of offsets to subtract
1580 intrinsic("xfb_counter_sub_amd", src_comp=[0], indices=[WRITE_MASK], bit_sizes=[32])
1582 # Provoking vertex index in a primitive
1583 system_value("provoking_vtx_in_prim_amd", 1)
1585 # Atomically add current wave's primitive count to query result
1586 # * GS emitted primitive is primitive emitted by any GS stream
1587 # * generated primitive is primitive that has been produced for that stream by VS/TES/GS
1588 # * streamout primitve is primitve that has been written to xfb buffer, may be different
1589 # than generated primitive when xfb buffer is too small to hold more primitives
1590 # src[] = { primitive_count }.
1591 intrinsic("atomic_add_gs_emit_prim_count_amd", [1])
1592 intrinsic("atomic_add_gen_prim_count_amd", [1], indices=[STREAM_ID])
1593 intrinsic("atomic_add_xfb_prim_count_amd", [1], indices=[STREAM_ID])
1595 # Atomically add current wave's invocation count to query result
1596 # src[] = { invocation_count }.
1597 intrinsic("atomic_add_gs_invocation_count_amd", [1])
1599 # LDS offset for scratch section in NGG shader
1600 system_value("lds_ngg_scratch_base_amd", 1)
1601 # LDS offset for NGG GS shader vertex emit
1602 system_value("lds_ngg_gs_out_vertex_base_amd", 1)
1604 # AMD GPU shader output export instruction
1605 # src[] = { export_value }
1606 # BASE = export target
1607 # FLAGS = AC_EXP_FLAG_*
1608 intrinsic("export_amd", [0], indices=[BASE, WRITE_MASK, FLAGS])
1610 # Alpha test reference value
1611 system_value("alpha_reference_amd", 1)
1613 # Whether to enable barycentric optimization
1614 system_value("barycentric_optimize_amd", dest_comp=1, bit_sizes=[1])
1616 # V3D-specific instrinc for tile buffer color reads.
1618 # The hardware requires that we read the samples and components of a pixel
1619 # in order, so we cannot eliminate or remove any loads in a sequence.
1621 # src[] = { render_target }
1622 # BASE = sample index
1623 load("tlb_color_v3d", [1], [BASE, COMPONENT], [])
1625 # V3D-specific instrinc for per-sample tile buffer color writes.
1627 # The driver backend needs to identify per-sample color writes and emit
1628 # specific code for them.
1630 # src[] = { value, render_target }
1631 # BASE = sample index
1632 store("tlb_sample_color_v3d", [1], [BASE, COMPONENT, SRC_TYPE], [])
1634 # V3D-specific intrinsic to load the number of layers attached to
1635 # the target framebuffer
1636 intrinsic("load_fb_layers_v3d", dest_comp=1, flags=[CAN_ELIMINATE, CAN_REORDER])
1638 # Load/store a pixel in local memory. This operation is formatted, with
1639 # conversion between the specified format and the implied register format of the
1640 # source/destination (for store/loads respectively). This mostly matters for
1641 # converting between floating-point registers and normalized memory formats.
1643 # The format is the pipe_format of the local memory (the source), see
1644 # agx_internal_formats.h for the supported list.
1646 # Logically, this loads/stores a single sample. The sample to load is
1647 # specified by the bitfield sample mask source. However, for stores multiple
1648 # bits of the sample mask may be set, which will replicate the value. For
1649 # pixel rate shading, use 0xFF as the mask to store to all samples regardless of
1652 # All calculations are relative to an immediate byte offset into local
1653 # memory, which acts relative to the start of the sample. These instructions
1656 # (((((y * tile_width) + x) * nr_samples) + sample) * sample_stride) + offset
1658 # src[] = { sample mask }
1660 load("local_pixel_agx", [1], [BASE, FORMAT], [CAN_REORDER, CAN_ELIMINATE])
1661 # src[] = { value, sample mask }
1663 store("local_pixel_agx", [1], [BASE, WRITE_MASK, FORMAT], [CAN_REORDER])
1665 # Combined depth/stencil emit, applying to a mask of samples. base indicates
1666 # which to write (1 = depth, 2 = stencil, 3 = both).
1668 # src[] = { sample mask, depth, stencil }
1669 intrinsic("store_zs_agx", [1, 1, 1], indices=[BASE], flags=[])
1671 # Store a block from local memory into a bound image. Used to write out render
1672 # targets within the end-of-tile shader, although it is valid in general compute
1675 # The format is the pipe_format of the local memory (the source), see
1676 # agx_internal_formats.h for the supported list. The image format is
1677 # specified in the PBE descriptor.
1679 # The image dimension is used to distinguish multisampled images from
1680 # non-multisampled images. It must be 2D or MS.
1682 # src[] = { image index, logical offset within shared memory }
1683 intrinsic("block_image_store_agx", [1, 1], bit_sizes=[32, 16],
1684 indices=[FORMAT, IMAGE_DIM], flags=[CAN_REORDER])
1686 # Formatted load/store. The format is the pipe_format in memory (see
1687 # agx_internal_formats.h for the supported list). This accesses:
1689 # address + extend(index) << (format shift + shift)
1691 # The nir_intrinsic_base() index encodes the shift. The sign_extend index
1692 # determines whether sign- or zero-extension is used for the index.
1694 # All loads and stores on AGX uses these hardware instructions, so while these are
1695 # logically load_global_agx/load_global_constant_agx/store_global_agx, the
1696 # _global is omitted as it adds nothing.
1698 # src[] = { address, index }.
1699 load("agx", [1, 1], [ACCESS, BASE, FORMAT, SIGN_EXTEND], [CAN_ELIMINATE])
1700 load("constant_agx", [1, 1], [ACCESS, BASE, FORMAT, SIGN_EXTEND],
1701 [CAN_ELIMINATE, CAN_REORDER])
1702 # src[] = { value, address, index }.
1703 store("agx", [1, 1], [ACCESS, BASE, FORMAT, SIGN_EXTEND])
1705 # Logical complement of load_front_face, mapping to an AGX system value
1706 system_value("back_face_agx", 1, bit_sizes=[1, 32])
1708 # Loads the texture descriptor base for indexed (non-bindless) textures. On G13,
1709 # the referenced array has stride 24.
1710 system_value("texture_base_agx", 1, bit_sizes=[64])
1712 # Load the base address of an indexed UBO/VBO (for lowering UBOs/VBOs)
1713 intrinsic("load_ubo_base_agx", src_comp=[1], dest_comp=1, bit_sizes=[64],
1714 flags=[CAN_ELIMINATE, CAN_REORDER])
1715 intrinsic("load_vbo_base_agx", src_comp=[1], dest_comp=1, bit_sizes=[64],
1716 flags=[CAN_ELIMINATE, CAN_REORDER])
1718 # Intel-specific query for loading from the brw_image_param struct passed
1719 # into the shader as a uniform. The variable is a deref to the image
1720 # variable. The const index specifies which of the six parameters to load.
1721 intrinsic("image_deref_load_param_intel", src_comp=[1], dest_comp=0,
1722 indices=[BASE], flags=[CAN_ELIMINATE, CAN_REORDER])
1723 image("load_raw_intel", src_comp=[1], dest_comp=0,
1724 flags=[CAN_ELIMINATE])
1725 image("store_raw_intel", src_comp=[1, 0])
1727 # Intrinsic to load a block of at least 32B of constant data from a 64-bit
1728 # global memory address. The memory address must be uniform and 32B-aligned.
1729 # The second source is a predicate which indicates whether or not to actually
1731 # src[] = { address, predicate }.
1732 intrinsic("load_global_const_block_intel", src_comp=[1, 1], dest_comp=0,
1733 bit_sizes=[32], indices=[BASE], flags=[CAN_ELIMINATE, CAN_REORDER])
1735 # Number of data items being operated on for a SIMD program.
1736 system_value("simd_width_intel", 1)
1738 # Load a relocatable 32-bit value
1739 intrinsic("load_reloc_const_intel", dest_comp=1, bit_sizes=[32],
1740 indices=[PARAM_IDX], flags=[CAN_ELIMINATE, CAN_REORDER])
1742 # 64-bit global address for a Vulkan descriptor set
1744 intrinsic("load_desc_set_address_intel", dest_comp=1, bit_sizes=[64],
1745 src_comp=[1], flags=[CAN_ELIMINATE, CAN_REORDER])
1747 # Base offset for a given set in the flatten array of dynamic offsets
1749 intrinsic("load_desc_set_dynamic_index_intel", dest_comp=1, bit_sizes=[32],
1750 src_comp=[1], flags=[CAN_ELIMINATE, CAN_REORDER])
1752 # OpSubgroupBlockReadINTEL and OpSubgroupBlockWriteINTEL from SPV_INTEL_subgroups.
1753 intrinsic("load_deref_block_intel", dest_comp=0, src_comp=[-1],
1754 indices=[ACCESS], flags=[CAN_ELIMINATE])
1755 intrinsic("store_deref_block_intel", src_comp=[-1, 0], indices=[WRITE_MASK, ACCESS])
1757 # src[] = { address }.
1758 load("global_block_intel", [1], [ACCESS, ALIGN_MUL, ALIGN_OFFSET], [CAN_ELIMINATE])
1760 # src[] = { buffer_index, offset }.
1761 load("ssbo_block_intel", [-1, 1], [ACCESS, ALIGN_MUL, ALIGN_OFFSET], [CAN_ELIMINATE])
1763 # src[] = { offset }.
1764 load("shared_block_intel", [1], [BASE, ALIGN_MUL, ALIGN_OFFSET], [CAN_ELIMINATE])
1766 # src[] = { value, address }.
1767 store("global_block_intel", [1], [WRITE_MASK, ACCESS, ALIGN_MUL, ALIGN_OFFSET])
1769 # src[] = { value, block_index, offset }
1770 store("ssbo_block_intel", [-1, 1], [WRITE_MASK, ACCESS, ALIGN_MUL, ALIGN_OFFSET])
1772 # src[] = { value, offset }.
1773 store("shared_block_intel", [1], [BASE, WRITE_MASK, ALIGN_MUL, ALIGN_OFFSET])
1775 # Similar to load_global_const_block_intel but for SSBOs
1776 # offset should be uniform
1777 # src[] = { buffer_index, offset }.
1778 load("ssbo_uniform_block_intel", [-1, 1], [ACCESS, ALIGN_MUL, ALIGN_OFFSET], [CAN_ELIMINATE])
1780 # Similar to load_global_const_block_intel but for shared memory
1781 # src[] = { offset }.
1782 load("shared_uniform_block_intel", [1], [BASE, ALIGN_MUL, ALIGN_OFFSET], [CAN_ELIMINATE])
1784 # Intrinsics for Intel mesh shading
1785 system_value("mesh_inline_data_intel", 1, [ALIGN_OFFSET], bit_sizes=[32, 64])
1787 # Intrinsics for Intel bindless thread dispatch
1788 # BASE=brw_topoloy_id
1789 system_value("topology_id_intel", 1, indices=[BASE])
1790 system_value("btd_stack_id_intel", 1)
1791 system_value("btd_global_arg_addr_intel", 1, bit_sizes=[64])
1792 system_value("btd_local_arg_addr_intel", 1, bit_sizes=[64])
1793 system_value("btd_resume_sbt_addr_intel", 1, bit_sizes=[64])
1794 # src[] = { global_arg_addr, btd_record }
1795 intrinsic("btd_spawn_intel", src_comp=[1, 1])
1797 intrinsic("btd_stack_push_intel", indices=[STACK_SIZE])
1799 intrinsic("btd_retire_intel")
1801 # Intel-specific ray-tracing intrinsic
1802 # src[] = { globals, level, operation } SYNCHRONOUS=synchronous
1803 intrinsic("trace_ray_intel", src_comp=[1, 1, 1], indices=[SYNCHRONOUS])
1805 # System values used for ray-tracing on Intel
1806 system_value("ray_base_mem_addr_intel", 1, bit_sizes=[64])
1807 system_value("ray_hw_stack_size_intel", 1)
1808 system_value("ray_sw_stack_size_intel", 1)
1809 system_value("ray_num_dss_rt_stacks_intel", 1)
1810 system_value("ray_hit_sbt_addr_intel", 1, bit_sizes=[64])
1811 system_value("ray_hit_sbt_stride_intel", 1, bit_sizes=[16])
1812 system_value("ray_miss_sbt_addr_intel", 1, bit_sizes=[64])
1813 system_value("ray_miss_sbt_stride_intel", 1, bit_sizes=[16])
1814 system_value("callable_sbt_addr_intel", 1, bit_sizes=[64])
1815 system_value("callable_sbt_stride_intel", 1, bit_sizes=[16])
1816 system_value("leaf_opaque_intel", 1, bit_sizes=[1])
1817 system_value("leaf_procedural_intel", 1, bit_sizes=[1])
1823 system_value("btd_shader_type_intel", 1)
1824 system_value("ray_query_global_intel", 1, bit_sizes=[64])
1826 # In order to deal with flipped render targets, gl_PointCoord may be flipped
1827 # in the shader requiring a shader key or extra instructions or it may be
1828 # flipped in hardware based on a state bit. This version of gl_PointCoord
1829 # is defined to be whatever thing the hardware can easily give you, so long as
1830 # it's in normalized coordinates in the range [0, 1] across the point.
1831 intrinsic("load_point_coord_maybe_flipped", dest_comp=2, bit_sizes=[32])