3

テクスチャ メモリを使用して補間の問題を解決しようとしていますが、できればグローバル メモリを使用するよりも高速な方法です。テクスチャ メモリを使用するのは初めてなので、補間の問題を線形補間の問題に単純化しすぎています。したがって、以下に報告されている方法よりも、線形補間を行うためのよりスマートで高速な方法があることは既に認識しています。Kernels_Interpolation.cuh ファイルは次のとおりです。__device__ 関数 linear_kernel_GPU は簡単にするために省略されていますが、正しいです。

texture<cuFloatComplex,1> data_d_texture;

__global__ void linear_interpolation_kernel_function_GPU_texture(cuComplex* result_d, float* x_in_d, float* x_out_d, int M, int N)
{    
   int j = threadIdx.x + blockDim.x * blockIdx.x;

   cuComplex datum;

   if(j<N)
   {
       result_d[j] = make_cuComplex(0.,0.);
       for(int k=0; k<M; k++)
       {
           datum = tex1Dfetch(data_d_texture,k);
           if (fabs(x_out_d[j]-x_in_d[k])<1.) result_d[j] = cuCaddf(result_d[j],cuCmulf(make_cuComplex(linear_kernel_GPU(x_out_d[j]-x_in_d[k]),0.),datum));
       }  
   } 
}

Kernels_Interpolation.cu 関数は次のとおりです。

extern "C" void linear_interpolation_function_GPU_texture(cuComplex* result_d, cuComplex* data_d, float* x_in_d, float* x_out_d, int M, int N){

   cudaBindTexture(NULL, data_d_texture, data_d, M);

   dim3 dimBlock(BLOCK_SIZE,1); dim3 dimGrid(N/BLOCK_SIZE + (N%BLOCK_SIZE == 0 ? 0:1),1);
   linear_interpolation_kernel_function_GPU_texture<<<dimGrid,dimBlock>>>(result_d, x_in_d, x_out_d, M, N);

}

最後に、メイン プログラムでは、次のように data_d 配列が割り当てられ、初期化されます。

cuComplex* data_d;      cudaMalloc((void**)&data_d,sizeof(cuComplex)*M);
cudaMemcpy(data_d,data,sizeof(cuComplex)*M,cudaMemcpyHostToDevice);

result_d 配列の長さは N です。

奇妙なことに、出力は最初の 16 の場所でのみ正しく計算されますが、N>16 であり、他の場所は 0 などです。

result.r[0] 0.563585 result.i[0] 0.001251 
result.r[1] 0.481203 result.i[1] 0.584259
result.r[2] 0.746924 result.i[2] 0.820994
result.r[3] 0.510477 result.i[3] 0.708008
result.r[4] 0.362980 result.i[4] 0.091818
result.r[5] 0.443626 result.i[5] 0.984452
result.r[6] 0.378992 result.i[6] 0.011919
result.r[7] 0.607517 result.i[7] 0.599023
result.r[8] 0.353575 result.i[8] 0.448551
result.r[9] 0.798026 result.i[9] 0.780909
result.r[10] 0.728561 result.i[10] 0.876729
result.r[11] 0.143276 result.i[11] 0.538575
result.r[12] 0.216170 result.i[12] 0.861384
result.r[13] 0.994566 result.i[13] 0.993541
result.r[14] 0.295192 result.i[14] 0.270596
result.r[15] 0.092388 result.i[15] 0.377816
result.r[16] 0.000000 result.i[16] 0.000000
result.r[17] 0.000000 result.i[17] 0.000000
result.r[18] 0.000000 result.i[18] 0.000000
result.r[19] 0.000000 result.i[19] 0.000000

コードの残りの部分は正しいです。つまり、linear_interpolation_kernel_function_GPU_texture と linear_interpolation_function_GPU_texture をグローバル メモリを使用する関数に置き換えれば、すべて問題ありません。

特定の位置 (M と N に依存)、たとえば 64 までテクスチャ メモリに正しくアクセスできることを確認しました。その後、0 が返されます。

cuComplex テクスチャを float テクスチャに置き換えた場合 (データを実数に強制する場合)、同じ問題が発生します。

何か案は？

4

1 に答える 1

3

プログラムの次の行に1つの論理エラーがあります。

cudaBindTexture(NULL, data_d_texture, data_d, M);

の最後の引数はcudaBindTextureバイト単位のデータのサイズを取り、要素の数を指定しています。

次のことを試してください。

cudaBindTexture(NULL, data_d_texture, data_d, M * sizeof(cuComplex));
于 2012-12-25T21:12:00.790 に答える