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) 中的某个值，那么显着高于该值将导致一些新操作“静默”失败（即返回空指针）。如果您随后尝试使用这些空指针，您将获得非法内存访问。

一些建议：

  1. 弄清楚你需要多少空间，并提前在设备堆上使用cudaDeviceSetLimit(cudaLimitMallocHeapSize, size_t size)
  2. 当您在使用newor的内核时遇到问题时，使用malloc调试宏检查返回的指针是否为 NULL 可能对调试有用。一般来说，这是一个很好的做法。
  3. 您可以使用此处描述的方法了解如何更清楚地调试非法内存访问（将其本地化到特定内核中的特定行）。
于 2015-02-03T01:30:31.880 回答