- 8.1. Memory
- 8.2. Integer Support
- 8.3. Floating-Point Support
- 8.4. Conditional Code
- 8.5. Textures and Surfaces
- 8.6. Miscellaneous Instructions
- 8.7. Instruction Sets
This chapter is from the book
This chapter is from the book
This chapter is from the book
8.6. Miscellaneous Instructions
8.6.1. Warp-Level Primitives
It did not take long for the importance of warps as a primitive unit of execution (naturally residing between threads and blocks) to become evident to CUDA programmers. Starting with SM 1.x, NVIDIA began adding instructions that specifically operate on warps.
Vote
That CUDA architectures are 32-bit and that warps are comprised of 32 threads made an irresistible match to instructions that can evaluate a condition and broadcast a 1-bit result to every thread in the warp. The VOTE instruction (first available in SM 1.2) evaluates a condition and broadcasts the result to all threads in the warp. The __any() intrinsic returns 1 if the predicate is true for any of the 32 threads in the warp. The __all() intrinsic returns 1 if the predicate is true for all of the 32 threads in the warp.
The Fermi architecture added a new variant of VOTE that passes back the predicate result for every thread in the warp. The __ballot() intrinsic evaluates a condition for all threads in the warp and returns a 32-bit value where each bit gives the condition for the corresponding thread in the warp.
Shuffle
Kepler added shuffle instructions that enable data interchange between threads within a warp without staging the data through shared memory. Although these instructions execute with the same latency as shared memory, they have the benefit of doing the exchange without performing both a read and a write, and they can reduce shared memory usage.
The following instruction is wrapped in a number of device functions that use inline PTX assembly defined in sm_30_intrinsics.h.
int __shfl(int var, int srcLane, int width=32); int __shfl_up(int var, unsigned int delta, int width=32); int __shfl_down(int var, unsigned int delta, int width=32); int __shfl_xor(int var, int laneMask, int width=32);
The width parameter, which defaults to the warp width of 32, must be a power of 2 in the range 2..32. It enables subdivision of the warp into segments; if width<32, each subsection of the warp behaves as a separate entity with a starting logical lane ID of 0. A thread may only exchange data with other threads in its subsection.
__shfl() returns the value of var held by the thread whose ID is given by srcLane. If srcLane is outside the range 0..width-1, the thread’s own value of var is returned. This variant of the instruction can be used to broadcast values within a warp. __shfl_up() calculates a source lane ID by subtracting delta from the caller’s lane ID and clamping to the range 0..width-1. __shfl_down() calculates a source lane ID by adding delta to the caller’s lane ID.
__shfl_up() and __shfl_down() enable warp-level scan and reverse scan operations, respectively. __shfl_xor() calculates a source lane ID by performing a bitwise XOR of the caller’s lane ID with laneMask; the value of var held by the resulting lane ID is returned. This variant can be used to do a reduction across the warps (or subwarps); each thread computes the reduction using a differently ordered series of the associative operator.
8.6.2. Block-Level Primitives
The __syncthreads() intrinsic serves as a barrier. It causes all threads to wait until every thread in the threadblock has arrived at the __syncthreads(). The Fermi instruction set (SM 2.x) added several new block-level barriers that aggregate information about the threads in the threadblock.
- __syncthreads_count(): evaluates a predicate and returns the sum of threads for which the predicate was true
- __syncthreads_or(): returns the OR of all the inputs across the threadblock
- __syncthreads_and(): returns the AND of all the inputs across the threadblock
8.6.3. Performance Counter
Developers can define their own set of performance counters and increment them in live code with the __prof_trigger() intrinsic.
void __prof_trigger(int counter);
Calling this function increments the corresponding counter by 1 per warp. counter must be in the range 0..7; counters 8..15 are reserved. The value of the counters may be obtained by listing prof_trigger_00..prof_trigger_07 in the profiler configuration file.
8.6.4. Video Instructions
The video instructions described in this section are accessible only via the inline PTX assembler. Their basic functionality is described here to help developers to decide whether they might be beneficial for their application. Anyone intending to use these instructions, however, should consult the PTX ISA specification.
Scalar Video Instructions
The scalar video instructions, added with SM 2.0 hardware, enable efficient operations on the short (8- and 16-bit) integer types needed for video processing. As described in the PTX 3.1 ISA Specification, the format of these instructions is as follows.
vop.dtype.atype.btype{.sat} d, a{.asel}, b{.bsel};
vop.dtype.atype.btype{.sat}.secop d, a{.asel}, b{.bsel}, c;
The source and destination operands are all 32-bit registers. dtype, atype, and btype may be .u32 or .s32 for unsigned and signed 32-bit integers, respectively. The asel/bsel specifiers select which 8- or 16-bit value to extract from the source operands: b0, b1, b2, and b3 select bytes (numbering from the least significant), and h0/h1 select the least significant and most significant 16 bits, respectively.
Once the input values are extracted, they are sign- or zero-extended internally to signed 33-bit integers, and the primary operation is performed, producing a 34-bit intermediate result whose sign depends on dtype. Finally, the result is clamped to the output range, and one of the following operations is performed.
- Apply a second operation (add, min or max) to the intermediate result and a third operand.
- Truncate the intermediate result to an 8- or 16-bit value and merge into a specified position in the third operand to produce the final result.
The lower 32 bits are then written to the destination operand.
The vset instruction performs a comparison between the 8-, 16-, or 32-bit input operands and generates the corresponding predicate (1 or 0) as output. The PTX scalar video instructions and the corresponding operations are given in Table 8.14.
Table 8.14 Scalar Video Instructions.
|
MNEMONIC |
OPERATION |
|
vabsdiff |
abs(a-b) |
|
vadd |
a+b |
|
vavrg |
(a+b)/2 |
|
vmad |
a*b+c |
|
vmax |
max(a,b) |
|
vmin |
min(a,b) |
|
vset |
Compare a and b |
|
vshl |
a< |
|
vshr |
a>>b |
|
vsub |
a-b |
Vector Video Instructions (SM 3.0 only)
These instructions, added with SM 3.0, are similar to the scalar video instructions in that they promote the inputs to a canonical integer format, perform the core operation, and then clamp and optionally merge the output. But they deliver higher performance by operating on pairs of 16-bit values or quads of 8-bit values.
Table 8.15 summarizes the PTX instructions and corresponding operations implemented by these instructions. They are most useful for video processing and certain image processing operations (such as the median filter).
Table 8.15 Vector Video Instructions
|
MNEMONIC |
OPERATION |
|
vabsdiff[2|4] |
abs(a-b) |
|
vadd[2|4] |
a+b |
|
vavrg[2|4] |
(a+b)/2 |
|
vmax[2|4] |
max(a,b) |
|
vmin[2|4] |
min(a,b) |
|
vset[2|4] |
Compare a and b |
|
vsub[2|4] |
a-b |
8.6.5. Special Registers
Many special registers are accessed by referencing the built-in variables threadIdx, blockIdx, blockDim, and gridDim. These pseudo-variables, described in detail in Section 7.3, are 3-dimensional structures that specify the thread ID, block ID, thread count, and block count, respectively.
Besides those, another special register is the SM’s clock register, which increments with each clock cycle. This counter can be read with the __clock() or __clock64() intrinsic. The counters are separately tracked for each SM and, like the time stamp counters on CPUs, are most useful for measuring relative performance of different code sequences and best avoided when trying to calculate wall clock times.