4 TGSI, Tungsten Graphics Shader Infrastructure, is an intermediate language
5 for describing shaders. Since Gallium is inherently shaderful, shaders are
6 an important part of the API. TGSI is the only intermediate representation
12 All TGSI instructions, known as *opcodes*, operate on arbitrary-precision
13 floating-point four-component vectors. An opcode may have up to one
14 destination register, known as *dst*, and between zero and three source
15 registers, called *src0* through *src2*, or simply *src* if there is only
18 Some instructions, like :opcode:`I2F`, permit re-interpretation of vector
19 components as integers. Other instructions permit using registers as
20 two-component vectors with double precision; see :ref:`doubleopcodes`.
22 When an instruction has a scalar result, the result is usually copied into
23 each of the components of *dst*. When this happens, the result is said to be
24 *replicated* to *dst*. :opcode:`RCP` is one such instruction.
29 TGSI supports modifiers on inputs (as well as saturate modifier on instructions).
31 For inputs which have a floating point type, both absolute value and negation
32 modifiers are supported (with absolute value being applied first).
33 TGSI_OPCODE_MOV is considered to have float input type for applying modifiers.
35 For inputs which have signed or unsigned type only the negate modifier is
42 ^^^^^^^^^^^^^^^^^^^^^^^^^
44 These opcodes are guaranteed to be available regardless of the driver being
47 .. opcode:: ARL - Address Register Load
51 dst.x = (int) \lfloor src.x\rfloor
53 dst.y = (int) \lfloor src.y\rfloor
55 dst.z = (int) \lfloor src.z\rfloor
57 dst.w = (int) \lfloor src.w\rfloor
60 .. opcode:: MOV - Move
73 .. opcode:: LIT - Light Coefficients
78 dst.y &= max(src.x, 0) \\
79 dst.z &= (src.x > 0) ? max(src.y, 0)^{clamp(src.w, -128, 128))} : 0 \\
83 .. opcode:: RCP - Reciprocal
85 This instruction replicates its result.
92 .. opcode:: RSQ - Reciprocal Square Root
94 This instruction replicates its result. The results are undefined for src <= 0.
98 dst = \frac{1}{\sqrt{src.x}}
101 .. opcode:: SQRT - Square Root
103 This instruction replicates its result. The results are undefined for src < 0.
110 .. opcode:: EXP - Approximate Exponential Base 2
114 dst.x &= 2^{\lfloor src.x\rfloor} \\
115 dst.y &= src.x - \lfloor src.x\rfloor \\
116 dst.z &= 2^{src.x} \\
120 .. opcode:: LOG - Approximate Logarithm Base 2
124 dst.x &= \lfloor\log_2{|src.x|}\rfloor \\
125 dst.y &= \frac{|src.x|}{2^{\lfloor\log_2{|src.x|}\rfloor}} \\
126 dst.z &= \log_2{|src.x|} \\
130 .. opcode:: MUL - Multiply
134 dst.x = src0.x \times src1.x
136 dst.y = src0.y \times src1.y
138 dst.z = src0.z \times src1.z
140 dst.w = src0.w \times src1.w
143 .. opcode:: ADD - Add
147 dst.x = src0.x + src1.x
149 dst.y = src0.y + src1.y
151 dst.z = src0.z + src1.z
153 dst.w = src0.w + src1.w
156 .. opcode:: DP3 - 3-component Dot Product
158 This instruction replicates its result.
162 dst = src0.x \times src1.x + src0.y \times src1.y + src0.z \times src1.z
165 .. opcode:: DP4 - 4-component Dot Product
167 This instruction replicates its result.
171 dst = src0.x \times src1.x + src0.y \times src1.y + src0.z \times src1.z + src0.w \times src1.w
174 .. opcode:: DST - Distance Vector
179 dst.y &= src0.y \times src1.y\\
184 .. opcode:: MIN - Minimum
188 dst.x = min(src0.x, src1.x)
190 dst.y = min(src0.y, src1.y)
192 dst.z = min(src0.z, src1.z)
194 dst.w = min(src0.w, src1.w)
197 .. opcode:: MAX - Maximum
201 dst.x = max(src0.x, src1.x)
203 dst.y = max(src0.y, src1.y)
205 dst.z = max(src0.z, src1.z)
207 dst.w = max(src0.w, src1.w)
210 .. opcode:: SLT - Set On Less Than
214 dst.x = (src0.x < src1.x) ? 1.0F : 0.0F
216 dst.y = (src0.y < src1.y) ? 1.0F : 0.0F
218 dst.z = (src0.z < src1.z) ? 1.0F : 0.0F
220 dst.w = (src0.w < src1.w) ? 1.0F : 0.0F
223 .. opcode:: SGE - Set On Greater Equal Than
227 dst.x = (src0.x >= src1.x) ? 1.0F : 0.0F
229 dst.y = (src0.y >= src1.y) ? 1.0F : 0.0F
231 dst.z = (src0.z >= src1.z) ? 1.0F : 0.0F
233 dst.w = (src0.w >= src1.w) ? 1.0F : 0.0F
236 .. opcode:: MAD - Multiply And Add
240 dst.x = src0.x \times src1.x + src2.x
242 dst.y = src0.y \times src1.y + src2.y
244 dst.z = src0.z \times src1.z + src2.z
246 dst.w = src0.w \times src1.w + src2.w
249 .. opcode:: LRP - Linear Interpolate
253 dst.x = src0.x \times src1.x + (1 - src0.x) \times src2.x
255 dst.y = src0.y \times src1.y + (1 - src0.y) \times src2.y
257 dst.z = src0.z \times src1.z + (1 - src0.z) \times src2.z
259 dst.w = src0.w \times src1.w + (1 - src0.w) \times src2.w
262 .. opcode:: FMA - Fused Multiply-Add
264 Perform a * b + c with no intermediate rounding step.
268 dst.x = src0.x \times src1.x + src2.x
270 dst.y = src0.y \times src1.y + src2.y
272 dst.z = src0.z \times src1.z + src2.z
274 dst.w = src0.w \times src1.w + src2.w
277 .. opcode:: DP2A - 2-component Dot Product And Add
281 dst.x = src0.x \times src1.x + src0.y \times src1.y + src2.x
283 dst.y = src0.x \times src1.x + src0.y \times src1.y + src2.x
285 dst.z = src0.x \times src1.x + src0.y \times src1.y + src2.x
287 dst.w = src0.x \times src1.x + src0.y \times src1.y + src2.x
290 .. opcode:: FRC - Fraction
294 dst.x = src.x - \lfloor src.x\rfloor
296 dst.y = src.y - \lfloor src.y\rfloor
298 dst.z = src.z - \lfloor src.z\rfloor
300 dst.w = src.w - \lfloor src.w\rfloor
303 .. opcode:: FLR - Floor
307 dst.x = \lfloor src.x\rfloor
309 dst.y = \lfloor src.y\rfloor
311 dst.z = \lfloor src.z\rfloor
313 dst.w = \lfloor src.w\rfloor
316 .. opcode:: ROUND - Round
329 .. opcode:: EX2 - Exponential Base 2
331 This instruction replicates its result.
338 .. opcode:: LG2 - Logarithm Base 2
340 This instruction replicates its result.
347 .. opcode:: POW - Power
349 This instruction replicates its result.
353 dst = src0.x^{src1.x}
355 .. opcode:: XPD - Cross Product
359 dst.x = src0.y \times src1.z - src1.y \times src0.z
361 dst.y = src0.z \times src1.x - src1.z \times src0.x
363 dst.z = src0.x \times src1.y - src1.x \times src0.y
368 .. opcode:: DPH - Homogeneous Dot Product
370 This instruction replicates its result.
374 dst = src0.x \times src1.x + src0.y \times src1.y + src0.z \times src1.z + src1.w
377 .. opcode:: COS - Cosine
379 This instruction replicates its result.
386 .. opcode:: DDX, DDX_FINE - Derivative Relative To X
388 The fine variant is only used when ``PIPE_CAP_TGSI_FS_FINE_DERIVATIVE`` is
389 advertised. When it is, the fine version guarantees one derivative per row
390 while DDX is allowed to be the same for the entire 2x2 quad.
394 dst.x = partialx(src.x)
396 dst.y = partialx(src.y)
398 dst.z = partialx(src.z)
400 dst.w = partialx(src.w)
403 .. opcode:: DDY, DDY_FINE - Derivative Relative To Y
405 The fine variant is only used when ``PIPE_CAP_TGSI_FS_FINE_DERIVATIVE`` is
406 advertised. When it is, the fine version guarantees one derivative per column
407 while DDY is allowed to be the same for the entire 2x2 quad.
411 dst.x = partialy(src.x)
413 dst.y = partialy(src.y)
415 dst.z = partialy(src.z)
417 dst.w = partialy(src.w)
420 .. opcode:: PK2H - Pack Two 16-bit Floats
422 This instruction replicates its result.
426 dst = f32\_to\_f16(src.x) | f32\_to\_f16(src.y) << 16
429 .. opcode:: PK2US - Pack Two Unsigned 16-bit Scalars
434 .. opcode:: PK4B - Pack Four Signed 8-bit Scalars
439 .. opcode:: PK4UB - Pack Four Unsigned 8-bit Scalars
444 .. opcode:: SEQ - Set On Equal
448 dst.x = (src0.x == src1.x) ? 1.0F : 0.0F
450 dst.y = (src0.y == src1.y) ? 1.0F : 0.0F
452 dst.z = (src0.z == src1.z) ? 1.0F : 0.0F
454 dst.w = (src0.w == src1.w) ? 1.0F : 0.0F
457 .. opcode:: SGT - Set On Greater Than
461 dst.x = (src0.x > src1.x) ? 1.0F : 0.0F
463 dst.y = (src0.y > src1.y) ? 1.0F : 0.0F
465 dst.z = (src0.z > src1.z) ? 1.0F : 0.0F
467 dst.w = (src0.w > src1.w) ? 1.0F : 0.0F
470 .. opcode:: SIN - Sine
472 This instruction replicates its result.
479 .. opcode:: SLE - Set On Less Equal Than
483 dst.x = (src0.x <= src1.x) ? 1.0F : 0.0F
485 dst.y = (src0.y <= src1.y) ? 1.0F : 0.0F
487 dst.z = (src0.z <= src1.z) ? 1.0F : 0.0F
489 dst.w = (src0.w <= src1.w) ? 1.0F : 0.0F
492 .. opcode:: SNE - Set On Not Equal
496 dst.x = (src0.x != src1.x) ? 1.0F : 0.0F
498 dst.y = (src0.y != src1.y) ? 1.0F : 0.0F
500 dst.z = (src0.z != src1.z) ? 1.0F : 0.0F
502 dst.w = (src0.w != src1.w) ? 1.0F : 0.0F
505 .. opcode:: TEX - Texture Lookup
507 for array textures src0.y contains the slice for 1D,
508 and src0.z contain the slice for 2D.
510 for shadow textures with no arrays (and not cube map),
511 src0.z contains the reference value.
513 for shadow textures with arrays, src0.z contains
514 the reference value for 1D arrays, and src0.w contains
515 the reference value for 2D arrays and cube maps.
517 for cube map array shadow textures, the reference value
518 cannot be passed in src0.w, and TEX2 must be used instead.
524 shadow_ref = src0.z or src0.w (optional)
528 dst = texture\_sample(unit, coord, shadow_ref)
531 .. opcode:: TEX2 - Texture Lookup (for shadow cube map arrays only)
533 this is the same as TEX, but uses another reg to encode the
544 dst = texture\_sample(unit, coord, shadow_ref)
549 .. opcode:: TXD - Texture Lookup with Derivatives
561 dst = texture\_sample\_deriv(unit, coord, ddx, ddy)
564 .. opcode:: TXP - Projective Texture Lookup
568 coord.x = src0.x / src0.w
570 coord.y = src0.y / src0.w
572 coord.z = src0.z / src0.w
578 dst = texture\_sample(unit, coord)
581 .. opcode:: UP2H - Unpack Two 16-Bit Floats
585 dst.x = f16\_to\_f32(src0.x \& 0xffff)
587 dst.y = f16\_to\_f32(src0.x >> 16)
589 dst.z = f16\_to\_f32(src0.x \& 0xffff)
591 dst.w = f16\_to\_f32(src0.x >> 16)
595 Considered for removal.
597 .. opcode:: UP2US - Unpack Two Unsigned 16-Bit Scalars
603 Considered for removal.
605 .. opcode:: UP4B - Unpack Four Signed 8-Bit Values
611 Considered for removal.
613 .. opcode:: UP4UB - Unpack Four Unsigned 8-Bit Scalars
619 Considered for removal.
622 .. opcode:: ARR - Address Register Load With Round
626 dst.x = (int) round(src.x)
628 dst.y = (int) round(src.y)
630 dst.z = (int) round(src.z)
632 dst.w = (int) round(src.w)
635 .. opcode:: SSG - Set Sign
639 dst.x = (src.x > 0) ? 1 : (src.x < 0) ? -1 : 0
641 dst.y = (src.y > 0) ? 1 : (src.y < 0) ? -1 : 0
643 dst.z = (src.z > 0) ? 1 : (src.z < 0) ? -1 : 0
645 dst.w = (src.w > 0) ? 1 : (src.w < 0) ? -1 : 0
648 .. opcode:: CMP - Compare
652 dst.x = (src0.x < 0) ? src1.x : src2.x
654 dst.y = (src0.y < 0) ? src1.y : src2.y
656 dst.z = (src0.z < 0) ? src1.z : src2.z
658 dst.w = (src0.w < 0) ? src1.w : src2.w
661 .. opcode:: KILL_IF - Conditional Discard
663 Conditional discard. Allowed in fragment shaders only.
667 if (src.x < 0 || src.y < 0 || src.z < 0 || src.w < 0)
672 .. opcode:: KILL - Discard
674 Unconditional discard. Allowed in fragment shaders only.
677 .. opcode:: SCS - Sine Cosine
690 .. opcode:: TXB - Texture Lookup With Bias
692 for cube map array textures and shadow cube maps, the bias value
693 cannot be passed in src0.w, and TXB2 must be used instead.
695 if the target is a shadow texture, the reference value is always
696 in src.z (this prevents shadow 3d and shadow 2d arrays from
697 using this instruction, but this is not needed).
713 dst = texture\_sample(unit, coord, bias)
716 .. opcode:: TXB2 - Texture Lookup With Bias (some cube maps only)
718 this is the same as TXB, but uses another reg to encode the
719 lod bias value for cube map arrays and shadow cube maps.
720 Presumably shadow 2d arrays and shadow 3d targets could use
721 this encoding too, but this is not legal.
723 shadow cube map arrays are neither possible nor required.
733 dst = texture\_sample(unit, coord, bias)
736 .. opcode:: DIV - Divide
740 dst.x = \frac{src0.x}{src1.x}
742 dst.y = \frac{src0.y}{src1.y}
744 dst.z = \frac{src0.z}{src1.z}
746 dst.w = \frac{src0.w}{src1.w}
749 .. opcode:: DP2 - 2-component Dot Product
751 This instruction replicates its result.
755 dst = src0.x \times src1.x + src0.y \times src1.y
758 .. opcode:: TEX_LZ - Texture Lookup With LOD = 0
760 This is the same as TXL with LOD = 0. Like every texture opcode, it obeys
761 pipe_sampler_view::u.tex.first_level and pipe_sampler_state::min_lod.
762 There is no way to override those two in shaders.
778 dst = texture\_sample(unit, coord, lod)
781 .. opcode:: TXL - Texture Lookup With explicit LOD
783 for cube map array textures, the explicit lod value
784 cannot be passed in src0.w, and TXL2 must be used instead.
786 if the target is a shadow texture, the reference value is always
787 in src.z (this prevents shadow 3d / 2d array / cube targets from
788 using this instruction, but this is not needed).
804 dst = texture\_sample(unit, coord, lod)
807 .. opcode:: TXL2 - Texture Lookup With explicit LOD (for cube map arrays only)
809 this is the same as TXL, but uses another reg to encode the
811 Presumably shadow 3d / 2d array / cube targets could use
812 this encoding too, but this is not legal.
814 shadow cube map arrays are neither possible nor required.
824 dst = texture\_sample(unit, coord, lod)
827 .. opcode:: PUSHA - Push Address Register On Stack
836 Considered for cleanup.
840 Considered for removal.
842 .. opcode:: POPA - Pop Address Register From Stack
851 Considered for cleanup.
855 Considered for removal.
858 .. opcode:: CALLNZ - Subroutine Call If Not Zero
864 Considered for cleanup.
868 Considered for removal.
872 ^^^^^^^^^^^^^^^^^^^^^^^^
874 These opcodes are primarily provided for special-use computational shaders.
875 Support for these opcodes indicated by a special pipe capability bit (TBD).
877 XXX doesn't look like most of the opcodes really belong here.
879 .. opcode:: CEIL - Ceiling
883 dst.x = \lceil src.x\rceil
885 dst.y = \lceil src.y\rceil
887 dst.z = \lceil src.z\rceil
889 dst.w = \lceil src.w\rceil
892 .. opcode:: TRUNC - Truncate
905 .. opcode:: MOD - Modulus
909 dst.x = src0.x \bmod src1.x
911 dst.y = src0.y \bmod src1.y
913 dst.z = src0.z \bmod src1.z
915 dst.w = src0.w \bmod src1.w
918 .. opcode:: UARL - Integer Address Register Load
920 Moves the contents of the source register, assumed to be an integer, into the
921 destination register, which is assumed to be an address (ADDR) register.
924 .. opcode:: SAD - Sum Of Absolute Differences
928 dst.x = |src0.x - src1.x| + src2.x
930 dst.y = |src0.y - src1.y| + src2.y
932 dst.z = |src0.z - src1.z| + src2.z
934 dst.w = |src0.w - src1.w| + src2.w
937 .. opcode:: TXF - Texel Fetch
939 As per NV_gpu_shader4, extract a single texel from a specified texture
940 image. The source sampler may not be a CUBE or SHADOW. src 0 is a
941 four-component signed integer vector used to identify the single texel
942 accessed. 3 components + level. Just like texture instructions, an optional
943 offset vector is provided, which is subject to various driver restrictions
944 (regarding range, source of offsets). This instruction ignores the sampler
947 TXF(uint_vec coord, int_vec offset).
950 .. opcode:: TXF_LZ - Texel Fetch
952 This is the same as TXF with level = 0. Like TXF, it obeys
953 pipe_sampler_view::u.tex.first_level.
956 .. opcode:: TXQ - Texture Size Query
958 As per NV_gpu_program4, retrieve the dimensions of the texture depending on
959 the target. For 1D (width), 2D/RECT/CUBE (width, height), 3D (width, height,
960 depth), 1D array (width, layers), 2D array (width, height, layers).
961 Also return the number of accessible levels (last_level - first_level + 1)
964 For components which don't return a resource dimension, their value
971 dst.x = texture\_width(unit, lod)
973 dst.y = texture\_height(unit, lod)
975 dst.z = texture\_depth(unit, lod)
977 dst.w = texture\_levels(unit)
980 .. opcode:: TXQS - Texture Samples Query
982 This retrieves the number of samples in the texture, and stores it
983 into the x component. The other components are undefined.
987 dst.x = texture\_samples(unit)
990 .. opcode:: TG4 - Texture Gather
992 As per ARB_texture_gather, gathers the four texels to be used in a bi-linear
993 filtering operation and packs them into a single register. Only works with
994 2D, 2D array, cubemaps, and cubemaps arrays. For 2D textures, only the
995 addressing modes of the sampler and the top level of any mip pyramid are
996 used. Set W to zero. It behaves like the TEX instruction, but a filtered
997 sample is not generated. The four samples that contribute to filtering are
998 placed into xyzw in clockwise order, starting with the (u,v) texture
999 coordinate delta at the following locations (-, +), (+, +), (+, -), (-, -),
1000 where the magnitude of the deltas are half a texel.
1002 PIPE_CAP_TEXTURE_SM5 enhances this instruction to support shadow per-sample
1003 depth compares, single component selection, and a non-constant offset. It
1004 doesn't allow support for the GL independent offset to get i0,j0. This would
1005 require another CAP is hw can do it natively. For now we lower that before
1014 dst = texture\_gather4 (unit, coord, component)
1016 (with SM5 - cube array shadow)
1024 dst = texture\_gather (uint, coord, compare)
1026 .. opcode:: LODQ - level of detail query
1028 Compute the LOD information that the texture pipe would use to access the
1029 texture. The Y component contains the computed LOD lambda_prime. The X
1030 component contains the LOD that will be accessed, based on min/max lod's
1037 dst.xy = lodq(uint, coord);
1040 ^^^^^^^^^^^^^^^^^^^^^^^^
1041 These opcodes are used for integer operations.
1042 Support for these opcodes indicated by PIPE_SHADER_CAP_INTEGERS (all of them?)
1045 .. opcode:: I2F - Signed Integer To Float
1047 Rounding is unspecified (round to nearest even suggested).
1051 dst.x = (float) src.x
1053 dst.y = (float) src.y
1055 dst.z = (float) src.z
1057 dst.w = (float) src.w
1060 .. opcode:: U2F - Unsigned Integer To Float
1062 Rounding is unspecified (round to nearest even suggested).
1066 dst.x = (float) src.x
1068 dst.y = (float) src.y
1070 dst.z = (float) src.z
1072 dst.w = (float) src.w
1075 .. opcode:: F2I - Float to Signed Integer
1077 Rounding is towards zero (truncate).
1078 Values outside signed range (including NaNs) produce undefined results.
1091 .. opcode:: F2U - Float to Unsigned Integer
1093 Rounding is towards zero (truncate).
1094 Values outside unsigned range (including NaNs) produce undefined results.
1098 dst.x = (unsigned) src.x
1100 dst.y = (unsigned) src.y
1102 dst.z = (unsigned) src.z
1104 dst.w = (unsigned) src.w
1107 .. opcode:: UADD - Integer Add
1109 This instruction works the same for signed and unsigned integers.
1110 The low 32bit of the result is returned.
1114 dst.x = src0.x + src1.x
1116 dst.y = src0.y + src1.y
1118 dst.z = src0.z + src1.z
1120 dst.w = src0.w + src1.w
1123 .. opcode:: UMAD - Integer Multiply And Add
1125 This instruction works the same for signed and unsigned integers.
1126 The multiplication returns the low 32bit (as does the result itself).
1130 dst.x = src0.x \times src1.x + src2.x
1132 dst.y = src0.y \times src1.y + src2.y
1134 dst.z = src0.z \times src1.z + src2.z
1136 dst.w = src0.w \times src1.w + src2.w
1139 .. opcode:: UMUL - Integer Multiply
1141 This instruction works the same for signed and unsigned integers.
1142 The low 32bit of the result is returned.
1146 dst.x = src0.x \times src1.x
1148 dst.y = src0.y \times src1.y
1150 dst.z = src0.z \times src1.z
1152 dst.w = src0.w \times src1.w
1155 .. opcode:: IMUL_HI - Signed Integer Multiply High Bits
1157 The high 32bits of the multiplication of 2 signed integers are returned.
1161 dst.x = (src0.x \times src1.x) >> 32
1163 dst.y = (src0.y \times src1.y) >> 32
1165 dst.z = (src0.z \times src1.z) >> 32
1167 dst.w = (src0.w \times src1.w) >> 32
1170 .. opcode:: UMUL_HI - Unsigned Integer Multiply High Bits
1172 The high 32bits of the multiplication of 2 unsigned integers are returned.
1176 dst.x = (src0.x \times src1.x) >> 32
1178 dst.y = (src0.y \times src1.y) >> 32
1180 dst.z = (src0.z \times src1.z) >> 32
1182 dst.w = (src0.w \times src1.w) >> 32
1185 .. opcode:: IDIV - Signed Integer Division
1187 TBD: behavior for division by zero.
1191 dst.x = src0.x \ src1.x
1193 dst.y = src0.y \ src1.y
1195 dst.z = src0.z \ src1.z
1197 dst.w = src0.w \ src1.w
1200 .. opcode:: UDIV - Unsigned Integer Division
1202 For division by zero, 0xffffffff is returned.
1206 dst.x = src0.x \ src1.x
1208 dst.y = src0.y \ src1.y
1210 dst.z = src0.z \ src1.z
1212 dst.w = src0.w \ src1.w
1215 .. opcode:: UMOD - Unsigned Integer Remainder
1217 If second arg is zero, 0xffffffff is returned.
1221 dst.x = src0.x \ src1.x
1223 dst.y = src0.y \ src1.y
1225 dst.z = src0.z \ src1.z
1227 dst.w = src0.w \ src1.w
1230 .. opcode:: NOT - Bitwise Not
1243 .. opcode:: AND - Bitwise And
1247 dst.x = src0.x \& src1.x
1249 dst.y = src0.y \& src1.y
1251 dst.z = src0.z \& src1.z
1253 dst.w = src0.w \& src1.w
1256 .. opcode:: OR - Bitwise Or
1260 dst.x = src0.x | src1.x
1262 dst.y = src0.y | src1.y
1264 dst.z = src0.z | src1.z
1266 dst.w = src0.w | src1.w
1269 .. opcode:: XOR - Bitwise Xor
1273 dst.x = src0.x \oplus src1.x
1275 dst.y = src0.y \oplus src1.y
1277 dst.z = src0.z \oplus src1.z
1279 dst.w = src0.w \oplus src1.w
1282 .. opcode:: IMAX - Maximum of Signed Integers
1286 dst.x = max(src0.x, src1.x)
1288 dst.y = max(src0.y, src1.y)
1290 dst.z = max(src0.z, src1.z)
1292 dst.w = max(src0.w, src1.w)
1295 .. opcode:: UMAX - Maximum of Unsigned Integers
1299 dst.x = max(src0.x, src1.x)
1301 dst.y = max(src0.y, src1.y)
1303 dst.z = max(src0.z, src1.z)
1305 dst.w = max(src0.w, src1.w)
1308 .. opcode:: IMIN - Minimum of Signed Integers
1312 dst.x = min(src0.x, src1.x)
1314 dst.y = min(src0.y, src1.y)
1316 dst.z = min(src0.z, src1.z)
1318 dst.w = min(src0.w, src1.w)
1321 .. opcode:: UMIN - Minimum of Unsigned Integers
1325 dst.x = min(src0.x, src1.x)
1327 dst.y = min(src0.y, src1.y)
1329 dst.z = min(src0.z, src1.z)
1331 dst.w = min(src0.w, src1.w)
1334 .. opcode:: SHL - Shift Left
1336 The shift count is masked with 0x1f before the shift is applied.
1340 dst.x = src0.x << (0x1f \& src1.x)
1342 dst.y = src0.y << (0x1f \& src1.y)
1344 dst.z = src0.z << (0x1f \& src1.z)
1346 dst.w = src0.w << (0x1f \& src1.w)
1349 .. opcode:: ISHR - Arithmetic Shift Right (of Signed Integer)
1351 The shift count is masked with 0x1f before the shift is applied.
1355 dst.x = src0.x >> (0x1f \& src1.x)
1357 dst.y = src0.y >> (0x1f \& src1.y)
1359 dst.z = src0.z >> (0x1f \& src1.z)
1361 dst.w = src0.w >> (0x1f \& src1.w)
1364 .. opcode:: USHR - Logical Shift Right
1366 The shift count is masked with 0x1f before the shift is applied.
1370 dst.x = src0.x >> (unsigned) (0x1f \& src1.x)
1372 dst.y = src0.y >> (unsigned) (0x1f \& src1.y)
1374 dst.z = src0.z >> (unsigned) (0x1f \& src1.z)
1376 dst.w = src0.w >> (unsigned) (0x1f \& src1.w)
1379 .. opcode:: UCMP - Integer Conditional Move
1383 dst.x = src0.x ? src1.x : src2.x
1385 dst.y = src0.y ? src1.y : src2.y
1387 dst.z = src0.z ? src1.z : src2.z
1389 dst.w = src0.w ? src1.w : src2.w
1393 .. opcode:: ISSG - Integer Set Sign
1397 dst.x = (src0.x < 0) ? -1 : (src0.x > 0) ? 1 : 0
1399 dst.y = (src0.y < 0) ? -1 : (src0.y > 0) ? 1 : 0
1401 dst.z = (src0.z < 0) ? -1 : (src0.z > 0) ? 1 : 0
1403 dst.w = (src0.w < 0) ? -1 : (src0.w > 0) ? 1 : 0
1407 .. opcode:: FSLT - Float Set On Less Than (ordered)
1409 Same comparison as SLT but returns integer instead of 1.0/0.0 float
1413 dst.x = (src0.x < src1.x) ? \sim 0 : 0
1415 dst.y = (src0.y < src1.y) ? \sim 0 : 0
1417 dst.z = (src0.z < src1.z) ? \sim 0 : 0
1419 dst.w = (src0.w < src1.w) ? \sim 0 : 0
1422 .. opcode:: ISLT - Signed Integer Set On Less Than
1426 dst.x = (src0.x < src1.x) ? \sim 0 : 0
1428 dst.y = (src0.y < src1.y) ? \sim 0 : 0
1430 dst.z = (src0.z < src1.z) ? \sim 0 : 0
1432 dst.w = (src0.w < src1.w) ? \sim 0 : 0
1435 .. opcode:: USLT - Unsigned Integer Set On Less Than
1439 dst.x = (src0.x < src1.x) ? \sim 0 : 0
1441 dst.y = (src0.y < src1.y) ? \sim 0 : 0
1443 dst.z = (src0.z < src1.z) ? \sim 0 : 0
1445 dst.w = (src0.w < src1.w) ? \sim 0 : 0
1448 .. opcode:: FSGE - Float Set On Greater Equal Than (ordered)
1450 Same comparison as SGE but returns integer instead of 1.0/0.0 float
1454 dst.x = (src0.x >= src1.x) ? \sim 0 : 0
1456 dst.y = (src0.y >= src1.y) ? \sim 0 : 0
1458 dst.z = (src0.z >= src1.z) ? \sim 0 : 0
1460 dst.w = (src0.w >= src1.w) ? \sim 0 : 0
1463 .. opcode:: ISGE - Signed Integer Set On Greater Equal Than
1467 dst.x = (src0.x >= src1.x) ? \sim 0 : 0
1469 dst.y = (src0.y >= src1.y) ? \sim 0 : 0
1471 dst.z = (src0.z >= src1.z) ? \sim 0 : 0
1473 dst.w = (src0.w >= src1.w) ? \sim 0 : 0
1476 .. opcode:: USGE - Unsigned Integer Set On Greater Equal Than
1480 dst.x = (src0.x >= src1.x) ? \sim 0 : 0
1482 dst.y = (src0.y >= src1.y) ? \sim 0 : 0
1484 dst.z = (src0.z >= src1.z) ? \sim 0 : 0
1486 dst.w = (src0.w >= src1.w) ? \sim 0 : 0
1489 .. opcode:: FSEQ - Float Set On Equal (ordered)
1491 Same comparison as SEQ but returns integer instead of 1.0/0.0 float
1495 dst.x = (src0.x == src1.x) ? \sim 0 : 0
1497 dst.y = (src0.y == src1.y) ? \sim 0 : 0
1499 dst.z = (src0.z == src1.z) ? \sim 0 : 0
1501 dst.w = (src0.w == src1.w) ? \sim 0 : 0
1504 .. opcode:: USEQ - Integer Set On Equal
1508 dst.x = (src0.x == src1.x) ? \sim 0 : 0
1510 dst.y = (src0.y == src1.y) ? \sim 0 : 0
1512 dst.z = (src0.z == src1.z) ? \sim 0 : 0
1514 dst.w = (src0.w == src1.w) ? \sim 0 : 0
1517 .. opcode:: FSNE - Float Set On Not Equal (unordered)
1519 Same comparison as SNE but returns integer instead of 1.0/0.0 float
1523 dst.x = (src0.x != src1.x) ? \sim 0 : 0
1525 dst.y = (src0.y != src1.y) ? \sim 0 : 0
1527 dst.z = (src0.z != src1.z) ? \sim 0 : 0
1529 dst.w = (src0.w != src1.w) ? \sim 0 : 0
1532 .. opcode:: USNE - Integer Set On Not Equal
1536 dst.x = (src0.x != src1.x) ? \sim 0 : 0
1538 dst.y = (src0.y != src1.y) ? \sim 0 : 0
1540 dst.z = (src0.z != src1.z) ? \sim 0 : 0
1542 dst.w = (src0.w != src1.w) ? \sim 0 : 0
1545 .. opcode:: INEG - Integer Negate
1560 .. opcode:: IABS - Integer Absolute Value
1574 These opcodes are used for bit-level manipulation of integers.
1576 .. opcode:: IBFE - Signed Bitfield Extract
1578 Like GLSL bitfieldExtract. Extracts a set of bits from the input, and
1579 sign-extends them if the high bit of the extracted window is set.
1583 def ibfe(value, offset, bits):
1584 if offset < 0 or bits < 0 or offset + bits > 32:
1586 if bits == 0: return 0
1587 # Note: >> sign-extends
1588 return (value << (32 - offset - bits)) >> (32 - bits)
1590 .. opcode:: UBFE - Unsigned Bitfield Extract
1592 Like GLSL bitfieldExtract. Extracts a set of bits from the input, without
1597 def ubfe(value, offset, bits):
1598 if offset < 0 or bits < 0 or offset + bits > 32:
1600 if bits == 0: return 0
1601 # Note: >> does not sign-extend
1602 return (value << (32 - offset - bits)) >> (32 - bits)
1604 .. opcode:: BFI - Bitfield Insert
1606 Like GLSL bitfieldInsert. Replaces a bit region of 'base' with the low bits
1611 def bfi(base, insert, offset, bits):
1612 if offset < 0 or bits < 0 or offset + bits > 32:
1614 # << defined such that mask == ~0 when bits == 32, offset == 0
1615 mask = ((1 << bits) - 1) << offset
1616 return ((insert << offset) & mask) | (base & ~mask)
1618 .. opcode:: BREV - Bitfield Reverse
1620 See SM5 instruction BFREV. Reverses the bits of the argument.
1622 .. opcode:: POPC - Population Count
1624 See SM5 instruction COUNTBITS. Counts the number of set bits in the argument.
1626 .. opcode:: LSB - Index of lowest set bit
1628 See SM5 instruction FIRSTBIT_LO. Computes the 0-based index of the first set
1629 bit of the argument. Returns -1 if none are set.
1631 .. opcode:: IMSB - Index of highest non-sign bit
1633 See SM5 instruction FIRSTBIT_SHI. Computes the 0-based index of the highest
1634 non-sign bit of the argument (i.e. highest 0 bit for negative numbers,
1635 highest 1 bit for positive numbers). Returns -1 if all bits are the same
1636 (i.e. for inputs 0 and -1).
1638 .. opcode:: UMSB - Index of highest set bit
1640 See SM5 instruction FIRSTBIT_HI. Computes the 0-based index of the highest
1641 set bit of the argument. Returns -1 if none are set.
1644 ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
1646 These opcodes are only supported in geometry shaders; they have no meaning
1647 in any other type of shader.
1649 .. opcode:: EMIT - Emit
1651 Generate a new vertex for the current primitive into the specified vertex
1652 stream using the values in the output registers.
1655 .. opcode:: ENDPRIM - End Primitive
1657 Complete the current primitive in the specified vertex stream (consisting of
1658 the emitted vertices), and start a new one.
1664 These opcodes are part of :term:`GLSL`'s opcode set. Support for these
1665 opcodes is determined by a special capability bit, ``GLSL``.
1666 Some require glsl version 1.30 (UIF/BREAKC/SWITCH/CASE/DEFAULT/ENDSWITCH).
1668 .. opcode:: CAL - Subroutine Call
1674 .. opcode:: RET - Subroutine Call Return
1679 .. opcode:: CONT - Continue
1681 Unconditionally moves the point of execution to the instruction after the
1682 last bgnloop. The instruction must appear within a bgnloop/endloop.
1686 Support for CONT is determined by a special capability bit,
1687 ``TGSI_CONT_SUPPORTED``. See :ref:`Screen` for more information.
1690 .. opcode:: BGNLOOP - Begin a Loop
1692 Start a loop. Must have a matching endloop.
1695 .. opcode:: BGNSUB - Begin Subroutine
1697 Starts definition of a subroutine. Must have a matching endsub.
1700 .. opcode:: ENDLOOP - End a Loop
1702 End a loop started with bgnloop.
1705 .. opcode:: ENDSUB - End Subroutine
1707 Ends definition of a subroutine.
1710 .. opcode:: NOP - No Operation
1715 .. opcode:: BRK - Break
1717 Unconditionally moves the point of execution to the instruction after the
1718 next endloop or endswitch. The instruction must appear within a loop/endloop
1719 or switch/endswitch.
1722 .. opcode:: BREAKC - Break Conditional
1724 Conditionally moves the point of execution to the instruction after the
1725 next endloop or endswitch. The instruction must appear within a loop/endloop
1726 or switch/endswitch.
1727 Condition evaluates to true if src0.x != 0 where src0.x is interpreted
1728 as an integer register.
1732 Considered for removal as it's quite inconsistent wrt other opcodes
1733 (could emulate with UIF/BRK/ENDIF).
1736 .. opcode:: IF - Float If
1738 Start an IF ... ELSE .. ENDIF block. Condition evaluates to true if
1742 where src0.x is interpreted as a floating point register.
1745 .. opcode:: UIF - Bitwise If
1747 Start an UIF ... ELSE .. ENDIF block. Condition evaluates to true if
1751 where src0.x is interpreted as an integer register.
1754 .. opcode:: ELSE - Else
1756 Starts an else block, after an IF or UIF statement.
1759 .. opcode:: ENDIF - End If
1761 Ends an IF or UIF block.
1764 .. opcode:: SWITCH - Switch
1766 Starts a C-style switch expression. The switch consists of one or multiple
1767 CASE statements, and at most one DEFAULT statement. Execution of a statement
1768 ends when a BRK is hit, but just like in C falling through to other cases
1769 without a break is allowed. Similarly, DEFAULT label is allowed anywhere not
1770 just as last statement, and fallthrough is allowed into/from it.
1771 CASE src arguments are evaluated at bit level against the SWITCH src argument.
1777 (some instructions here)
1780 (some instructions here)
1783 (some instructions here)
1788 .. opcode:: CASE - Switch case
1790 This represents a switch case label. The src arg must be an integer immediate.
1793 .. opcode:: DEFAULT - Switch default
1795 This represents the default case in the switch, which is taken if no other
1799 .. opcode:: ENDSWITCH - End of switch
1801 Ends a switch expression.
1807 The interpolation instructions allow an input to be interpolated in a
1808 different way than its declaration. This corresponds to the GLSL 4.00
1809 interpolateAt* functions. The first argument of each of these must come from
1810 ``TGSI_FILE_INPUT``.
1812 .. opcode:: INTERP_CENTROID - Interpolate at the centroid
1814 Interpolates the varying specified by src0 at the centroid
1816 .. opcode:: INTERP_SAMPLE - Interpolate at the specified sample
1818 Interpolates the varying specified by src0 at the sample id specified by
1819 src1.x (interpreted as an integer)
1821 .. opcode:: INTERP_OFFSET - Interpolate at the specified offset
1823 Interpolates the varying specified by src0 at the offset src1.xy from the
1824 pixel center (interpreted as floats)
1832 The double-precision opcodes reinterpret four-component vectors into
1833 two-component vectors with doubled precision in each component.
1835 .. opcode:: DABS - Absolute
1843 .. opcode:: DADD - Add
1847 dst.xy = src0.xy + src1.xy
1849 dst.zw = src0.zw + src1.zw
1851 .. opcode:: DSEQ - Set on Equal
1855 dst.x = src0.xy == src1.xy ? \sim 0 : 0
1857 dst.z = src0.zw == src1.zw ? \sim 0 : 0
1859 .. opcode:: DSNE - Set on Equal
1863 dst.x = src0.xy != src1.xy ? \sim 0 : 0
1865 dst.z = src0.zw != src1.zw ? \sim 0 : 0
1867 .. opcode:: DSLT - Set on Less than
1871 dst.x = src0.xy < src1.xy ? \sim 0 : 0
1873 dst.z = src0.zw < src1.zw ? \sim 0 : 0
1875 .. opcode:: DSGE - Set on Greater equal
1879 dst.x = src0.xy >= src1.xy ? \sim 0 : 0
1881 dst.z = src0.zw >= src1.zw ? \sim 0 : 0
1883 .. opcode:: DFRAC - Fraction
1887 dst.xy = src.xy - \lfloor src.xy\rfloor
1889 dst.zw = src.zw - \lfloor src.zw\rfloor
1891 .. opcode:: DTRUNC - Truncate
1895 dst.xy = trunc(src.xy)
1897 dst.zw = trunc(src.zw)
1899 .. opcode:: DCEIL - Ceiling
1903 dst.xy = \lceil src.xy\rceil
1905 dst.zw = \lceil src.zw\rceil
1907 .. opcode:: DFLR - Floor
1911 dst.xy = \lfloor src.xy\rfloor
1913 dst.zw = \lfloor src.zw\rfloor
1915 .. opcode:: DROUND - Fraction
1919 dst.xy = round(src.xy)
1921 dst.zw = round(src.zw)
1923 .. opcode:: DSSG - Set Sign
1927 dst.xy = (src.xy > 0) ? 1.0 : (src.xy < 0) ? -1.0 : 0.0
1929 dst.zw = (src.zw > 0) ? 1.0 : (src.zw < 0) ? -1.0 : 0.0
1931 .. opcode:: DFRACEXP - Convert Number to Fractional and Integral Components
1933 Like the ``frexp()`` routine in many math libraries, this opcode stores the
1934 exponent of its source to ``dst0``, and the significand to ``dst1``, such that
1935 :math:`dst1 \times 2^{dst0} = src` .
1939 dst0.xy = exp(src.xy)
1941 dst1.xy = frac(src.xy)
1943 dst0.zw = exp(src.zw)
1945 dst1.zw = frac(src.zw)
1947 .. opcode:: DLDEXP - Multiply Number by Integral Power of 2
1949 This opcode is the inverse of :opcode:`DFRACEXP`. The second
1950 source is an integer.
1954 dst.xy = src0.xy \times 2^{src1.x}
1956 dst.zw = src0.zw \times 2^{src1.y}
1958 .. opcode:: DMIN - Minimum
1962 dst.xy = min(src0.xy, src1.xy)
1964 dst.zw = min(src0.zw, src1.zw)
1966 .. opcode:: DMAX - Maximum
1970 dst.xy = max(src0.xy, src1.xy)
1972 dst.zw = max(src0.zw, src1.zw)
1974 .. opcode:: DMUL - Multiply
1978 dst.xy = src0.xy \times src1.xy
1980 dst.zw = src0.zw \times src1.zw
1983 .. opcode:: DMAD - Multiply And Add
1987 dst.xy = src0.xy \times src1.xy + src2.xy
1989 dst.zw = src0.zw \times src1.zw + src2.zw
1992 .. opcode:: DFMA - Fused Multiply-Add
1994 Perform a * b + c with no intermediate rounding step.
1998 dst.xy = src0.xy \times src1.xy + src2.xy
2000 dst.zw = src0.zw \times src1.zw + src2.zw
2003 .. opcode:: DDIV - Divide
2007 dst.xy = \frac{src0.xy}{src1.xy}
2009 dst.zw = \frac{src0.zw}{src1.zw}
2012 .. opcode:: DRCP - Reciprocal
2016 dst.xy = \frac{1}{src.xy}
2018 dst.zw = \frac{1}{src.zw}
2020 .. opcode:: DSQRT - Square Root
2024 dst.xy = \sqrt{src.xy}
2026 dst.zw = \sqrt{src.zw}
2028 .. opcode:: DRSQ - Reciprocal Square Root
2032 dst.xy = \frac{1}{\sqrt{src.xy}}
2034 dst.zw = \frac{1}{\sqrt{src.zw}}
2036 .. opcode:: F2D - Float to Double
2040 dst.xy = double(src0.x)
2042 dst.zw = double(src0.y)
2044 .. opcode:: D2F - Double to Float
2048 dst.x = float(src0.xy)
2050 dst.y = float(src0.zw)
2052 .. opcode:: I2D - Int to Double
2056 dst.xy = double(src0.x)
2058 dst.zw = double(src0.y)
2060 .. opcode:: D2I - Double to Int
2064 dst.x = int(src0.xy)
2066 dst.y = int(src0.zw)
2068 .. opcode:: U2D - Unsigned Int to Double
2072 dst.xy = double(src0.x)
2074 dst.zw = double(src0.y)
2076 .. opcode:: D2U - Double to Unsigned Int
2080 dst.x = unsigned(src0.xy)
2082 dst.y = unsigned(src0.zw)
2087 The 64-bit integer opcodes reinterpret four-component vectors into
2088 two-component vectors with 64-bits in each component.
2090 .. opcode:: I64ABS - 64-bit Integer Absolute Value
2098 .. opcode:: I64NEG - 64-bit Integer Negate
2108 .. opcode:: I64SSG - 64-bit Integer Set Sign
2112 dst.xy = (src0.xy < 0) ? -1 : (src0.xy > 0) ? 1 : 0
2114 dst.zw = (src0.zw < 0) ? -1 : (src0.zw > 0) ? 1 : 0
2116 .. opcode:: U64ADD - 64-bit Integer Add
2120 dst.xy = src0.xy + src1.xy
2122 dst.zw = src0.zw + src1.zw
2124 .. opcode:: U64MUL - 64-bit Integer Multiply
2128 dst.xy = src0.xy * src1.xy
2130 dst.zw = src0.zw * src1.zw
2132 .. opcode:: U64SEQ - 64-bit Integer Set on Equal
2136 dst.x = src0.xy == src1.xy ? \sim 0 : 0
2138 dst.z = src0.zw == src1.zw ? \sim 0 : 0
2140 .. opcode:: U64SNE - 64-bit Integer Set on Not Equal
2144 dst.x = src0.xy != src1.xy ? \sim 0 : 0
2146 dst.z = src0.zw != src1.zw ? \sim 0 : 0
2148 .. opcode:: U64SLT - 64-bit Unsigned Integer Set on Less Than
2152 dst.x = src0.xy < src1.xy ? \sim 0 : 0
2154 dst.z = src0.zw < src1.zw ? \sim 0 : 0
2156 .. opcode:: U64SGE - 64-bit Unsigned Integer Set on Greater Equal
2160 dst.x = src0.xy >= src1.xy ? \sim 0 : 0
2162 dst.z = src0.zw >= src1.zw ? \sim 0 : 0
2164 .. opcode:: I64SLT - 64-bit Signed Integer Set on Less Than
2168 dst.x = src0.xy < src1.xy ? \sim 0 : 0
2170 dst.z = src0.zw < src1.zw ? \sim 0 : 0
2172 .. opcode:: I64SGE - 64-bit Signed Integer Set on Greater Equal
2176 dst.x = src0.xy >= src1.xy ? \sim 0 : 0
2178 dst.z = src0.zw >= src1.zw ? \sim 0 : 0
2180 .. opcode:: I64MIN - Minimum of 64-bit Signed Integers
2184 dst.xy = min(src0.xy, src1.xy)
2186 dst.zw = min(src0.zw, src1.zw)
2188 .. opcode:: U64MIN - Minimum of 64-bit Unsigned Integers
2192 dst.xy = min(src0.xy, src1.xy)
2194 dst.zw = min(src0.zw, src1.zw)
2196 .. opcode:: I64MAX - Maximum of 64-bit Signed Integers
2200 dst.xy = max(src0.xy, src1.xy)
2202 dst.zw = max(src0.zw, src1.zw)
2204 .. opcode:: U64MAX - Maximum of 64-bit Unsigned Integers
2208 dst.xy = max(src0.xy, src1.xy)
2210 dst.zw = max(src0.zw, src1.zw)
2212 .. opcode:: U64SHL - Shift Left 64-bit Unsigned Integer
2214 The shift count is masked with 0x3f before the shift is applied.
2218 dst.xy = src0.xy << (0x3f \& src1.x)
2220 dst.zw = src0.zw << (0x3f \& src1.y)
2222 .. opcode:: I64SHR - Arithmetic Shift Right (of 64-bit Signed Integer)
2224 The shift count is masked with 0x3f before the shift is applied.
2228 dst.xy = src0.xy >> (0x3f \& src1.x)
2230 dst.zw = src0.zw >> (0x3f \& src1.y)
2232 .. opcode:: U64SHR - Logical Shift Right (of 64-bit Unsigned Integer)
2234 The shift count is masked with 0x3f before the shift is applied.
2238 dst.xy = src0.xy >> (unsigned) (0x3f \& src1.x)
2240 dst.zw = src0.zw >> (unsigned) (0x3f \& src1.y)
2242 .. opcode:: I64DIV - 64-bit Signed Integer Division
2246 dst.xy = src0.xy \ src1.xy
2248 dst.zw = src0.zw \ src1.zw
2250 .. opcode:: U64DIV - 64-bit Unsigned Integer Division
2254 dst.xy = src0.xy \ src1.xy
2256 dst.zw = src0.zw \ src1.zw
2258 .. opcode:: U64MOD - 64-bit Unsigned Integer Remainder
2262 dst.xy = src0.xy \bmod src1.xy
2264 dst.zw = src0.zw \bmod src1.zw
2266 .. opcode:: I64MOD - 64-bit Signed Integer Remainder
2270 dst.xy = src0.xy \bmod src1.xy
2272 dst.zw = src0.zw \bmod src1.zw
2274 .. opcode:: F2U64 - Float to 64-bit Unsigned Int
2278 dst.xy = (uint64_t) src0.x
2280 dst.zw = (uint64_t) src0.y
2282 .. opcode:: F2I64 - Float to 64-bit Int
2286 dst.xy = (int64_t) src0.x
2288 dst.zw = (int64_t) src0.y
2290 .. opcode:: U2I64 - Unsigned Integer to 64-bit Integer
2292 This is a zero extension.
2296 dst.xy = (uint64_t) src0.x
2298 dst.zw = (uint64_t) src0.y
2300 .. opcode:: I2I64 - Signed Integer to 64-bit Integer
2302 This is a sign extension.
2306 dst.xy = (int64_t) src0.x
2308 dst.zw = (int64_t) src0.y
2310 .. opcode:: D2U64 - Double to 64-bit Unsigned Int
2314 dst.xy = (uint64_t) src0.xy
2316 dst.zw = (uint64_t) src0.zw
2318 .. opcode:: D2I64 - Double to 64-bit Int
2322 dst.xy = (int64_t) src0.xy
2324 dst.zw = (int64_t) src0.zw
2326 .. opcode:: U642F - 64-bit unsigned integer to float
2330 dst.x = (float) src0.xy
2332 dst.y = (float) src0.zw
2334 .. opcode:: I642F - 64-bit Int to Float
2338 dst.x = (float) src0.xy
2340 dst.y = (float) src0.zw
2342 .. opcode:: U642D - 64-bit unsigned integer to double
2346 dst.xy = (double) src0.xy
2348 dst.zw = (double) src0.zw
2350 .. opcode:: I642D - 64-bit Int to double
2354 dst.xy = (double) src0.xy
2356 dst.zw = (double) src0.zw
2358 .. _samplingopcodes:
2360 Resource Sampling Opcodes
2361 ^^^^^^^^^^^^^^^^^^^^^^^^^
2363 Those opcodes follow very closely semantics of the respective Direct3D
2364 instructions. If in doubt double check Direct3D documentation.
2365 Note that the swizzle on SVIEW (src1) determines texel swizzling
2370 Using provided address, sample data from the specified texture using the
2371 filtering mode identified by the given sampler. The source data may come from
2372 any resource type other than buffers.
2374 Syntax: ``SAMPLE dst, address, sampler_view, sampler``
2376 Example: ``SAMPLE TEMP[0], TEMP[1], SVIEW[0], SAMP[0]``
2378 .. opcode:: SAMPLE_I
2380 Simplified alternative to the SAMPLE instruction. Using the provided
2381 integer address, SAMPLE_I fetches data from the specified sampler view
2382 without any filtering. The source data may come from any resource type
2385 Syntax: ``SAMPLE_I dst, address, sampler_view``
2387 Example: ``SAMPLE_I TEMP[0], TEMP[1], SVIEW[0]``
2389 The 'address' is specified as unsigned integers. If the 'address' is out of
2390 range [0...(# texels - 1)] the result of the fetch is always 0 in all
2391 components. As such the instruction doesn't honor address wrap modes, in
2392 cases where that behavior is desirable 'SAMPLE' instruction should be used.
2393 address.w always provides an unsigned integer mipmap level. If the value is
2394 out of the range then the instruction always returns 0 in all components.
2395 address.yz are ignored for buffers and 1d textures. address.z is ignored
2396 for 1d texture arrays and 2d textures.
2398 For 1D texture arrays address.y provides the array index (also as unsigned
2399 integer). If the value is out of the range of available array indices
2400 [0... (array size - 1)] then the opcode always returns 0 in all components.
2401 For 2D texture arrays address.z provides the array index, otherwise it
2402 exhibits the same behavior as in the case for 1D texture arrays. The exact
2403 semantics of the source address are presented in the table below:
2405 +---------------------------+----+-----+-----+---------+
2406 | resource type | X | Y | Z | W |
2407 +===========================+====+=====+=====+=========+
2408 | ``PIPE_BUFFER`` | x | | | ignored |
2409 +---------------------------+----+-----+-----+---------+
2410 | ``PIPE_TEXTURE_1D`` | x | | | mpl |
2411 +---------------------------+----+-----+-----+---------+
2412 | ``PIPE_TEXTURE_2D`` | x | y | | mpl |
2413 +---------------------------+----+-----+-----+---------+
2414 | ``PIPE_TEXTURE_3D`` | x | y | z | mpl |
2415 +---------------------------+----+-----+-----+---------+
2416 | ``PIPE_TEXTURE_RECT`` | x | y | | mpl |
2417 +---------------------------+----+-----+-----+---------+
2418 | ``PIPE_TEXTURE_CUBE`` | not allowed as source |
2419 +---------------------------+----+-----+-----+---------+
2420 | ``PIPE_TEXTURE_1D_ARRAY`` | x | idx | | mpl |
2421 +---------------------------+----+-----+-----+---------+
2422 | ``PIPE_TEXTURE_2D_ARRAY`` | x | y | idx | mpl |
2423 +---------------------------+----+-----+-----+---------+
2425 Where 'mpl' is a mipmap level and 'idx' is the array index.
2427 .. opcode:: SAMPLE_I_MS
2429 Just like SAMPLE_I but allows fetch data from multi-sampled surfaces.
2431 Syntax: ``SAMPLE_I_MS dst, address, sampler_view, sample``
2433 .. opcode:: SAMPLE_B
2435 Just like the SAMPLE instruction with the exception that an additional bias
2436 is applied to the level of detail computed as part of the instruction
2439 Syntax: ``SAMPLE_B dst, address, sampler_view, sampler, lod_bias``
2441 Example: ``SAMPLE_B TEMP[0], TEMP[1], SVIEW[0], SAMP[0], TEMP[2].x``
2443 .. opcode:: SAMPLE_C
2445 Similar to the SAMPLE instruction but it performs a comparison filter. The
2446 operands to SAMPLE_C are identical to SAMPLE, except that there is an
2447 additional float32 operand, reference value, which must be a register with
2448 single-component, or a scalar literal. SAMPLE_C makes the hardware use the
2449 current samplers compare_func (in pipe_sampler_state) to compare reference
2450 value against the red component value for the surce resource at each texel
2451 that the currently configured texture filter covers based on the provided
2454 Syntax: ``SAMPLE_C dst, address, sampler_view.r, sampler, ref_value``
2456 Example: ``SAMPLE_C TEMP[0], TEMP[1], SVIEW[0].r, SAMP[0], TEMP[2].x``
2458 .. opcode:: SAMPLE_C_LZ
2460 Same as SAMPLE_C, but LOD is 0 and derivatives are ignored. The LZ stands
2463 Syntax: ``SAMPLE_C_LZ dst, address, sampler_view.r, sampler, ref_value``
2465 Example: ``SAMPLE_C_LZ TEMP[0], TEMP[1], SVIEW[0].r, SAMP[0], TEMP[2].x``
2468 .. opcode:: SAMPLE_D
2470 SAMPLE_D is identical to the SAMPLE opcode except that the derivatives for
2471 the source address in the x direction and the y direction are provided by
2474 Syntax: ``SAMPLE_D dst, address, sampler_view, sampler, der_x, der_y``
2476 Example: ``SAMPLE_D TEMP[0], TEMP[1], SVIEW[0], SAMP[0], TEMP[2], TEMP[3]``
2478 .. opcode:: SAMPLE_L
2480 SAMPLE_L is identical to the SAMPLE opcode except that the LOD is provided
2481 directly as a scalar value, representing no anisotropy.
2483 Syntax: ``SAMPLE_L dst, address, sampler_view, sampler, explicit_lod``
2485 Example: ``SAMPLE_L TEMP[0], TEMP[1], SVIEW[0], SAMP[0], TEMP[2].x``
2489 Gathers the four texels to be used in a bi-linear filtering operation and
2490 packs them into a single register. Only works with 2D, 2D array, cubemaps,
2491 and cubemaps arrays. For 2D textures, only the addressing modes of the
2492 sampler and the top level of any mip pyramid are used. Set W to zero. It
2493 behaves like the SAMPLE instruction, but a filtered sample is not
2494 generated. The four samples that contribute to filtering are placed into
2495 xyzw in counter-clockwise order, starting with the (u,v) texture coordinate
2496 delta at the following locations (-, +), (+, +), (+, -), (-, -), where the
2497 magnitude of the deltas are half a texel.
2500 .. opcode:: SVIEWINFO
2502 Query the dimensions of a given sampler view. dst receives width, height,
2503 depth or array size and number of mipmap levels as int4. The dst can have a
2504 writemask which will specify what info is the caller interested in.
2506 Syntax: ``SVIEWINFO dst, src_mip_level, sampler_view``
2508 Example: ``SVIEWINFO TEMP[0], TEMP[1].x, SVIEW[0]``
2510 src_mip_level is an unsigned integer scalar. If it's out of range then
2511 returns 0 for width, height and depth/array size but the total number of
2512 mipmap is still returned correctly for the given sampler view. The returned
2513 width, height and depth values are for the mipmap level selected by the
2514 src_mip_level and are in the number of texels. For 1d texture array width
2515 is in dst.x, array size is in dst.y and dst.z is 0. The number of mipmaps is
2516 still in dst.w. In contrast to d3d10 resinfo, there's no way in the tgsi
2517 instruction encoding to specify the return type (float/rcpfloat/uint), hence
2518 always using uint. Also, unlike the SAMPLE instructions, the swizzle on src1
2519 resinfo allowing swizzling dst values is ignored (due to the interaction
2520 with rcpfloat modifier which requires some swizzle handling in the state
2523 .. opcode:: SAMPLE_POS
2525 Query the position of a given sample. dst receives float4 (x, y, 0, 0)
2526 indicated where the sample is located. If the resource is not a multi-sample
2527 resource and not a render target, the result is 0.
2529 .. opcode:: SAMPLE_INFO
2531 dst receives number of samples in x. If the resource is not a multi-sample
2532 resource and not a render target, the result is 0.
2535 .. _resourceopcodes:
2537 Resource Access Opcodes
2538 ^^^^^^^^^^^^^^^^^^^^^^^
2540 .. opcode:: LOAD - Fetch data from a shader buffer or image
2542 Syntax: ``LOAD dst, resource, address``
2544 Example: ``LOAD TEMP[0], BUFFER[0], TEMP[1]``
2546 Using the provided integer address, LOAD fetches data
2547 from the specified buffer or texture without any
2550 The 'address' is specified as a vector of unsigned
2551 integers. If the 'address' is out of range the result
2554 Only the first mipmap level of a resource can be read
2555 from using this instruction.
2557 For 1D or 2D texture arrays, the array index is
2558 provided as an unsigned integer in address.y or
2559 address.z, respectively. address.yz are ignored for
2560 buffers and 1D textures. address.z is ignored for 1D
2561 texture arrays and 2D textures. address.w is always
2564 A swizzle suffix may be added to the resource argument
2565 this will cause the resource data to be swizzled accordingly.
2567 .. opcode:: STORE - Write data to a shader resource
2569 Syntax: ``STORE resource, address, src``
2571 Example: ``STORE BUFFER[0], TEMP[0], TEMP[1]``
2573 Using the provided integer address, STORE writes data
2574 to the specified buffer or texture.
2576 The 'address' is specified as a vector of unsigned
2577 integers. If the 'address' is out of range the result
2580 Only the first mipmap level of a resource can be
2581 written to using this instruction.
2583 For 1D or 2D texture arrays, the array index is
2584 provided as an unsigned integer in address.y or
2585 address.z, respectively. address.yz are ignored for
2586 buffers and 1D textures. address.z is ignored for 1D
2587 texture arrays and 2D textures. address.w is always
2590 .. opcode:: RESQ - Query information about a resource
2592 Syntax: ``RESQ dst, resource``
2594 Example: ``RESQ TEMP[0], BUFFER[0]``
2596 Returns information about the buffer or image resource. For buffer
2597 resources, the size (in bytes) is returned in the x component. For
2598 image resources, .xyz will contain the width/height/layers of the
2599 image, while .w will contain the number of samples for multi-sampled
2602 .. opcode:: FBFETCH - Load data from framebuffer
2604 Syntax: ``FBFETCH dst, output``
2606 Example: ``FBFETCH TEMP[0], OUT[0]``
2608 This is only valid on ``COLOR`` semantic outputs. Returns the color
2609 of the current position in the framebuffer from before this fragment
2610 shader invocation. May return the same value from multiple calls for
2611 a particular output within a single invocation. Note that result may
2612 be undefined if a fragment is drawn multiple times without a blend
2616 .. _threadsyncopcodes:
2618 Inter-thread synchronization opcodes
2619 ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
2621 These opcodes are intended for communication between threads running
2622 within the same compute grid. For now they're only valid in compute
2625 .. opcode:: MFENCE - Memory fence
2627 Syntax: ``MFENCE resource``
2629 Example: ``MFENCE RES[0]``
2631 This opcode forces strong ordering between any memory access
2632 operations that affect the specified resource. This means that
2633 previous loads and stores (and only those) will be performed and
2634 visible to other threads before the program execution continues.
2637 .. opcode:: LFENCE - Load memory fence
2639 Syntax: ``LFENCE resource``
2641 Example: ``LFENCE RES[0]``
2643 Similar to MFENCE, but it only affects the ordering of memory loads.
2646 .. opcode:: SFENCE - Store memory fence
2648 Syntax: ``SFENCE resource``
2650 Example: ``SFENCE RES[0]``
2652 Similar to MFENCE, but it only affects the ordering of memory stores.
2655 .. opcode:: BARRIER - Thread group barrier
2659 This opcode suspends the execution of the current thread until all
2660 the remaining threads in the working group reach the same point of
2661 the program. Results are unspecified if any of the remaining
2662 threads terminates or never reaches an executed BARRIER instruction.
2664 .. opcode:: MEMBAR - Memory barrier
2668 This opcode waits for the completion of all memory accesses based on
2669 the type passed in. The type is an immediate bitfield with the following
2672 Bit 0: Shader storage buffers
2673 Bit 1: Atomic buffers
2675 Bit 3: Shared memory
2678 These may be passed in in any combination. An implementation is free to not
2679 distinguish between these as it sees fit. However these map to all the
2680 possibilities made available by GLSL.
2687 These opcodes provide atomic variants of some common arithmetic and
2688 logical operations. In this context atomicity means that another
2689 concurrent memory access operation that affects the same memory
2690 location is guaranteed to be performed strictly before or after the
2691 entire execution of the atomic operation. The resource may be a buffer
2692 or an image. In the case of an image, the offset works the same as for
2693 ``LOAD`` and ``STORE``, specified above. These atomic operations may
2694 only be used with 32-bit integer image formats.
2696 .. opcode:: ATOMUADD - Atomic integer addition
2698 Syntax: ``ATOMUADD dst, resource, offset, src``
2700 Example: ``ATOMUADD TEMP[0], BUFFER[0], TEMP[1], TEMP[2]``
2702 The following operation is performed atomically:
2706 dst_x = resource[offset]
2708 resource[offset] = dst_x + src_x
2711 .. opcode:: ATOMXCHG - Atomic exchange
2713 Syntax: ``ATOMXCHG dst, resource, offset, src``
2715 Example: ``ATOMXCHG TEMP[0], BUFFER[0], TEMP[1], TEMP[2]``
2717 The following operation is performed atomically:
2721 dst_x = resource[offset]
2723 resource[offset] = src_x
2726 .. opcode:: ATOMCAS - Atomic compare-and-exchange
2728 Syntax: ``ATOMCAS dst, resource, offset, cmp, src``
2730 Example: ``ATOMCAS TEMP[0], BUFFER[0], TEMP[1], TEMP[2], TEMP[3]``
2732 The following operation is performed atomically:
2736 dst_x = resource[offset]
2738 resource[offset] = (dst_x == cmp_x ? src_x : dst_x)
2741 .. opcode:: ATOMAND - Atomic bitwise And
2743 Syntax: ``ATOMAND dst, resource, offset, src``
2745 Example: ``ATOMAND TEMP[0], BUFFER[0], TEMP[1], TEMP[2]``
2747 The following operation is performed atomically:
2751 dst_x = resource[offset]
2753 resource[offset] = dst_x \& src_x
2756 .. opcode:: ATOMOR - Atomic bitwise Or
2758 Syntax: ``ATOMOR dst, resource, offset, src``
2760 Example: ``ATOMOR TEMP[0], BUFFER[0], TEMP[1], TEMP[2]``
2762 The following operation is performed atomically:
2766 dst_x = resource[offset]
2768 resource[offset] = dst_x | src_x
2771 .. opcode:: ATOMXOR - Atomic bitwise Xor
2773 Syntax: ``ATOMXOR dst, resource, offset, src``
2775 Example: ``ATOMXOR TEMP[0], BUFFER[0], TEMP[1], TEMP[2]``
2777 The following operation is performed atomically:
2781 dst_x = resource[offset]
2783 resource[offset] = dst_x \oplus src_x
2786 .. opcode:: ATOMUMIN - Atomic unsigned minimum
2788 Syntax: ``ATOMUMIN dst, resource, offset, src``
2790 Example: ``ATOMUMIN TEMP[0], BUFFER[0], TEMP[1], TEMP[2]``
2792 The following operation is performed atomically:
2796 dst_x = resource[offset]
2798 resource[offset] = (dst_x < src_x ? dst_x : src_x)
2801 .. opcode:: ATOMUMAX - Atomic unsigned maximum
2803 Syntax: ``ATOMUMAX dst, resource, offset, src``
2805 Example: ``ATOMUMAX TEMP[0], BUFFER[0], TEMP[1], TEMP[2]``
2807 The following operation is performed atomically:
2811 dst_x = resource[offset]
2813 resource[offset] = (dst_x > src_x ? dst_x : src_x)
2816 .. opcode:: ATOMIMIN - Atomic signed minimum
2818 Syntax: ``ATOMIMIN dst, resource, offset, src``
2820 Example: ``ATOMIMIN TEMP[0], BUFFER[0], TEMP[1], TEMP[2]``
2822 The following operation is performed atomically:
2826 dst_x = resource[offset]
2828 resource[offset] = (dst_x < src_x ? dst_x : src_x)
2831 .. opcode:: ATOMIMAX - Atomic signed maximum
2833 Syntax: ``ATOMIMAX dst, resource, offset, src``
2835 Example: ``ATOMIMAX TEMP[0], BUFFER[0], TEMP[1], TEMP[2]``
2837 The following operation is performed atomically:
2841 dst_x = resource[offset]
2843 resource[offset] = (dst_x > src_x ? dst_x : src_x)
2851 These opcodes compare the given value across the shader invocations
2852 running in the current SIMD group. The details of exactly which
2853 invocations get compared are implementation-defined, and it would be a
2854 correct implementation to only ever consider the current thread's
2855 value. (i.e. SIMD group of 1). The argument is treated as a boolean.
2857 .. opcode:: VOTE_ANY - Value is set in any of the current invocations
2859 .. opcode:: VOTE_ALL - Value is set in all of the current invocations
2861 .. opcode:: VOTE_EQ - Value is the same in all of the current invocations
2864 Explanation of symbols used
2865 ------------------------------
2872 :math:`|x|` Absolute value of `x`.
2874 :math:`\lceil x \rceil` Ceiling of `x`.
2876 clamp(x,y,z) Clamp x between y and z.
2877 (x < y) ? y : (x > z) ? z : x
2879 :math:`\lfloor x\rfloor` Floor of `x`.
2881 :math:`\log_2{x}` Logarithm of `x`, base 2.
2883 max(x,y) Maximum of x and y.
2886 min(x,y) Minimum of x and y.
2889 partialx(x) Derivative of x relative to fragment's X.
2891 partialy(x) Derivative of x relative to fragment's Y.
2893 pop() Pop from stack.
2895 :math:`x^y` `x` to the power `y`.
2897 push(x) Push x on stack.
2901 trunc(x) Truncate x, i.e. drop the fraction bits.
2908 discard Discard fragment.
2912 target Label of target instruction.
2923 Declares a register that is will be referenced as an operand in Instruction
2926 File field contains register file that is being declared and is one
2929 UsageMask field specifies which of the register components can be accessed
2930 and is one of TGSI_WRITEMASK.
2932 The Local flag specifies that a given value isn't intended for
2933 subroutine parameter passing and, as a result, the implementation
2934 isn't required to give any guarantees of it being preserved across
2935 subroutine boundaries. As it's merely a compiler hint, the
2936 implementation is free to ignore it.
2938 If Dimension flag is set to 1, a Declaration Dimension token follows.
2940 If Semantic flag is set to 1, a Declaration Semantic token follows.
2942 If Interpolate flag is set to 1, a Declaration Interpolate token follows.
2944 If file is TGSI_FILE_RESOURCE, a Declaration Resource token follows.
2946 If Array flag is set to 1, a Declaration Array token follows.
2949 ^^^^^^^^^^^^^^^^^^^^^^^^
2951 Declarations can optional have an ArrayID attribute which can be referred by
2952 indirect addressing operands. An ArrayID of zero is reserved and treated as
2953 if no ArrayID is specified.
2955 If an indirect addressing operand refers to a specific declaration by using
2956 an ArrayID only the registers in this declaration are guaranteed to be
2957 accessed, accessing any register outside this declaration results in undefined
2958 behavior. Note that for compatibility the effective index is zero-based and
2959 not relative to the specified declaration
2961 If no ArrayID is specified with an indirect addressing operand the whole
2962 register file might be accessed by this operand. This is strongly discouraged
2963 and will prevent packing of scalar/vec2 arrays and effective alias analysis.
2964 This is only legal for TEMP and CONST register files.
2966 Declaration Semantic
2967 ^^^^^^^^^^^^^^^^^^^^^^^^
2969 Vertex and fragment shader input and output registers may be labeled
2970 with semantic information consisting of a name and index.
2972 Follows Declaration token if Semantic bit is set.
2974 Since its purpose is to link a shader with other stages of the pipeline,
2975 it is valid to follow only those Declaration tokens that declare a register
2976 either in INPUT or OUTPUT file.
2978 SemanticName field contains the semantic name of the register being declared.
2979 There is no default value.
2981 SemanticIndex is an optional subscript that can be used to distinguish
2982 different register declarations with the same semantic name. The default value
2985 The meanings of the individual semantic names are explained in the following
2988 TGSI_SEMANTIC_POSITION
2989 """"""""""""""""""""""
2991 For vertex shaders, TGSI_SEMANTIC_POSITION indicates the vertex shader
2992 output register which contains the homogeneous vertex position in the clip
2993 space coordinate system. After clipping, the X, Y and Z components of the
2994 vertex will be divided by the W value to get normalized device coordinates.
2996 For fragment shaders, TGSI_SEMANTIC_POSITION is used to indicate that
2997 fragment shader input (or system value, depending on which one is
2998 supported by the driver) contains the fragment's window position. The X
2999 component starts at zero and always increases from left to right.
3000 The Y component starts at zero and always increases but Y=0 may either
3001 indicate the top of the window or the bottom depending on the fragment
3002 coordinate origin convention (see TGSI_PROPERTY_FS_COORD_ORIGIN).
3003 The Z coordinate ranges from 0 to 1 to represent depth from the front
3004 to the back of the Z buffer. The W component contains the interpolated
3005 reciprocal of the vertex position W component (corresponding to gl_Fragcoord,
3006 but unlike d3d10 which interpolates the same 1/w but then gives back
3007 the reciprocal of the interpolated value).
3009 Fragment shaders may also declare an output register with
3010 TGSI_SEMANTIC_POSITION. Only the Z component is writable. This allows
3011 the fragment shader to change the fragment's Z position.
3018 For vertex shader outputs or fragment shader inputs/outputs, this
3019 label indicates that the register contains an R,G,B,A color.
3021 Several shader inputs/outputs may contain colors so the semantic index
3022 is used to distinguish them. For example, color[0] may be the diffuse
3023 color while color[1] may be the specular color.
3025 This label is needed so that the flat/smooth shading can be applied
3026 to the right interpolants during rasterization.
3030 TGSI_SEMANTIC_BCOLOR
3031 """"""""""""""""""""
3033 Back-facing colors are only used for back-facing polygons, and are only valid
3034 in vertex shader outputs. After rasterization, all polygons are front-facing
3035 and COLOR and BCOLOR end up occupying the same slots in the fragment shader,
3036 so all BCOLORs effectively become regular COLORs in the fragment shader.
3042 Vertex shader inputs and outputs and fragment shader inputs may be
3043 labeled with TGSI_SEMANTIC_FOG to indicate that the register contains
3044 a fog coordinate. Typically, the fragment shader will use the fog coordinate
3045 to compute a fog blend factor which is used to blend the normal fragment color
3046 with a constant fog color. But fog coord really is just an ordinary vec4
3047 register like regular semantics.
3053 Vertex shader input and output registers may be labeled with
3054 TGIS_SEMANTIC_PSIZE to indicate that the register contains a point size
3055 in the form (S, 0, 0, 1). The point size controls the width or diameter
3056 of points for rasterization. This label cannot be used in fragment
3059 When using this semantic, be sure to set the appropriate state in the
3060 :ref:`rasterizer` first.
3063 TGSI_SEMANTIC_TEXCOORD
3064 """"""""""""""""""""""
3066 Only available if PIPE_CAP_TGSI_TEXCOORD is exposed !
3068 Vertex shader outputs and fragment shader inputs may be labeled with
3069 this semantic to make them replaceable by sprite coordinates via the
3070 sprite_coord_enable state in the :ref:`rasterizer`.
3071 The semantic index permitted with this semantic is limited to <= 7.
3073 If the driver does not support TEXCOORD, sprite coordinate replacement
3074 applies to inputs with the GENERIC semantic instead.
3076 The intended use case for this semantic is gl_TexCoord.
3079 TGSI_SEMANTIC_PCOORD
3080 """"""""""""""""""""
3082 Only available if PIPE_CAP_TGSI_TEXCOORD is exposed !
3084 Fragment shader inputs may be labeled with TGSI_SEMANTIC_PCOORD to indicate
3085 that the register contains sprite coordinates in the form (x, y, 0, 1), if
3086 the current primitive is a point and point sprites are enabled. Otherwise,
3087 the contents of the register are undefined.
3089 The intended use case for this semantic is gl_PointCoord.
3092 TGSI_SEMANTIC_GENERIC
3093 """""""""""""""""""""
3095 All vertex/fragment shader inputs/outputs not labeled with any other
3096 semantic label can be considered to be generic attributes. Typical
3097 uses of generic inputs/outputs are texcoords and user-defined values.
3100 TGSI_SEMANTIC_NORMAL
3101 """"""""""""""""""""
3103 Indicates that a vertex shader input is a normal vector. This is
3104 typically only used for legacy graphics APIs.
3110 This label applies to fragment shader inputs (or system values,
3111 depending on which one is supported by the driver) and indicates that
3112 the register contains front/back-face information.
3114 If it is an input, it will be a floating-point vector in the form (F, 0, 0, 1),
3115 where F will be positive when the fragment belongs to a front-facing polygon,
3116 and negative when the fragment belongs to a back-facing polygon.
3118 If it is a system value, it will be an integer vector in the form (F, 0, 0, 1),
3119 where F is 0xffffffff when the fragment belongs to a front-facing polygon and
3120 0 when the fragment belongs to a back-facing polygon.
3123 TGSI_SEMANTIC_EDGEFLAG
3124 """"""""""""""""""""""
3126 For vertex shaders, this sematic label indicates that an input or
3127 output is a boolean edge flag. The register layout is [F, x, x, x]
3128 where F is 0.0 or 1.0 and x = don't care. Normally, the vertex shader
3129 simply copies the edge flag input to the edgeflag output.
3131 Edge flags are used to control which lines or points are actually
3132 drawn when the polygon mode converts triangles/quads/polygons into
3136 TGSI_SEMANTIC_STENCIL
3137 """""""""""""""""""""
3139 For fragment shaders, this semantic label indicates that an output
3140 is a writable stencil reference value. Only the Y component is writable.
3141 This allows the fragment shader to change the fragments stencilref value.
3144 TGSI_SEMANTIC_VIEWPORT_INDEX
3145 """"""""""""""""""""""""""""
3147 For geometry shaders, this semantic label indicates that an output
3148 contains the index of the viewport (and scissor) to use.
3149 This is an integer value, and only the X component is used.
3155 For geometry shaders, this semantic label indicates that an output
3156 contains the layer value to use for the color and depth/stencil surfaces.
3157 This is an integer value, and only the X component is used.
3158 (Also known as rendertarget array index.)
3161 TGSI_SEMANTIC_CULLDIST
3162 """"""""""""""""""""""
3164 Used as distance to plane for performing application-defined culling
3165 of individual primitives against a plane. When components of vertex
3166 elements are given this label, these values are assumed to be a
3167 float32 signed distance to a plane. Primitives will be completely
3168 discarded if the plane distance for all of the vertices in the
3169 primitive are < 0. If a vertex has a cull distance of NaN, that
3170 vertex counts as "out" (as if its < 0);
3171 The limits on both clip and cull distances are bound
3172 by the PIPE_MAX_CLIP_OR_CULL_DISTANCE_COUNT define which defines
3173 the maximum number of components that can be used to hold the
3174 distances and by the PIPE_MAX_CLIP_OR_CULL_DISTANCE_ELEMENT_COUNT
3175 which specifies the maximum number of registers which can be
3176 annotated with those semantics.
3179 TGSI_SEMANTIC_CLIPDIST
3180 """"""""""""""""""""""
3182 Note this covers clipping and culling distances.
3184 When components of vertex elements are identified this way, these
3185 values are each assumed to be a float32 signed distance to a plane.
3188 Primitive setup only invokes rasterization on pixels for which
3189 the interpolated plane distances are >= 0.
3192 Primitives will be completely discarded if the plane distance
3193 for all of the vertices in the primitive are < 0.
3194 If a vertex has a cull distance of NaN, that vertex counts as "out"
3197 Multiple clip/cull planes can be implemented simultaneously, by
3198 annotating multiple components of one or more vertex elements with
3199 the above specified semantic.
3200 The limits on both clip and cull distances are bound
3201 by the PIPE_MAX_CLIP_OR_CULL_DISTANCE_COUNT define which defines
3202 the maximum number of components that can be used to hold the
3203 distances and by the PIPE_MAX_CLIP_OR_CULL_DISTANCE_ELEMENT_COUNT
3204 which specifies the maximum number of registers which can be
3205 annotated with those semantics.
3206 The properties NUM_CLIPDIST_ENABLED and NUM_CULLDIST_ENABLED
3207 are used to divide up the 2 x vec4 space between clipping and culling.
3209 TGSI_SEMANTIC_SAMPLEID
3210 """"""""""""""""""""""
3212 For fragment shaders, this semantic label indicates that a system value
3213 contains the current sample id (i.e. gl_SampleID).
3214 This is an integer value, and only the X component is used.
3216 TGSI_SEMANTIC_SAMPLEPOS
3217 """""""""""""""""""""""
3219 For fragment shaders, this semantic label indicates that a system value
3220 contains the current sample's position (i.e. gl_SamplePosition). Only the X
3221 and Y values are used.
3223 TGSI_SEMANTIC_SAMPLEMASK
3224 """"""""""""""""""""""""
3226 For fragment shaders, this semantic label indicates that an output contains
3227 the sample mask used to disable further sample processing
3228 (i.e. gl_SampleMask). Only the X value is used, up to 32x MS.
3230 TGSI_SEMANTIC_INVOCATIONID
3231 """"""""""""""""""""""""""
3233 For geometry shaders, this semantic label indicates that a system value
3234 contains the current invocation id (i.e. gl_InvocationID).
3235 This is an integer value, and only the X component is used.
3237 TGSI_SEMANTIC_INSTANCEID
3238 """"""""""""""""""""""""
3240 For vertex shaders, this semantic label indicates that a system value contains
3241 the current instance id (i.e. gl_InstanceID). It does not include the base
3242 instance. This is an integer value, and only the X component is used.
3244 TGSI_SEMANTIC_VERTEXID
3245 """"""""""""""""""""""
3247 For vertex shaders, this semantic label indicates that a system value contains
3248 the current vertex id (i.e. gl_VertexID). It does (unlike in d3d10) include the
3249 base vertex. This is an integer value, and only the X component is used.
3251 TGSI_SEMANTIC_VERTEXID_NOBASE
3252 """""""""""""""""""""""""""""""
3254 For vertex shaders, this semantic label indicates that a system value contains
3255 the current vertex id without including the base vertex (this corresponds to
3256 d3d10 vertex id, so TGSI_SEMANTIC_VERTEXID_NOBASE + TGSI_SEMANTIC_BASEVERTEX
3257 == TGSI_SEMANTIC_VERTEXID). This is an integer value, and only the X component
3260 TGSI_SEMANTIC_BASEVERTEX
3261 """"""""""""""""""""""""
3263 For vertex shaders, this semantic label indicates that a system value contains
3264 the base vertex (i.e. gl_BaseVertex). Note that for non-indexed draw calls,
3265 this contains the first (or start) value instead.
3266 This is an integer value, and only the X component is used.
3268 TGSI_SEMANTIC_PRIMID
3269 """"""""""""""""""""
3271 For geometry and fragment shaders, this semantic label indicates the value
3272 contains the primitive id (i.e. gl_PrimitiveID). This is an integer value,
3273 and only the X component is used.
3274 FIXME: This right now can be either a ordinary input or a system value...
3280 For tessellation evaluation/control shaders, this semantic label indicates a
3281 generic per-patch attribute. Such semantics will not implicitly be per-vertex
3284 TGSI_SEMANTIC_TESSCOORD
3285 """""""""""""""""""""""
3287 For tessellation evaluation shaders, this semantic label indicates the
3288 coordinates of the vertex being processed. This is available in XYZ; W is
3291 TGSI_SEMANTIC_TESSOUTER
3292 """""""""""""""""""""""
3294 For tessellation evaluation/control shaders, this semantic label indicates the
3295 outer tessellation levels of the patch. Isoline tessellation will only have XY
3296 defined, triangle will have XYZ and quads will have XYZW defined. This
3297 corresponds to gl_TessLevelOuter.
3299 TGSI_SEMANTIC_TESSINNER
3300 """""""""""""""""""""""
3302 For tessellation evaluation/control shaders, this semantic label indicates the
3303 inner tessellation levels of the patch. The X value is only defined for
3304 triangle tessellation, while quads will have XY defined. This is entirely
3305 undefined for isoline tessellation.
3307 TGSI_SEMANTIC_VERTICESIN
3308 """"""""""""""""""""""""
3310 For tessellation evaluation/control shaders, this semantic label indicates the
3311 number of vertices provided in the input patch. Only the X value is defined.
3313 TGSI_SEMANTIC_HELPER_INVOCATION
3314 """""""""""""""""""""""""""""""
3316 For fragment shaders, this semantic indicates whether the current
3317 invocation is covered or not. Helper invocations are created in order
3318 to properly compute derivatives, however it may be desirable to skip
3319 some of the logic in those cases. See ``gl_HelperInvocation`` documentation.
3321 TGSI_SEMANTIC_BASEINSTANCE
3322 """"""""""""""""""""""""""
3324 For vertex shaders, the base instance argument supplied for this
3325 draw. This is an integer value, and only the X component is used.
3327 TGSI_SEMANTIC_DRAWID
3328 """"""""""""""""""""
3330 For vertex shaders, the zero-based index of the current draw in a
3331 ``glMultiDraw*`` invocation. This is an integer value, and only the X
3335 TGSI_SEMANTIC_WORK_DIM
3336 """"""""""""""""""""""
3338 For compute shaders started via opencl this retrieves the work_dim
3339 parameter to the clEnqueueNDRangeKernel call with which the shader
3343 TGSI_SEMANTIC_GRID_SIZE
3344 """""""""""""""""""""""
3346 For compute shaders, this semantic indicates the maximum (x, y, z) dimensions
3347 of a grid of thread blocks.
3350 TGSI_SEMANTIC_BLOCK_ID
3351 """"""""""""""""""""""
3353 For compute shaders, this semantic indicates the (x, y, z) coordinates of the
3354 current block inside of the grid.
3357 TGSI_SEMANTIC_BLOCK_SIZE
3358 """"""""""""""""""""""""
3360 For compute shaders, this semantic indicates the maximum (x, y, z) dimensions
3361 of a block in threads.
3364 TGSI_SEMANTIC_THREAD_ID
3365 """""""""""""""""""""""
3367 For compute shaders, this semantic indicates the (x, y, z) coordinates of the
3368 current thread inside of the block.
3371 Declaration Interpolate
3372 ^^^^^^^^^^^^^^^^^^^^^^^
3374 This token is only valid for fragment shader INPUT declarations.
3376 The Interpolate field specifes the way input is being interpolated by
3377 the rasteriser and is one of TGSI_INTERPOLATE_*.
3379 The Location field specifies the location inside the pixel that the
3380 interpolation should be done at, one of ``TGSI_INTERPOLATE_LOC_*``. Note that
3381 when per-sample shading is enabled, the implementation may choose to
3382 interpolate at the sample irrespective of the Location field.
3384 The CylindricalWrap bitfield specifies which register components
3385 should be subject to cylindrical wrapping when interpolating by the
3386 rasteriser. If TGSI_CYLINDRICAL_WRAP_X is set to 1, the X component
3387 should be interpolated according to cylindrical wrapping rules.
3390 Declaration Sampler View
3391 ^^^^^^^^^^^^^^^^^^^^^^^^
3393 Follows Declaration token if file is TGSI_FILE_SAMPLER_VIEW.
3395 DCL SVIEW[#], resource, type(s)
3397 Declares a shader input sampler view and assigns it to a SVIEW[#]
3400 resource can be one of BUFFER, 1D, 2D, 3D, 1DArray and 2DArray.
3402 type must be 1 or 4 entries (if specifying on a per-component
3403 level) out of UNORM, SNORM, SINT, UINT and FLOAT.
3405 For TEX\* style texture sample opcodes (as opposed to SAMPLE\* opcodes
3406 which take an explicit SVIEW[#] source register), there may be optionally
3407 SVIEW[#] declarations. In this case, the SVIEW index is implied by the
3408 SAMP index, and there must be a corresponding SVIEW[#] declaration for
3409 each SAMP[#] declaration. Drivers are free to ignore this if they wish.
3410 But note in particular that some drivers need to know the sampler type
3411 (float/int/unsigned) in order to generate the correct code, so cases
3412 where integer textures are sampled, SVIEW[#] declarations should be
3415 NOTE: It is NOT legal to mix SAMPLE\* style opcodes and TEX\* opcodes
3418 Declaration Resource
3419 ^^^^^^^^^^^^^^^^^^^^
3421 Follows Declaration token if file is TGSI_FILE_RESOURCE.
3423 DCL RES[#], resource [, WR] [, RAW]
3425 Declares a shader input resource and assigns it to a RES[#]
3428 resource can be one of BUFFER, 1D, 2D, 3D, CUBE, 1DArray and
3431 If the RAW keyword is not specified, the texture data will be
3432 subject to conversion, swizzling and scaling as required to yield
3433 the specified data type from the physical data format of the bound
3436 If the RAW keyword is specified, no channel conversion will be
3437 performed: the values read for each of the channels (X,Y,Z,W) will
3438 correspond to consecutive words in the same order and format
3439 they're found in memory. No element-to-address conversion will be
3440 performed either: the value of the provided X coordinate will be
3441 interpreted in byte units instead of texel units. The result of
3442 accessing a misaligned address is undefined.
3444 Usage of the STORE opcode is only allowed if the WR (writable) flag
3449 ^^^^^^^^^^^^^^^^^^^^^^^^
3451 Properties are general directives that apply to the whole TGSI program.
3456 Specifies the fragment shader TGSI_SEMANTIC_POSITION coordinate origin.
3457 The default value is UPPER_LEFT.
3459 If UPPER_LEFT, the position will be (0,0) at the upper left corner and
3460 increase downward and rightward.
3461 If LOWER_LEFT, the position will be (0,0) at the lower left corner and
3462 increase upward and rightward.
3464 OpenGL defaults to LOWER_LEFT, and is configurable with the
3465 GL_ARB_fragment_coord_conventions extension.
3467 DirectX 9/10 use UPPER_LEFT.
3469 FS_COORD_PIXEL_CENTER
3470 """""""""""""""""""""
3472 Specifies the fragment shader TGSI_SEMANTIC_POSITION pixel center convention.
3473 The default value is HALF_INTEGER.
3475 If HALF_INTEGER, the fractionary part of the position will be 0.5
3476 If INTEGER, the fractionary part of the position will be 0.0
3478 Note that this does not affect the set of fragments generated by
3479 rasterization, which is instead controlled by half_pixel_center in the
3482 OpenGL defaults to HALF_INTEGER, and is configurable with the
3483 GL_ARB_fragment_coord_conventions extension.
3485 DirectX 9 uses INTEGER.
3486 DirectX 10 uses HALF_INTEGER.
3488 FS_COLOR0_WRITES_ALL_CBUFS
3489 """"""""""""""""""""""""""
3490 Specifies that writes to the fragment shader color 0 are replicated to all
3491 bound cbufs. This facilitates OpenGL's fragColor output vs fragData[0] where
3492 fragData is directed to a single color buffer, but fragColor is broadcast.
3495 """"""""""""""""""""""""""
3496 If this property is set on the program bound to the shader stage before the
3497 fragment shader, user clip planes should have no effect (be disabled) even if
3498 that shader does not write to any clip distance outputs and the rasterizer's
3499 clip_plane_enable is non-zero.
3500 This property is only supported by drivers that also support shader clip
3502 This is useful for APIs that don't have UCPs and where clip distances written
3503 by a shader cannot be disabled.
3508 Specifies the number of times a geometry shader should be executed for each
3509 input primitive. Each invocation will have a different
3510 TGSI_SEMANTIC_INVOCATIONID system value set. If not specified, assumed to
3513 VS_WINDOW_SPACE_POSITION
3514 """"""""""""""""""""""""""
3515 If this property is set on the vertex shader, the TGSI_SEMANTIC_POSITION output
3516 is assumed to contain window space coordinates.
3517 Division of X,Y,Z by W and the viewport transformation are disabled, and 1/W is
3518 directly taken from the 4-th component of the shader output.
3519 Naturally, clipping is not performed on window coordinates either.
3520 The effect of this property is undefined if a geometry or tessellation shader
3526 The number of vertices written by the tessellation control shader. This
3527 effectively defines the patch input size of the tessellation evaluation shader
3533 This sets the tessellation primitive mode, one of ``PIPE_PRIM_TRIANGLES``,
3534 ``PIPE_PRIM_QUADS``, or ``PIPE_PRIM_LINES``. (Unlike in GL, there is no
3535 separate isolines settings, the regular lines is assumed to mean isolines.)
3540 This sets the spacing mode of the tessellation generator, one of
3541 ``PIPE_TESS_SPACING_*``.
3546 This sets the vertex order to be clockwise if the value is 1, or
3547 counter-clockwise if set to 0.
3552 If set to a non-zero value, this turns on point mode for the tessellator,
3553 which means that points will be generated instead of primitives.
3555 NUM_CLIPDIST_ENABLED
3556 """"""""""""""""""""
3558 How many clip distance scalar outputs are enabled.
3560 NUM_CULLDIST_ENABLED
3561 """"""""""""""""""""
3563 How many cull distance scalar outputs are enabled.
3565 FS_EARLY_DEPTH_STENCIL
3566 """"""""""""""""""""""
3568 Whether depth test, stencil test, and occlusion query should run before
3569 the fragment shader (regardless of fragment shader side effects). Corresponds
3570 to GLSL early_fragment_tests.
3575 Which shader stage will MOST LIKELY follow after this shader when the shader
3576 is bound. This is only a hint to the driver and doesn't have to be precise.
3577 Only set for VS and TES.
3579 CS_FIXED_BLOCK_WIDTH / HEIGHT / DEPTH
3580 """""""""""""""""""""""""""""""""""""
3582 Threads per block in each dimension, if known at compile time. If the block size
3583 is known all three should be at least 1. If it is unknown they should all be set
3589 The MUL TGSI operation (FP32 multiplication) will return 0 if either
3590 of the operands are equal to 0. That means that 0 * Inf = 0. This
3591 should be set the same way for an entire pipeline. Note that this
3592 applies not only to the literal MUL TGSI opcode, but all FP32
3593 multiplications implied by other operations, such as MAD, FMA, DP2,
3594 DP3, DP4, DPH, DST, LOG, LRP, XPD, and possibly others. If there is a
3595 mismatch between shaders, then it is unspecified whether this behavior
3599 Texture Sampling and Texture Formats
3600 ------------------------------------
3602 This table shows how texture image components are returned as (x,y,z,w) tuples
3603 by TGSI texture instructions, such as :opcode:`TEX`, :opcode:`TXD`, and
3604 :opcode:`TXP`. For reference, OpenGL and Direct3D conventions are shown as
3607 +--------------------+--------------+--------------------+--------------+
3608 | Texture Components | Gallium | OpenGL | Direct3D 9 |
3609 +====================+==============+====================+==============+
3610 | R | (r, 0, 0, 1) | (r, 0, 0, 1) | (r, 1, 1, 1) |
3611 +--------------------+--------------+--------------------+--------------+
3612 | RG | (r, g, 0, 1) | (r, g, 0, 1) | (r, g, 1, 1) |
3613 +--------------------+--------------+--------------------+--------------+
3614 | RGB | (r, g, b, 1) | (r, g, b, 1) | (r, g, b, 1) |
3615 +--------------------+--------------+--------------------+--------------+
3616 | RGBA | (r, g, b, a) | (r, g, b, a) | (r, g, b, a) |
3617 +--------------------+--------------+--------------------+--------------+
3618 | A | (0, 0, 0, a) | (0, 0, 0, a) | (0, 0, 0, a) |
3619 +--------------------+--------------+--------------------+--------------+
3620 | L | (l, l, l, 1) | (l, l, l, 1) | (l, l, l, 1) |
3621 +--------------------+--------------+--------------------+--------------+
3622 | LA | (l, l, l, a) | (l, l, l, a) | (l, l, l, a) |
3623 +--------------------+--------------+--------------------+--------------+
3624 | I | (i, i, i, i) | (i, i, i, i) | N/A |
3625 +--------------------+--------------+--------------------+--------------+
3626 | UV | XXX TBD | (0, 0, 0, 1) | (u, v, 1, 1) |
3627 | | | [#envmap-bumpmap]_ | |
3628 +--------------------+--------------+--------------------+--------------+
3629 | Z | XXX TBD | (z, z, z, 1) | (0, z, 0, 1) |
3630 | | | [#depth-tex-mode]_ | |
3631 +--------------------+--------------+--------------------+--------------+
3632 | S | (s, s, s, s) | unknown | unknown |
3633 +--------------------+--------------+--------------------+--------------+
3635 .. [#envmap-bumpmap] http://www.opengl.org/registry/specs/ATI/envmap_bumpmap.txt
3636 .. [#depth-tex-mode] the default is (z, z, z, 1) but may also be (0, 0, 0, z)
3637 or (z, z, z, z) depending on the value of GL_DEPTH_TEXTURE_MODE.