CUDA Reduzierung Minimalwert und Index

Question

I implementiert ein Minimum CUDA 8 durch folgende this große Erklärung reduce verwenden und es

__inline__ __device__ int warpReduceMin(int val) 
{ 
    for (int offset = warpSize/2; offset > 0; offset /= 2) 
    { 
     int tmpVal = __shfl_down(val, offset); 
     if (tmpVal < val) 
     { 
      val = tmpVal; 
     } 
    } 
    return val; 
} 

__inline__ __device__ int blockReduceMin(int val) 
{ 

    static __shared__ int shared[32]; // Shared mem for 32 partial mins 
    int lane = threadIdx.x % warpSize; 
    int wid = threadIdx.x/warpSize; 

    val = warpReduceMin(val);  // Each warp performs partial reduction 

    if (lane == 0) 
    { 
     shared[wid] = val; // Write reduced value to shared memory 
    } 

    __syncthreads();    // Wait for all partial reductions 

    //read from shared memory only if that warp existed 
    val = (threadIdx.x < blockDim.x/warpSize) ? shared[lane] : INT_MAX; 

    if (wid == 0) 
    { 
     val = warpReduceMin(val); //Final reduce within first warp 
    } 

    return val; 
} 

__global__ void deviceReduceBlockAtomicKernel(int *in, int* out, int N) { 
    int minVal = INT_MAX; 
    for (int i = blockIdx.x * blockDim.x + threadIdx.x; 
     i < N; 
     i += blockDim.x * gridDim.x) 
    { 
     minVal = min(minVal, in[i]); 
    } 
    minVal = blockReduceMin(minVal); 
    if (threadIdx.x == 0) 
    { 
     atomicMin(out, minVal); 
    } 
} 

modifizieren und es funktioniert super und ich bin den Minimalwert zu bekommen. Der minimale Wert ist mir jedoch egal, nur der Index im ursprünglichen Eingabe-Array.

Ich habe versucht, meinen Code ein bisschen

__inline__ __device__ int warpReduceMin(int val, int* idx) // Adding output idx 
{ 
    for (int offset = warpSize/2; offset > 0; offset /= 2) 
    { 
     int tmpVal = __shfl_down(val, offset); 
     if (tmpVal < val) 
     { 
      *idx = blockIdx.x * blockDim.x + threadIdx.x + offset; // I guess I'm missing something here 
      val = tmpVal; 
     } 
    } 
    return val; 
} 

... 
blockReduceMin stayed the same only adding idx to function calls 
... 

__global__ void deviceReduceBlockAtomicKernel(int *in, int* out, int N) { 
    int minVal = INT_MAX; 
    int minIdx = 0; // Added this 
    for (int i = blockIdx.x * blockDim.x + threadIdx.x; 
     i < N; 
     i += blockDim.x * gridDim.x) 
    { 
     if (in[i] < minVal) 
     { 
      minVal = in[i]; 
      minIdx = i; // Added this 
     } 
    } 
    minVal = blockReduceMin(minVal, &minIdx); 
    if (threadIdx.x == 0) 
    { 
     int old = atomicMin(out, minVal); 
     if (old != minVal) // value was updated 
     { 
      atomicExch(out + 1, minIdx); 
     } 
    } 
} 

Aber es funktioniert nicht ändern. Ich habe das Gefühl, dass mir etwas Wichtiges fehlt und dass dies nicht der richtige Weg ist, aber meine Suche ergab keine Ergebnisse.

Answer 1

Hier gibt es mehrere Probleme. Sie müssen sowohl die Warp- als auch die Blockminimumfunktionen ändern, um sowohl den Minimalwert als auch den Index jedes Mal zu verbreiten, wenn ein neues lokales Minimum gefunden wird. so etwas wie dies vielleicht:

__inline__ __device__ void warpReduceMin(int& val, int& idx) 
{ 
    for (int offset = warpSize/2; offset > 0; offset /= 2) { 
     int tmpVal = __shfl_down(val, offset); 
     int tmpIdx = __shfl_down(idx, offset); 
     if (tmpVal < val) { 
      val = tmpVal; 
      idx = tmpIdx; 
     } 
    } 
} 

__inline__ __device__ void blockReduceMin(int& val, int& idx) 
{ 

    static __shared__ int values[32], indices[32]; // Shared mem for 32 partial mins 
    int lane = threadIdx.x % warpSize; 
    int wid = threadIdx.x/warpSize; 

    warpReduceMin(val, idx);  // Each warp performs partial reduction 

    if (lane == 0) { 
     values[wid] = val; // Write reduced value to shared memory 
     indices[wid] = idx; // Write reduced value to shared memory 
    } 

    __syncthreads();    // Wait for all partial reductions 

    //read from shared memory only if that warp existed 
    if (threadIdx.x < blockDim.x/warpSize) { 
     val = values[lane]; 
     idx = indices[lane]; 
    } else { 
     val = INT_MAX; 
     idx = 0; 
    } 

    if (wid == 0) { 
     warpReduceMin(val, idx); //Final reduce within first warp 
    } 
} 

[Anmerkung: geschrieben in Browser, kompiliert oder nie getestet, verwenden Sie auf eigene Gefahr]

, dass jeder Block verlassen sollte es korrekt lokalen Minimum und Index zu halten. Dann hast du ein zweites Problem. Dies ist:

int old = atomicMin(out, minVal); 
if (old != minVal) // value was updated 
{ 
    atomicExch(out + 1, minIdx); 
} 

ist kaputt. Es gibt keine Garantie, dass der Mindestwert und sein Index in diesem Code korrekt festgelegt werden. Dies liegt daran, dass es keine Garantie dafür gibt, dass beide atomaren Operationen eine Synchronisation haben, und es gibt eine potentielle Wettlaufbahn, bei der ein Block den Minimalwert eines anderen Blocks korrekt überschreiben kann, dessen Index jedoch durch den ersetzten Block überschrieben wird. Die einzige Lösung wäre hier irgendeine Art von Mutex, oder man führt einen zweiten Reduzierungskern auf den Ergebnissen jedes Blocks aus.