How to compare value within a warp is positive or negative using warp-level primitive?

thread 0, 1, 2, 3 has 8 values, I want to check each value whether positive or negative. Like, 1. tell me in each line, how many positive number or negative number? 2. The value will be from small (negative) to large(positive), how to find the sign change position(last one is negative and next one is positive)?

Thank you!!!

According to the diagram, each of the 4 threads have 4 values (2 for row 0, 2 for row 8). Probably you were referring only to the row 0 values. For each thread those are stored in the variables c0 and c1. Please ask, if you have difficulty to access those result variables of the previous matrix operation. It depends on how you call the matrix operation.

For counting the number of negative values in each line=row, you get two results (1 result for row 0-7, 1 result for row 8-15) for each 4 threads.

First each thread adds up its two variables: int c01 = (c0 < 0) + (c1 < 0); int c23 = (c2 < 0) + (c3 < 0);

You can combine data within a warp with warp reduce and warp shuffle functions.

You could use __reduce_add_sync, but it can only do one reduction at a time. So only 4 threads could participate.

So better shuffle the data for each row into threads 0, 4, 8, 12, 16, 20, 24, 28:

int shuffled1 = c01 + __shfl_down_sync(0xFFFFFFFF, c01, 1, 32);
int shuffled2 = shuffled1 + __shfl_down_sync(0xFFFFFFFF, shuffled1, 1, 32);

Now the mentioned threads have the row results for rows 0…7.

You can do the same with rows 8…15. Or, more efficiently: Put c01 in the lower half and c23 in the upper half of a 32 bit value. And shuffle and add this 32 bit value containing both halfs at the same time. This is possible, because c01 and c23 contain counts, which are positive and small.

1 Like