cudaDeviceSynchronize 上的非法内存访问

Question

我遇到了一个非常奇怪的错误，当运行进行特定尺寸的 Heat 2D 模拟时，我得到一个 'illegal memory access' 错误，但是模拟运行如果我运行完全相同的模拟，只是元素更少。

是否有增加数组大小会导致此异常的原因？我使用的是 Titan Black GPU（6 GB 内存），但我运行ning 的模拟远不及那个大小。我计算过我可以运行模拟 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) 中的某个值，那么大大超过该范围将导致某些新操作 "silently" 失败（即 return 空指针）。如果您随后尝试使用这些空指针，您将获得非法的内存访问。

几点建议：

弄清楚你需要多少 space，并使用 cudaDeviceSetLimit(cudaLimitMallocHeapSize, size_t size)
当您在使用 new 或 malloc 的内核上遇到问题时，使用调试宏来检查 returned 指针可能对调试有用为空。这通常是一个很好的做法。
您可以使用 here.

cudaDeviceSynchronize 上的非法内存访问

Illegal Memory Access on cudaDeviceSynchronize

cuda

cuda-gdb