2

特定のサイズの Heat 2D シミュレーションを実行すると「不正なメモリ アクセス」エラーが発生するという非常に奇妙なバグに遭遇していますが、要素が少ないだけでまったく同じシミュレーションを実行すると、シミュレーションは正常に実行されます。

配列のサイズを大きくするとこの例外が発生する理由はありますか? Titan Black GPU (メモリ 6 GB) を使用していますが、実行しているシミュレーションはそのサイズにはほど遠いものです。4000 x 4000 のシミュレーションを実行できると計算しましたが、250 x 250 を超えるとエラーになります。

デバイスでシミュレーション オブジェクトの配列をインスタンス化した直後にエラーが発生します。インスタンス化コードは次のとおりです。

template<typename PlaceType, typename StateType>
__global__ void instantiatePlacesKernel(Place** places, StateType *state,
        void *arg, int *dims, int nDims, int qty) {
    unsigned idx = blockDim.x * blockIdx.x + threadIdx.x;

    if (idx < qty) {
        // set pointer to corresponding state object
        places[idx] = new PlaceType(&(state[idx]), arg);
        places[idx]->setIndex(idx);
        places[idx]->setSize(dims, nDims);
    }
}

template<typename PlaceType, typename StateType>
Place** DeviceConfig::instantiatePlaces(int handle, void *argument, int argSize,
        int dimensions, int size[], int qty) {

    // add global constants to the GPU
    memcpy(glob.globalDims,size, sizeof(int) * dimensions);
    updateConstants(glob);

    // create places tracking
    PlaceArray p; // a struct to track qty, 
    p.qty = qty;

    // create state array on device
    StateType* d_state = NULL;
    int Sbytes = sizeof(StateType);
    CATCH(cudaMalloc((void** ) &d_state, qty * Sbytes));
    p.devState = d_state; // save device pointer

    // allocate device pointers
    Place** tmpPlaces = NULL;
    int ptrbytes = sizeof(Place*);
    CATCH(cudaMalloc((void** ) &tmpPlaces, qty * ptrbytes));
    p.devPtr = tmpPlaces; // save device pointer

    // handle arg if necessary
    void *d_arg = NULL;
    if (NULL != argument) {
        CATCH(cudaMalloc((void** ) &d_arg, argSize));
        CATCH(cudaMemcpy(d_arg, argument, argSize, H2D));
    }

    // load places dimensions
    int *d_dims;
    int dimBytes = sizeof(int) * dimensions;
    CATCH(cudaMalloc((void** ) &d_dims, dimBytes));
    CATCH(cudaMemcpy(d_dims, size, dimBytes, H2D));

    // launch instantiation kernel
    int blockDim = (qty - 1) / BLOCK_SIZE + 1;
    int threadDim = (qty - 1) / blockDim + 1;
    Logger::debug("Launching instantiation kernel");
    instantiatePlacesKernel<PlaceType, StateType> <<<blockDim, threadDim>>>(tmpPlaces, d_state,
            d_arg, d_dims, dimensions, qty);
    CHECK();

    CATCH(cudaDeviceSynchronize()); // ERROR OCCURS HERE

    // clean up memory
    if (NULL != argument) {
        CATCH(cudaFree(d_arg));
    }
    CATCH(cudaFree(d_dims));
    CATCH(cudaMemGetInfo(&freeMem, &allMem));

    return p.devPtr;
}

このコードは十分に小さいシミュレーションでエラーなしで実行されるため、表示されるカスタム型はすべて機能していると想定してください。サイズが 250 x 250 要素を超えると、カーネル関数の場所と状態配列の要素数がエラーを引き起こすように見えることに不満を感じています。どんな洞察も素晴らしいでしょう。

ありがとうございました！

4

1 に答える 1

9

newメモリを割り当てすぎているため、カーネル内で失敗している可能性が高いと思います。

カーネル内には、カーネル内 mallocnewと同様の動作と制限があります。これらの割り当ては、デフォルトで 8MB から始まるデバイス ヒープに制限されます。250x250 の配列サイズがその範囲 (8MB) に相当する場合、それを大幅に超えると、新しい操作の一部が「静かに」失敗します (つまり、null ポインターを返します)。これらのヌル ポインターを使用しようとすると、不正なメモリ アクセスが発生します。

いくつかの推奨事項:

  1. 必要なスペースの量を把握し、事前にデバイス ヒープに次のコマンドを使用して予約します。cudaDeviceSetLimit(cudaLimitMallocHeapSize, size_t size)
  2. newまたはを使用するカーネルで問題が発生した場合malloc、おそらくデバッグ マクロを使用して、返されたポインターの NULL をチェックすると、デバッグの目的で役立つ場合があります。これは一般的に良い方法です。
  3. こちら で説明されている方法を使用して、不正なメモリ アクセスをより明確に (特定のカーネルの特定の行にローカライズして) デバッグする方法を学習できます。
于 2015-02-03T01:30:31.880 に答える