Снижение CUDA минимальное значение и индекс

Question

я реализовал минимум уменьшают с помощью CUDA 8 следуя this большое объяснение и изменение его

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

и он прекрасно работает, и я получаю минимальное значение. Тем не менее, я не забочусь о минимальном значении, только о его индексе в исходном массиве ввода.

Я попытался модифицировать мой код чуток

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

Но это не работает. Я чувствую, что мне не хватает чего-то важного и что это не тот способ, но мой поиск не принес результатов.

Answer 1

Здесь есть несколько проблем. Вам необходимо изменить как функцию минимального значения, так и минимальное значение для распространения как минимального значения, так и его индекса каждый раз, когда будет найден новый локальный минимум. Возможно, что-то вроде этого:

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

[Примечание: записано в браузере, не компиляции или тестирования, использовать на свой страх и риск]

Это должно оставить каждый блок держит это правильный локальный минимум и индекс. Тогда у вас есть вторая проблема. Это:

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

не работает. Нет гарантии, что минимальное значение и его индекс будут правильно установлены в этом коде. Это связано с тем, что нет никакой гарантии, что обе атомные операции имеют какую-либо синхронизацию, и существует потенциальная гонка, когда один блок может правильно перезаписать минимальное значение другого блока, но затем его индекс будет перезаписан замененным блоком. Единственное решение здесь было бы своего рода мьютексом или запустить второе сокращение ядра по результатам каждого блока.