dst.w = src0.w \times src1.w + src2.w
-.. opcode:: SUB - Subtract
-
-.. math::
-
- dst.x = src0.x - src1.x
-
- dst.y = src0.y - src1.y
-
- dst.z = src0.z - src1.z
-
- dst.w = src0.w - src1.w
-
-
.. opcode:: LRP - Linear Interpolate
.. math::
dst.w = src.w - \lfloor src.w\rfloor
-.. opcode:: CLAMP - Clamp
-
-.. math::
-
- dst.x = clamp(src0.x, src1.x, src2.x)
-
- dst.y = clamp(src0.y, src1.y, src2.y)
-
- dst.z = clamp(src0.z, src1.z, src2.z)
-
- dst.w = clamp(src0.w, src1.w, src2.w)
-
-
.. opcode:: FLR - Floor
.. math::
dst.w = 1
-.. opcode:: ABS - Absolute
-
-.. math::
-
- dst.x = |src.x|
-
- dst.y = |src.y|
-
- dst.z = |src.z|
-
- dst.w = |src.w|
-
-
.. opcode:: DPH - Homogeneous Dot Product
This instruction replicates its result.
.. opcode:: IBFE - Signed Bitfield Extract
- See SM5 instruction of the same name. Extracts a set of bits from the input,
- and sign-extends them if the high bit of the extracted window is set.
+ Like GLSL bitfieldExtract. Extracts a set of bits from the input, and
+ sign-extends them if the high bit of the extracted window is set.
Pseudocode::
def ibfe(value, offset, bits):
- offset = offset & 0x1f
- bits = bits & 0x1f
+ if offset < 0 or bits < 0 or offset + bits > 32:
+ return undefined
if bits == 0: return 0
# Note: >> sign-extends
- if width + offset < 32:
- return (value << (32 - offset - bits)) >> (32 - bits)
- else:
- return value >> offset
+ return (value << (32 - offset - bits)) >> (32 - bits)
.. opcode:: UBFE - Unsigned Bitfield Extract
- See SM5 instruction of the same name. Extracts a set of bits from the input,
- without any sign-extension.
+ Like GLSL bitfieldExtract. Extracts a set of bits from the input, without
+ any sign-extension.
Pseudocode::
def ubfe(value, offset, bits):
- offset = offset & 0x1f
- bits = bits & 0x1f
+ if offset < 0 or bits < 0 or offset + bits > 32:
+ return undefined
if bits == 0: return 0
# Note: >> does not sign-extend
- if width + offset < 32:
- return (value << (32 - offset - bits)) >> (32 - bits)
- else:
- return value >> offset
+ return (value << (32 - offset - bits)) >> (32 - bits)
.. opcode:: BFI - Bitfield Insert
- See SM5 instruction of the same name. Replaces a bit region of 'base' with
- the low bits of 'insert'.
+ Like GLSL bitfieldInsert. Replaces a bit region of 'base' with the low bits
+ of 'insert'.
Pseudocode::
def bfi(base, insert, offset, bits):
- offset = offset & 0x1f
- bits = bits & 0x1f
+ if offset < 0 or bits < 0 or offset + bits > 32:
+ return undefined
+ # << defined such that mask == ~0 when bits == 32, offset == 0
mask = ((1 << bits) - 1) << offset
return ((insert << offset) & mask) | (base & ~mask)
.. opcode:: DABS - Absolute
+.. math::
+
dst.xy = |src0.xy|
dst.zw = |src0.zw|
dst.zw = src0.zw \times src1.zw + src2.zw
+.. opcode:: DDIV - Divide
+
+.. math::
+
+ dst.xy = \frac{src0.xy}{src1.xy}
+
+ dst.zw = \frac{src0.zw}{src1.zw}
+
+
.. opcode:: DRCP - Reciprocal
.. math::
dst.y = unsigned(src0.zw)
+64-bit Integer ISA
+^^^^^^^^^^^^^^^^^^
+
+The 64-bit integer opcodes reinterpret four-component vectors into
+two-component vectors with 64-bits in each component.
+
+.. opcode:: I64ABS - 64-bit Integer Absolute Value
+
+.. math::
+
+ dst.xy = |src0.xy|
+ dst.zw = |src0.zw|
+
+.. opcode:: I64NEG - 64-bit Integer Negate
+
+ Two's complement.
+
+.. math::
+
+ dst.xy = -src.xy
+ dst.zw = -src.zw
+
+.. opcode:: I64SSG - 64-bit Integer Set Sign
+
+.. math::
+
+ dst.xy = (src0.xy < 0) ? -1 : (src0.xy > 0) ? 1 : 0
+ dst.zw = (src0.zw < 0) ? -1 : (src0.zw > 0) ? 1 : 0
+
+.. opcode:: U64ADD - 64-bit Integer Add
+
+.. math::
+
+ dst.xy = src0.xy + src1.xy
+ dst.zw = src0.zw + src1.zw
+
+.. opcode:: U64MUL - 64-bit Integer Multiply
+
+.. math::
+
+ dst.xy = src0.xy * src1.xy
+ dst.zw = src0.zw * src1.zw
+
+.. opcode:: U64SEQ - 64-bit Integer Set on Equal
+
+.. math::
+
+ dst.x = src0.xy == src1.xy ? \sim 0 : 0
+ dst.z = src0.zw == src1.zw ? \sim 0 : 0
+
+.. opcode:: U64SNE - 64-bit Integer Set on Not Equal
+
+.. math::
+
+ dst.x = src0.xy != src1.xy ? \sim 0 : 0
+ dst.z = src0.zw != src1.zw ? \sim 0 : 0
+
+.. opcode:: U64SLT - 64-bit Unsigned Integer Set on Less Than
+
+.. math::
+
+ dst.x = src0.xy < src1.xy ? \sim 0 : 0
+ dst.z = src0.zw < src1.zw ? \sim 0 : 0
+
+.. opcode:: U64SGE - 64-bit Unsigned Integer Set on Greater Equal
+
+.. math::
+
+ dst.x = src0.xy >= src1.xy ? \sim 0 : 0
+ dst.z = src0.zw >= src1.zw ? \sim 0 : 0
+
+.. opcode:: I64SLT - 64-bit Signed Integer Set on Less Than
+
+.. math::
+
+ dst.x = src0.xy < src1.xy ? \sim 0 : 0
+ dst.z = src0.zw < src1.zw ? \sim 0 : 0
+
+.. opcode:: I64SGE - 64-bit Signed Integer Set on Greater Equal
+
+.. math::
+
+ dst.x = src0.xy >= src1.xy ? \sim 0 : 0
+ dst.z = src0.zw >= src1.zw ? \sim 0 : 0
+
+.. opcode:: I64MIN - Minimum of 64-bit Signed Integers
+
+.. math::
+
+ dst.xy = min(src0.xy, src1.xy)
+ dst.zw = min(src0.zw, src1.zw)
+
+.. opcode:: U64MIN - Minimum of 64-bit Unsigned Integers
+
+.. math::
+
+ dst.xy = min(src0.xy, src1.xy)
+ dst.zw = min(src0.zw, src1.zw)
+
+.. opcode:: I64MAX - Maximum of 64-bit Signed Integers
+
+.. math::
+
+ dst.xy = max(src0.xy, src1.xy)
+ dst.zw = max(src0.zw, src1.zw)
+
+.. opcode:: U64MAX - Maximum of 64-bit Unsigned Integers
+
+.. math::
+
+ dst.xy = max(src0.xy, src1.xy)
+ dst.zw = max(src0.zw, src1.zw)
+
+.. opcode:: U64SHL - Shift Left 64-bit Unsigned Integer
+
+ The shift count is masked with 0x3f before the shift is applied.
+
+.. math::
+
+ dst.xy = src0.xy << (0x3f \& src1.x)
+ dst.zw = src0.zw << (0x3f \& src1.y)
+
+.. opcode:: I64SHR - Arithmetic Shift Right (of 64-bit Signed Integer)
+
+ The shift count is masked with 0x3f before the shift is applied.
+
+.. math::
+
+ dst.xy = src0.xy >> (0x3f \& src1.x)
+ dst.zw = src0.zw >> (0x3f \& src1.y)
+
+.. opcode:: U64SHR - Logical Shift Right (of 64-bit Unsigned Integer)
+
+ The shift count is masked with 0x3f before the shift is applied.
+
+.. math::
+
+ dst.xy = src0.xy >> (unsigned) (0x3f \& src1.x)
+ dst.zw = src0.zw >> (unsigned) (0x3f \& src1.y)
+
+.. opcode:: I64DIV - 64-bit Signed Integer Division
+
+.. math::
+
+ dst.xy = src0.xy \ src1.xy
+ dst.zw = src0.zw \ src1.zw
+
+.. opcode:: U64DIV - 64-bit Unsigned Integer Division
+
+.. math::
+
+ dst.xy = src0.xy \ src1.xy
+ dst.zw = src0.zw \ src1.zw
+
+.. opcode:: U64MOD - 64-bit Unsigned Integer Remainder
+
+.. math::
+
+ dst.xy = src0.xy \bmod src1.xy
+ dst.zw = src0.zw \bmod src1.zw
+
+.. opcode:: I64MOD - 64-bit Signed Integer Remainder
+
+.. math::
+
+ dst.xy = src0.xy \bmod src1.xy
+ dst.zw = src0.zw \bmod src1.zw
+
+.. opcode:: F2U64 - Float to 64-bit Unsigned Int
+
+.. math::
+
+ dst.xy = (uint64_t) src0.x
+ dst.zw = (uint64_t) src0.y
+
+.. opcode:: F2I64 - Float to 64-bit Int
+
+.. math::
+
+ dst.xy = (int64_t) src0.x
+ dst.zw = (int64_t) src0.y
+
+.. opcode:: U2I64 - Unsigned Integer to 64-bit Integer
+
+ This is a zero extension.
+
+.. math::
+
+ dst.xy = (uint64_t) src0.x
+ dst.zw = (uint64_t) src0.y
+
+.. opcode:: I2I64 - Signed Integer to 64-bit Integer
+
+ This is a sign extension.
+
+.. math::
+
+ dst.xy = (int64_t) src0.x
+ dst.zw = (int64_t) src0.y
+
+.. opcode:: D2U64 - Double to 64-bit Unsigned Int
+
+.. math::
+
+ dst.xy = (uint64_t) src0.xy
+ dst.zw = (uint64_t) src0.zw
+
+.. opcode:: D2I64 - Double to 64-bit Int
+
+.. math::
+
+ dst.xy = (int64_t) src0.xy
+ dst.zw = (int64_t) src0.zw
+
+.. opcode:: U642F - 64-bit unsigned integer to float
+
+.. math::
+
+ dst.x = (float) src0.xy
+ dst.y = (float) src0.zw
+
+.. opcode:: I642F - 64-bit Int to Float
+
+.. math::
+
+ dst.x = (float) src0.xy
+ dst.y = (float) src0.zw
+
+.. opcode:: U642D - 64-bit unsigned integer to double
+
+.. math::
+
+ dst.xy = (double) src0.xy
+ dst.zw = (double) src0.zw
+
+.. opcode:: I642D - 64-bit Int to double
+
+.. math::
+
+ dst.xy = (double) src0.xy
+ dst.zw = (double) src0.zw
+
.. _samplingopcodes:
Resource Sampling Opcodes
.. opcode:: SAMPLE
Using provided address, sample data from the specified texture using the
- filtering mode identified by the gven sampler. The source data may come from
+ filtering mode identified by the given sampler. The source data may come from
any resource type other than buffers.
Syntax: ``SAMPLE dst, address, sampler_view, sampler``
texture arrays and 2D textures. address.w is always
ignored.
+ A swizzle suffix may be added to the resource argument
+ this will cause the resource data to be swizzled accordingly.
+
.. opcode:: STORE - Write data to a shader resource
Syntax: ``STORE resource, address, src``
image, while .w will contain the number of samples for multi-sampled
images.
+.. opcode:: FBFETCH - Load data from framebuffer
+
+ Syntax: ``FBFETCH dst, output``
+
+ Example: ``FBFETCH TEMP[0], OUT[0]``
+
+ This is only valid on ``COLOR`` semantic outputs. Returns the color
+ of the current position in the framebuffer from before this fragment
+ shader invocation. May return the same value from multiple calls for
+ a particular output within a single invocation. Note that result may
+ be undefined if a fragment is drawn multiple times without a blend
+ barrier in between.
+
.. _threadsyncopcodes:
resource[offset] = (dst_x > src_x ? dst_x : src_x)
+.. _voteopcodes:
+
+Vote opcodes
+^^^^^^^^^^^^
+
+These opcodes compare the given value across the shader invocations
+running in the current SIMD group. The details of exactly which
+invocations get compared are implementation-defined, and it would be a
+correct implementation to only ever consider the current thread's
+value. (i.e. SIMD group of 1). The argument is treated as a boolean.
+
+.. opcode:: VOTE_ANY - Value is set in any of the current invocations
+
+.. opcode:: VOTE_ALL - Value is set in all of the current invocations
+
+.. opcode:: VOTE_EQ - Value is the same in all of the current invocations
+
Explanation of symbols used
------------------------------
"""""""""""""""""""
For vertex shader outputs or fragment shader inputs/outputs, this
-label indicates that the resister contains an R,G,B,A color.
+label indicates that the register contains an R,G,B,A color.
Several shader inputs/outputs may contain colors so the semantic index
is used to distinguish them. For example, color[0] may be the diffuse
TGSI_SEMANTIC_CLIPDIST
""""""""""""""""""""""
+Note this covers clipping and culling distances.
+
When components of vertex elements are identified this way, these
values are each assumed to be a float32 signed distance to a plane.
+
+For clip distances:
Primitive setup only invokes rasterization on pixels for which
-the interpolated plane distances are >= 0. Multiple clip planes
-can be implemented simultaneously, by annotating multiple
-components of one or more vertex elements with the above specified
-semantic. The limits on both clip and cull distances are bound
+the interpolated plane distances are >= 0.
+
+For cull distances:
+Primitives will be completely discarded if the plane distance
+for all of the vertices in the primitive are < 0.
+If a vertex has a cull distance of NaN, that vertex counts as "out"
+(as if its < 0);
+
+Multiple clip/cull planes can be implemented simultaneously, by
+annotating multiple components of one or more vertex elements with
+the above specified semantic.
+The limits on both clip and cull distances are bound
by the PIPE_MAX_CLIP_OR_CULL_DISTANCE_COUNT define which defines
the maximum number of components that can be used to hold the
distances and by the PIPE_MAX_CLIP_OR_CULL_DISTANCE_ELEMENT_COUNT
which specifies the maximum number of registers which can be
annotated with those semantics.
+The properties NUM_CLIPDIST_ENABLED and NUM_CULLDIST_ENABLED
+are used to divide up the 2 x vec4 space between clipping and culling.
TGSI_SEMANTIC_SAMPLEID
""""""""""""""""""""""
component is used.
+TGSI_SEMANTIC_WORK_DIM
+""""""""""""""""""""""
+
+For compute shaders started via opencl this retrieves the work_dim
+parameter to the clEnqueueNDRangeKernel call with which the shader
+was started.
+
+
+TGSI_SEMANTIC_GRID_SIZE
+"""""""""""""""""""""""
+
+For compute shaders, this semantic indicates the maximum (x, y, z) dimensions
+of a grid of thread blocks.
+
+
+TGSI_SEMANTIC_BLOCK_ID
+""""""""""""""""""""""
+
+For compute shaders, this semantic indicates the (x, y, z) coordinates of the
+current block inside of the grid.
+
+
+TGSI_SEMANTIC_BLOCK_SIZE
+""""""""""""""""""""""""
+
+For compute shaders, this semantic indicates the maximum (x, y, z) dimensions
+of a block in threads.
+
+
+TGSI_SEMANTIC_THREAD_ID
+"""""""""""""""""""""""
+
+For compute shaders, this semantic indicates the (x, y, z) coordinates of the
+current thread inside of the block.
+
+
Declaration Interpolate
^^^^^^^^^^^^^^^^^^^^^^^
the fragment shader (regardless of fragment shader side effects). Corresponds
to GLSL early_fragment_tests.
+NEXT_SHADER
+"""""""""""
+
+Which shader stage will MOST LIKELY follow after this shader when the shader
+is bound. This is only a hint to the driver and doesn't have to be precise.
+Only set for VS and TES.
+
+CS_FIXED_BLOCK_WIDTH / HEIGHT / DEPTH
+"""""""""""""""""""""""""""""""""""""
+
+Threads per block in each dimension, if known at compile time. If the block size
+is known all three should be at least 1. If it is unknown they should all be set
+to 0 or not set.
+
+MUL_ZERO_WINS
+"""""""""""""
+
+The MUL TGSI operation (FP32 multiplication) will return 0 if either
+of the operands are equal to 0. That means that 0 * Inf = 0. This
+should be set the same way for an entire pipeline. Note that this
+applies not only to the literal MUL TGSI opcode, but all FP32
+multiplications implied by other operations, such as MAD, FMA, DP2,
+DP3, DP4, DPH, DST, LOG, LRP, XPD, and possibly others. If there is a
+mismatch between shaders, then it is unspecified whether this behavior
+will be enabled.
+
+
Texture Sampling and Texture Formats
------------------------------------
projects / mesa.git / blobdiff