reduction in registers

Accelerated Computing CUDA CUDA Programming and Performance
1

Hello, I have three arrays of integers, one with all values, another where I keep the values resulting from the reduction of array0, and another that keeps the positions of the maximum values of the reduction of each warp but with reference to its position in array0

array0 → all values
array1 → maximum value for each warp
array2 → the position of each maximum value of each warp

example

array1[0] = 2234 → maximum value of a warp
array2[0] = 4 → the position of 2234 in array0

how save the maximum value and position at the same time?

this is my code:

for (int i=1; i<=width; i*=2) {
        int n = __shfl_xor(value, i, width);
        if (lane_id >= i && value<n) { value = n; }
    } 

    if (threadIdx.x % warpSize == warpSize-1) {
        maximums[warp_id] = value;
	d_pos_out[warp_id] = positions[id];
    }

values are saved correctly, but the positions do not match with the values

any idea?

many thanks

2

I don’t really understand your code, but how about

int position = threadIdx.x;
    for (int i=warpSize/2; i>0; i/=2) {
        int val_n = __shfl_up(value, i, width);
        int pos_n = __shfl_up(position, i, width);
        if (val_n > value) {
            value = val_n;
            position = pos_n;
        }
    }

(disclaimer: written entirely in browser, untested!)

3

In the way you suggest not match indexes, and my code always gets the last thread of each warp with
int id = ((blockIdx.x * blockDim.x) + threadIdx.x);

How i can get the specific thread i want? that matches with the variable ‘value’

thanks.