質問に対する私の理解に基づいて、N個の異なる信号の平均値を格納するためにN個のレジスタは必要ないと言います。
すでに N 個のスレッドがある場合 [各スレッドが 1 つのシグナルのみでリダクションを行う場合]、1 つのシグナルのリダクションを格納するために N 個のレジスタは必要ありません。必要なのは、平均値を格納するためのレジスタ 1 つだけです。
dim3 threads (N,1);
reduction<<<threads,1>>>(signals); // signals is the [N*M] array
__global__ reduction (int *signals)
{
int id = threadIdx.x;
float meanValue = 0.0;
for(int i = 0; i < M; i++)
meanValue = signals[id*M +i];
meanValue = meanValue/M;
// Then do the subtraction
for(int i = 0; i < M; i++)
signals[id*M +i] -= meanValue;
}
N 個の異なる信号のすべての meanValues の一種のグローバル リダクションを実行する必要がある場合は、2 つのレジスタ [1 つはローカル平均を格納し、もう 1 つはグローバル平均を格納する] と共有メモリを使用する必要があります。
dim3 threads (N,1);
reduction<<<threads,1>>>(signals); // signals is the [N*M] array
__global__ reduction (int *signals)
{
__shared__ float means[N]; // shared value
int id = threadIdx.x;
float meanValue = 0.0;
float globalMean = 0.0;
for(int i = 0; i < M; i++)
meanValue += signals[id*M +i];
means[id] = meanValue/M;
__syncthreads();
// do the global reduction
for(int i = 0; i < N; i++)
globalMean += means[i];
globalMean = globalMean/N;
// Then do the subtraction
for(int i = 0; i < M; i++)
signals[id*M +i] -= globalMean;
}
これがお役に立てば幸いです。ご不明な点がございましたら、お知らせください。