CUDA matriz de memoria compartida - comportamiento extraño

Question

En un kernel CUDA, tengo un código similar a lo siguiente. Estoy tratando de calcular un numerador por hilo, y se acumulan los numeradores sobre el bloque para calcular un denominador, y luego regresar la relación. Sin embargo, CUDA está fijando el valor de denom a lo que se calcula el valor de numer por el hilo en el bloque con el mayor threadIdx.x, en lugar de la suma del valor numer calculada a través de todos los hilos en el bloque. ¿Alguien sabe lo que está pasando?

extern __shared__ float s_shared[];

float numer = //calculate numerator

s_shared[threadIdx.x] = numer;
s_shared[blockDim.x] += numer;
__syncthreads();

float denom = s_shared[blockDim.x];
float result = numer/denom;

"resultado" siempre debe estar entre 0 y 1 y debe sumar 1 a través del bloque, sino que es igual a 1,0 para cada hilo donde threadIdx.x es el máximo, y algún otro valor no se limita a la gama para los otros hilos en el bloque.

Answer 1

No está sincronizando el resumen adecuadamente a la blockDim.x ubicación. Ninguno de los hilos están esperando a ver lo que otros han escrito antes de añadir su suma. Algo así como

• Todo el mundo lee cero,
• se va a casa, calcula + numer cero.
• Everone escribe + numer cero a la posición de memoria

La alta ThreadID gana b / c que tiene una alta probabilidad de actuar pasado, supongo.

Lo que se quiere hacer en su lugar, con el fin de hacer una suma rápida, es hacer una suma binaria de s_shared[threadIdx.x]

• cada uno escribe su numer
• la mitad de los hilos de los cálculos de sumas de pares y escriben los a una nueva ubicación
• una cuarta parte de los hilos caluclate las sumas de pares de pares, y escribir los a una nueva ubicación
• etc.
• hasta que sólo hay un hilo y una suma

Esto toma O (n) de trabajo y tiempo O (log n).