CUDA matriz de memória compartilhada - comportamento estranho

Question

Em um kernel CUDA, tenho código semelhante ao seguinte. Eu estou tentando calcular um numerador por thread, e acumular os numeradores sobre o bloco para calcular um denominador e, em seguida, retornar a relação. No entanto, CUDA está definindo o valor da denom para qualquer valor é calculado para numer pelo segmento no bloco com o maior threadIdx.x, em vez do que a soma do valor numer calculada em todos os tópicos no bloco. Alguém sabe o que está acontecendo?

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" deve ser sempre entre 0 e 1 e deve somar 1 ao longo de bloco, mas em vez disso, é igual a 1,0 para cada segmento onde threadIdx.x é o máximo, e algum outro valor não confinado ao intervalo para os outros fios no bloco.

Answer 1

Você não está sincronizando a soma corretamente para o local blockDim.x. Nenhum dos segmentos estão aguardando para ver o que outros têm escrito antes de adicionar sua soma. Mais ou menos como

• Todos lê zero,
• vai para casa, calcula zerar + numer.
• Everone escreve de zero + numer para a localização de memória

A alta ThreadID vitórias b / c ele tem uma alta probabilidade de agir por último, eu suponho.

O que você quer fazer em vez disso, a fim de fazer uma soma rápida, é fazer uma soma binária s_shared[threadIdx.x]

• todos escreve sua numer
• metade dos tópicos calcular somas de pares e escrever os para um novo local
• um quarto dos tópicos caluclate as somas de pares de pares, e escrever os para um novo local
• etc
• até que você só tem um segmento e um sum

Isso leva tempo O (n) trabalho e O (log n).