配列を使用した操作で値が破損するのはなぜですか？

Question

私はParticle Swarm OptimizationをCUDAに実装しようとしています。ホスト上のデータ配列を部分的に初期化してから、CUDAにメモリを割り当ててそこにコピーしてから、初期化を続行します。

問題は、私はそう

__global__ void kernelInit(
    float* X, 
    size_t pitch, 
    int width, 
    float X_high, 
    float X_low 
) { 
    // Silly, but pretty reliable way to address array elements 
    unsigned int tid = blockIdx.x * blockDim.x + threadIdx.x; 
    int r = tid/width; 
    int c = tid % width; 
    float* pElement = (float*)((char*)X + r * pitch) + c; 
    *pElement = *pElement * (X_high - X_low) - X_low; 
    //*pElement = (X_high - X_low) - X_low; 
} 

のような配列要素を変更しようとしているとき、それは価値を破壊し、配列の要素として私1.#INF00与え、です。最後の行のコメントを外して、先にコメントすると、期待どおりに動作します。15.36などの値が得られます。

私の問題は、私のメモリの割り当てとコピー、および/または特定の配列要素の指定です。私はこれらの両方のトピックについてのCUDAのマニュアルを読んでいますが、私はエラーを見つけることはできません：私はまだ何かを壊れた配列を取得配列の要素です。たとえば、*pElement = *pElement * 2は、pElementが[0;1]にちょうど浮動小数点であると予想される場合、779616...00000000.00000のような不当な大きな結果をもたらします。

ここは完全な情報源です。アレイの初期化main（ソースの底部）で始まり、その後f1関数はCUDAのために作業を行い、初期化カーネルkernelInitを起動：

#include <stdio.h> 
#include <stdlib.h> 
#include <time.h> 
#include <math.h> 
#include <cuda.h> 
#include <cuda_runtime.h> 

const unsigned f_n = 3; 
const unsigned n = 2; 
const unsigned p = 64; 

typedef struct { 
    unsigned k_max; 
    float c1; 
    float c2; 
    unsigned p; 
    float inertia_factor; 
    float Ef; 
    float X_low[f_n]; 
    float X_high[f_n]; 
    float X_min[n][f_n]; 
} params_t; 

typedef void (*kernelWrapperType) (
    float *X, 
    float *X_highVec, 
    float *V, 
    float *X_best, 
    float *Y, 
    float *Y_best, 
    float *X_swarmBest, 
    bool &termination, 
    const float &inertia, 
    const params_t *params, 
    const unsigned &f 
); 

typedef float (*twoArgsFuncType) (
    float x1, 
    float x2 
); 

__global__ void kernelInit(
    float* X, 
    size_t pitch, 
    int width, 
    float X_high, 
    float X_low 
) { 
    // Silly, but pretty reliable way to address array elements 
    unsigned int tid = blockIdx.x * blockDim.x + threadIdx.x; 
    int r = tid/width; 
    int c = tid % width; 
    float* pElement = (float*)((char*)X + r * pitch) + c; 
    *pElement = *pElement * (X_high - X_low) - X_low; 
    //*pElement = (X_high - X_low) - X_low; 
} 

__device__ float kernelF1(
    float x1, 
    float x2 
) { 
    float y = pow(x1, 2.f) + pow(x2, 2.f); 
    return y; 
} 

void f1(
    float *X, 
    float *X_highVec, 
    float *V, 
    float *X_best, 
    float *Y, 
    float *Y_best, 
    float *X_swarmBest, 
    bool &termination, 
    const float &inertia, 
    const params_t *params, 
    const unsigned &f 
) { 
    float *X_d = NULL; 
    float *Y_d = NULL; 
    unsigned length = n * p; 
    const cudaChannelFormatDesc desc = cudaCreateChannelDesc<float4>(); 
    size_t pitch; 
    size_t dpitch; 
    cudaError_t err; 
    unsigned width = n; 
    unsigned height = p; 

    err = cudaMallocPitch (&X_d, &dpitch, width * sizeof(float), height); 
    pitch = n * sizeof(float); 
    err = cudaMemcpy2D(X_d, dpitch, X, pitch, width * sizeof(float), height, cudaMemcpyHostToDevice); 

    err = cudaMalloc (&Y_d, sizeof(float) * p); 
    err = cudaMemcpy (Y_d, Y, sizeof(float) * p, cudaMemcpyHostToDevice); 

    dim3 threads; threads.x = 32; 
    dim3 blocks; blocks.x = (length/threads.x) + 1; 

    kernelInit<<<threads,blocks>>>(X_d, dpitch, width, params->X_high[f], params->X_low[f]); 

    err = cudaMemcpy2D(X, pitch, X_d, dpitch, n*sizeof(float), p, cudaMemcpyDeviceToHost); 
    err = cudaFree(X_d); 

    err = cudaMemcpy(Y, Y_d, sizeof(float) * p, cudaMemcpyDeviceToHost); 
    err = cudaFree(Y_d); 
} 

float F1(
    float x1, 
    float x2 
) { 
    float y = pow(x1, 2.f) + pow(x2, 2.f); 
    return y; 
} 

/* 
* Generates random float in [0.0; 1.0] 
*/ 
float frand(){ 
    return (float)rand()/(float)RAND_MAX; 
} 

/* 
* This is the main routine which declares and initializes the integer vector, moves it to the device, launches kernel 
* brings the result vector back to host and dumps it on the console. 
*/ 
int main() { 
    const params_t params = { 
     100, 
     0.5, 
     0.5, 
     p, 
     0.98, 
     0.01, 
     {-5.12, -2.048, -5.12}, 
     {5.12, 2.048, 5.12}, 
     {{0, 1, 0}, {0, 1, 0}} 
    }; 
    float X[p][n]; 
    float X_highVec[n]; 
    float V[p][n]; 
    float X_best[p][n]; 
    float Y[p] = {0}; 
    float Y_best[p] = {0}; 
    float X_swarmBest[n]; 

    kernelWrapperType F_wrapper[f_n] = {&f1, &f1, &f1}; 
    twoArgsFuncType F[f_n] = {&F1, &F1, &F1}; 

    for (unsigned f = 0; f < f_n; f++) { 
     printf("Optimizing function #%u\n", f); 

     srand (time(NULL)); 
     for (unsigned i = 0; i < p; i++) 
      for (unsigned j = 0; j < n; j++) 
       X[i][j] = X_best[i][j] = frand(); 
     for (int i = 0; i < n; i++) 
      X_highVec[i] = params.X_high[f]; 
     for (unsigned i = 0; i < p; i++) 
      for (unsigned j = 0; j < n; j++) 
       V[i][j] = frand(); 
     for (unsigned i = 0; i < p; i++) 
      Y_best[i] = F[f](X[i][0], X[i][1]); 
     for (unsigned i = 0; i < n; i++) 
      X_swarmBest[i] = params.X_high[f]; 
     float y_swarmBest = F[f](X_highVec[0], X_highVec[1]); 

     bool termination = false; 
     float inertia = 1.; 

     for (unsigned k = 0; k < params.k_max; k++) { 
      F_wrapper[f]((float *)X, X_highVec, (float *)V, (float *)X_best, Y, Y_best, X_swarmBest, termination, inertia, &params, f); 
     } 

     for (unsigned i = 0; i < p; i++) 
     { 
      for (unsigned j = 0; j < n; j++) 
      { 
       printf("%f\t", X[i][j]); 
      } 
      printf("F = %f\n", Y[i]); 
     } 
     getchar(); 
    } 
} 

更新：私は

err = cudaMallocPitch (&X_d, &dpitch, width * sizeof(float), height); 
if (err != cudaSuccess) { 
    fprintf(stderr, cudaGetErrorString(err)); 
    exit(1); 
} 

ようにエラー処理を追加しようとしましたAPI呼び出し後に

が返されましたが、は返されませんでした。（まだすべての結果とプログラムが終了します）。

Answer 1

これは、単純なREPRO場合どうあるべきかのコードの不必要に複雑な作品ですが、これはすぐに飛び出し：

const unsigned n = 2; 
const unsigned p = 64; 

unsigned length = n * p 

dim3 threads; threads.x = 32; 
dim3 blocks; blocks.x = (length/threads.x) + 1; 

kernelInit<<<threads,blocks>>>(X_d, dpitch, width, params->X_high[f], params->X_low[f]); 

ですから、まずブロックの数が誤って計算して、順序を逆にしています1つのグリッドあたりのブロック数と、カーネルの起動時の1ブロックあたりのスレッド数です。これは、GPUメモリに何かをぶつけたり、不特定の起動失敗を引き起こしたりして、エラー処理が不足している可能性があります。 cuda-memcheckというツールがあり、CUDA 3.0以降、ツールキットに同梱されています。それを実行すると、valgrindスタイルのメモリアクセス違反レポートが表示されます。あなたがまだそうしていないなら、あなたはそれを使う習慣に入るべきです。

無限の値は、それが期待されるのですか？コードは、（0,1）の値で開始し、次いで

X[i] = X[i] * (2.048--2.048) - -2.048 

100に続いて10^100を掛けるの粗い等価である

X[i] = X[i] * (5.12--5.12) - -5.12 

100回、い最終的には4^100を乗算するのに相当する時間であり、最後に

X[i] = X[i] * (5.12--5.12) - -5.12 

が再び続きます。したがって、結果は1E250のオーダになるはずです。これは最大3よりはるかに大きくなります。IEEE 754単精度における表現可能な数値の大まかな上限である4E38を使用する。