Functions for computing lane masks.
More...
|
| __device__ lanemask_t | sp::lanemaskLt (int laneID) |
| | Get a bitmask with a 1 in every position lower than this thread's lane ID. More...
|
| |
| __device__ lanemask_t | sp::lanemaskLe (int laneID) |
| | Get a bitmask with a 1 in every position lower than or equal to this thread's lane ID. More...
|
| |
| __device__ lanemask_t | sp::lanemaskEq (int laneID) |
| | Get a bitmask with a 1 only in the position equal to this thread's lane ID. More...
|
| |
| __device__ lanemask_t | sp::lanemaskGe (int laneID) |
| | Get a bitmask with a 1 in every position greater than this thread's lane ID. More...
|
| |
| __device__ lanemask_t | sp::lanemaskGt (int laneID) |
| | Get a bitmask with a 1 in every position greater than or equal to this thread's lane ID. More...
|
| |
Functions for computing lane masks.
These are preferable to using inline PTX to read the corresponding special registers because:
- The optimiser can use them in value-tracking optimisations. If it knows things about your lane ID, then it also can deduce things about the values of these, and hence potentially optimise other calculations. The optimiser also has the freedom to do things like choose between calculating lanemask from laneID or using the special instruction depending on context. If some of the intermediate values have already been calculated for another reason, it may be cheaper not to use the special instruction.
- These functions will be emulated using arithemtic instructions in cases where no direct instruction exists, or when the available instruction is slower than the math.
◆ lanemaskEq()
| __device__ lanemask_t sp::lanemaskEq |
( |
int |
laneID | ) |
|
Get a bitmask with a 1 only in the position equal to this thread's lane ID.
- Parameters
-
| laneID | The laneID for this thread. This may not be used (and hence optimised out) if the target has an instruction for directly getting the desired mask. Otherwise this will be used to calculate the result. |
◆ lanemaskGe()
| __device__ lanemask_t sp::lanemaskGe |
( |
int |
laneID | ) |
|
Get a bitmask with a 1 in every position greater than this thread's lane ID.
- Parameters
-
| laneID | The laneID for this thread. This may not be used (and hence optimised out) if the target has an instruction for directly getting the desired mask. Otherwise this will be used to calculate the result. |
◆ lanemaskGt()
| __device__ lanemask_t sp::lanemaskGt |
( |
int |
laneID | ) |
|
Get a bitmask with a 1 in every position greater than or equal to this thread's lane ID.
- Parameters
-
| laneID | The laneID for this thread. This may not be used (and hence optimised out) if the target has an instruction for directly getting the desired mask. Otherwise this will be used to calculate the result. |
◆ lanemaskLe()
| __device__ lanemask_t sp::lanemaskLe |
( |
int |
laneID | ) |
|
Get a bitmask with a 1 in every position lower than or equal to this thread's lane ID.
- Parameters
-
| laneID | The laneID for this thread. This may not be used (and hence optimised out) if the target has an instruction for directly getting the desired mask. Otherwise this will be used to calculate the result. |
◆ lanemaskLt()
| __device__ lanemask_t sp::lanemaskLt |
( |
int |
laneID | ) |
|
Get a bitmask with a 1 in every position lower than this thread's lane ID.
- Parameters
-
| laneID | The laneID for this thread. This may not be used (and hence optimised out) if the target has an instruction for directly getting the desired mask. Otherwise this will be used to calculate the result. |
