cudaDeviceSynchronize 上的非法内存访问答案

【问题标题】：Illegal Memory Access on cudaDeviceSynchronizecudaDeviceSynchronize 上的非法内存访问
【发布时间】：2015-04-02 01:58:19
【问题描述】：

我遇到了一个非常奇怪的错误，即在运行特定大小的 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 个元素时，内核函数位置和状态数组中的元素数量似乎会导致错误，这让我感到很沮丧。任何见解都会很棒。

谢谢！

【问题讨论】：

请注意，SO 希望您提供完整的MCVE 这样的问题（“为什么这段代码不起作用？”）。根据 SO 的期望，这并不是真正的可选。未能提供 MCVE 是投票结束此类问题的正当理由。

标签： cuda cuda-gdb

【解决方案1】：

我认为内核中的new 可能会失败，因为您分配了太多内存。

内核内new 具有与in-kernel malloc 类似的行为和限制。这些分配仅限于 设备堆，默认情况下以 8MB 开始。如果 250x250 数组大小对应于该范围 (8MB) 中的某个值，那么显着高于该值将导致一些新操作“静默”失败（即返回空指针）。如果您随后尝试使用这些空指针，您将获得非法内存访问。

一些建议：

计算出您需要多少空间，并使用cudaDeviceSetLimit(cudaLimitMallocHeapSize, size_t size) 提前在设备堆上预先保留它
当您在使用new 或malloc 的内核时遇到问题时，可能出于调试目的使用调试宏来检查返回的指针是否为NULL 可能很有用。总的来说，这是一个很好的做法。
您可以使用here 中描述的方法了解如何更清楚地调试非法内存访问（将其本地化到特定内核中的特定行）。

【讨论】：