Visual Studio 2010 を使用して CUDA で FFT の実装を開発しようとしていますが、これまでのところ、1 つのブロック内で 1024 ポイントまで動作させることができました。問題は、複数のブロックを使用するたびに、ブロック 1 の結果は問題なく、他のブロックは間違った値を返すことです (ランダムではないようで、複数回実行しても変化しません)。これが私のカーネルです。
__device__ void FFT(int idxS,int bfsize, Complex* data1, Complex* data0, int k, int N ){
Complex alpha;
if((idxS % bfsize) < (bfsize/2)){
data1[idxS] = ComplexAdd(data0[idxS],data0[idxS+bfsize/2]);
}
else
{
float angle = -PI*2*((idxS*(1<<k)%(bfsize/2)))/N;
alpha.x = cos(angle);
alpha.y= sin(angle);
Complex v0;
v0 = ComplexAdd(data0[idxS-bfsize/2] ,ComplexScale(data0[idxS],-1));
data1[idxS] = ComplexMul(v0, alpha);
}
}
__device__ void Ordenador(int r, int idxS ,Complex* data1, Complex* data0 ){
int p = 0;
for(int k = 0;k < r;k++)
{
if(idxS & (1<<k))
p+=1<<(r - k - 1);
}
data1[idxS] = data0[p];
__syncthreads();
}
__global__ void GPU_FFT(int N, int r, Complex* data0, Complex* data1, int k) {
int idxS = threadIdx.x+ blockIdx.x * blockDim.x;
__syncthreads;
int bfsize = 1<<(r - k);
FFT(idxS, bfsize, data1, data0, k, N);
data0[idxS] = data1[idxS];
}
int prepFFT(float *Entrada, Complex* saida, int N ){
if(ceilf(log2((float)N)) == log2((float)N) ){
for (int i=0; i<N; i++){
saida[i].x = Entrada[i];
saida[i].y = 0;
}
Complex *d_saida;
int m = (int)log2((float)N);
Complex *data1 = new Complex[N];
Complex *data1_d;
if (N<1024){
HANDLE_ERROR (cudaMalloc((void**)&d_saida, sizeof(Complex) * N));
HANDLE_ERROR (cudaMemcpy(d_saida,saida, sizeof(Complex)*N, cudaMemcpyHostToDevice));
HANDLE_ERROR (cudaMalloc((void**)&data1_d, sizeof(Complex) * N));
HANDLE_ERROR (cudaMemcpy(data1_d,data1, sizeof(Complex)*N, cudaMemcpyHostToDevice));
const dim3 numThreads (N,1,1);
const dim3 numBlocks(1,1,1);
for(int k = 0 ;k < m ; k++)
{
GPU_FFT<<<numBlocks,numThreads, N*2>>>( N, m, d_saida, data1_d, k);
HANDLE_ERROR (cudaDeviceSynchronize());
}
HANDLE_ERROR (cudaDeviceSynchronize());
HANDLE_ERROR (cudaMemcpy(saida,data1_d, sizeof(Complex)*N, cudaMemcpyDeviceToHost));
HANDLE_ERROR (cudaDeviceSynchronize());
}
else{
HANDLE_ERROR (cudaMalloc((void**)&d_saida, sizeof(Complex) * N));
HANDLE_ERROR (cudaMemcpy(d_saida,saida, sizeof(Complex)*N, cudaMemcpyHostToDevice));
HANDLE_ERROR (cudaMalloc((void**)&data1_d, sizeof(Complex) * N));
HANDLE_ERROR (cudaMemcpy(data1_d,data1, sizeof(Complex)*N, cudaMemcpyHostToDevice));
const dim3 numThreads (1024,1,1);
const dim3 numBlocks(N/1024 +1,1,1);
for(int k = 0;k < m;k++)
{
GPU_FFT<<<numBlocks,numThreads, N*2>>>( N, m, d_saida, data1_d, k);
HANDLE_ERROR (cudaDeviceSynchronize());
}
HANDLE_ERROR (cudaMemcpy(saida,data1_d, sizeof(Complex)*N, cudaMemcpyDeviceToHost));
HANDLE_ERROR (cudaDeviceSynchronize());
cudaFree(data1_d);
cudaFree(d_saida);
delete data1;
}
return 1;
}
else
return 0;
}
共有メモリを使用してみましたが、すべて 0 が返され、CUDA がグローバルから共有にコピーされていないと考えました (NSight は、そのメモリ位置の値が ???? であると教えてくれます)。このコードは、現時点では概念実証にすぎません。最適化する必要はありません。正しい値を返すだけです。完全なコードが必要な場合は、提供します。私はこれに対する解決策を1か月以上探していましたが、これは私の必死の呼びかけです.
ありがとう、ジョン
- - - - アップデート - - - -
2 つのブロックのそれぞれで 2 つのスレッドを起動するデバッグ目的でコードを変更しました。
int prepFFT(float *Entrada, Complex* saida, int N ){
if(ceilf(log2((float)N)) == log2((float)N) ){
for (int i=0; i<N; i++){
saida[i].x = Entrada[i];
saida[i].y = 0;
}
Complex *d_saida;
int m = (int)log2((float)N);
Complex *data1 = new Complex[N];
Complex *data1_d;
if (N<1024){
HANDLE_ERROR (cudaMalloc((void**)&d_saida, sizeof(Complex) * N));
HANDLE_ERROR (cudaMemcpy(d_saida,saida, sizeof(Complex)*N, cudaMemcpyHostToDevice));
HANDLE_ERROR (cudaMalloc((void**)&data1_d, sizeof(Complex) * N));
HANDLE_ERROR (cudaMemcpy(data1_d,data1, sizeof(Complex)*N, cudaMemcpyHostToDevice));
const dim3 numThreads (2,1,1);
const dim3 numBlocks(2,1,1);
for(int k = 0 ;k < m ; k++)
{
GPU_FFT<<<numBlocks,numThreads, N*2>>>( N, m, d_saida, data1_d, k);
HANDLE_ERROR (cudaDeviceSynchronize());
}
HANDLE_ERROR (cudaDeviceSynchronize());
HANDLE_ERROR (cudaMemcpy(saida,data1_d, sizeof(Complex)*N, cudaMemcpyDeviceToHost));
HANDLE_ERROR (cudaDeviceSynchronize());
}
else{
HANDLE_ERROR (cudaMalloc((void**)&d_saida, sizeof(Complex) * N));
HANDLE_ERROR (cudaMemcpy(d_saida,saida, sizeof(Complex)*N, cudaMemcpyHostToDevice));
HANDLE_ERROR (cudaMalloc((void**)&data1_d, sizeof(Complex) * N));
HANDLE_ERROR (cudaMemcpy(data1_d,data1, sizeof(Complex)*N, cudaMemcpyHostToDevice));
const dim3 numThreads (1024,1,1);
const dim3 numBlocks(N/1024 +1,1,1);
for(int k = 0;k < m;k++)
{
GPU_FFT<<<numBlocks,numThreads, N*2>>>( N, m, d_saida, data1_d, k);
HANDLE_ERROR (cudaDeviceSynchronize());
}
HANDLE_ERROR (cudaMemcpy(saida,data1_d, sizeof(Complex)*N, cudaMemcpyDeviceToHost));
HANDLE_ERROR (cudaDeviceSynchronize());
cudaFree(data1_d);
cudaFree(d_saida);
delete data1;
}
return 1;
}
else
return 0;
}
---------------------編集2 ---------------------
本当に奇妙なのは、memcheck を (任意のモードで) 使用すると、プログラムが正しい結果を返すことです。
----最終編集 ---------------
問題はこのコードにあることがわかりました
FFT(idxS, bfsize, data1, data0, k, N);
data0[idxS] = data1[idxS];
新しい関数の最後の行を分離し、それを CPU で呼び出すと、正しい結果が得られることがわかりました。お手伝いありがとう!!よろしくお願いします!