Réduction CUDA valeur minimale et l'indice

Question

I mis en œuvre un minimum à l'aide à réduire CUDA 8 en suivant this grande explication et de la modifier

__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); 
    } 
} 

et il fonctionne très bien et je suis obtenir la valeur minimale. Cependant, je ne me soucie pas de la valeur minimale, seulement de son index dans le tableau d'entrée d'origine.

J'ai essayé de modifier mon code un peu

__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); 
     } 
    } 
} 

Mais il ne fonctionne pas. J'ai l'impression qu'il me manque quelque chose d'important et que ce n'est pas la bonne façon de procéder, mais ma recherche n'a donné aucun résultat.

Answer 1

Il y a plusieurs problèmes ici. Vous devez modifier les fonctions minimales de distorsion et de blocage pour propager à la fois la valeur minimale et son index chaque fois qu'un nouveau minimum local est trouvé. Peut-être quelque chose comme ceci:

__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 
    } 
} 

[Note: écrit dans le navigateur, jamais compilé ou testé, utilisez à vos propres risques]

Cela devrait laisser chaque bloc maintenant il est correct minimum local et de l'indice. Ensuite, vous avez un deuxième problème. Ceci:

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

est cassé. Il n'y a aucune garantie que la valeur minimale et son index seront correctement définis dans ce code. En effet, il n'y a aucune garantie que les deux opérations atomiques ont une synchronisation et il y a une course potentielle où un bloc peut remplacer correctement la valeur minimale d'un autre bloc, mais alors son index écrasé par le bloc qu'il a remplacé. La seule solution ici serait une sorte de mutex, ou exécuter un second noyau de réduction sur les résultats de chaque bloc.