cuda - cudaDeviceSynchronize 上的非法内存访问

Question

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

谢谢!

Answer 1

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

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

一些建议:

1. 找出您需要多少空间，并使用 cudaDeviceSetLimit(cudaLimitMallocHeapSize, size_t size)
提前在设备堆上预留空间
2. 当您遇到使用 new 或 malloc 的内核时，对于调试目的可能有用，也许使用调试宏来检查返回的指针无效的。这通常是一个很好的做法。
3. 您可以使用 here 中描述的方法更清楚地了解如何调试非法内存访问(将其定位到特定内核中的特定行) .