2

私のコードはエラーメッセージを表示していて、その原因を突き止めようとしています。問題を見つけやすくするために、エラーメッセージの発生に明らかに関係のないコードを削除しました。次の簡単なコードでエラーメッセージが表示される理由を教えていただければ、元のコードを修正できるはずです。

#include "cuComplex.h"
#include <cutil.h>

__device__ void compute_energy(void *data, int isample, int nsamples) {
  cuDoubleComplex * const nminusarray          = (cuDoubleComplex*)data;
  cuDoubleComplex * const f                    = (cuDoubleComplex*)(nminusarray+101);
  double          * const abs_est_errorrow_all = (double*)(f+3);
  double          * const rel_est_errorrow_all = (double*)(abs_est_errorrow_all+nsamples*51);
  int             * const iid_all              = (int*)(rel_est_errorrow_all+nsamples*51);
  int             * const iiu_all              = (int*)(iid_all+nsamples*21);
  int             * const piv_all              = (int*)(iiu_all+nsamples*21);
  cuDoubleComplex * const energyrow_all        = (cuDoubleComplex*)(piv_all+nsamples*12);
  cuDoubleComplex * const refinedenergyrow_all = (cuDoubleComplex*)(energyrow_all+nsamples*51);
  cuDoubleComplex * const btplus_all           = (cuDoubleComplex*)(refinedenergyrow_all+nsamples*51);

  cuDoubleComplex * const btplus           = btplus_all+isample*21021;

  btplus[0] = make_cuDoubleComplex(0.0, 0.0);
}

__global__ void computeLamHeight(void *data, int nlambda) {
  compute_energy(data, blockIdx.x, nlambda);
}

int main(int argc, char *argv[]) {
  void *device_data;

  CUT_DEVICE_INIT(argc, argv);
  CUDA_SAFE_CALL(cudaMalloc(&device_data, 184465640));
  computeLamHeight<<<dim3(101, 1, 1), dim3(512, 1, 1), 45000>>>(device_data, 101);
  CUDA_SAFE_CALL(cudaThreadSynchronize());
}

私はGeForceGTX480を使用しており、次のようにコードをコンパイルしています。

nvcc -L /soft/cuda-sdk/4.0.17/C/lib -I /soft/cuda-sdk/4.0.17/C/common/inc -lcutil_x86_64 -arch sm_13 -O3 -Xopencc "-Wall" Main.cu

出力は次のとおりです。

Using device 0: GeForce GTX 480
Cuda error in file 'Main.cu' in line 31 : unspecified launch failure.

編集：コードをさらに簡略化しました。次のより単純なコードでも、エラーメッセージが生成されます。

#include <cutil.h>

__global__ void compute_energy(void *data) {
  *(double*)((int*)data+101) = 0.0;
}

int main(int argc, char *argv[]) {
  void *device_data;

  CUT_DEVICE_INIT(argc, argv);
  CUDA_SAFE_CALL(cudaMalloc(&device_data, 101*sizeof(int)+sizeof(double)));
  compute_energy<<<dim3(1, 1, 1), dim3(1, 1, 1)>>>(device_data);
  CUDA_SAFE_CALL(cudaThreadSynchronize());
}

これで、オフセットが有効であることが簡単にわかります。cuda-memcheckを実行してみましたが、次のように表示されます。

========= CUDA-MEMCHECK
Using device 0: GeForce GTX 480
Cuda error in file 'Main.cu' in line 13 : unspecified launch failure.
========= Invalid __global__ write of size 8
=========     at 0x00000020 in compute_energy
=========     by thread (0,0,0) in block (0,0,0)
=========     Address 0x200200194 is misaligned
=========
========= ERROR SUMMARY: 1 error

インターネットを検索して、アドレスがずれているとはどういう意味かを調べてみましたが、説明が見つかりませんでした。契約は何ですか？

4

2 に答える 2

8

これらすべての魔法の定数を使用して元のコードを解析することは非常に困難でしたが、更新された再現ケースにより、問題がすぐに明らかになります。GPUアーキテクチャでは、すべてのポインターを単語の境界に揃える必要があります。カーネルに、単語が正しく整列されていないポインタアクセスが含まれています。Doubleは64ビットタイプであり、アドレス指定は64ビット境界にさえ整列されていません。これ：

*(double*)((int*)data+100) = 0.0; // 50th double

またはこれ：

*(double*)((int*)data+102) = 0.0; // 51st double

どちらも合法です。これ：

*(double*)((int*)data+101) = 0.0; // not aligned to a 64 bit boundary

ではありません。

于 2012-08-06T17:13:05.463 に答える
2

エラーはアウトオブバウンドメモリアクセスを示しています。オフセット値を確認してください。

于 2012-08-06T04:35:07.283 に答える