“内存不足”错误除了缺乏可用全局内存之外还有其他原因吗？

Question

正如标题所说，在

cudaCreateTextureObject

调用后，我收到“内存不足”错误（通过标准 CUDA 错误检查），但是当我打印设备上的可用内存量时，它几乎全部空闲（11GB/12GB） 。有趣的是，只有当我使用完全相同的输入从 python/MATLAB 调用 CUDA 代码约 34K 次后，才会出现这种情况。我担心的是，有一些我不知道的“其他”内存正在被填满并且没有被释放。或者，此错误是由其他原因引起/引发的，这导致了标题中的问题。 我很高兴对标题中的问题有一个好的答案，但我认为最好是我提供问题的整个背景。这就是现在的情况：

有关代码的更多详细信息

---

现在，请耐心等待我分享更多细节。我无法从我的复杂代码中重现这一点，我对最小示例的尝试是徒劳的。标题中的问题在没有代码的情况下独立存在，因为我可以使用正确的代码和足够的可用全局内存来获取错误。但是，更多有关代码的上下文可能会有所帮助。

内存不足错误出现在我的纹理内存分配中。功能如下：

void CreateTexture(const GpuIds& gpuids, float* projectiondata,Geometry geo,cudaArray** d_cuArrTex,unsigned int nangles, cudaTextureObject_t *texImage,cudaStream_t* stream,int nStreamDevice,bool allocate){


    const cudaExtent extent =make_cudaExtent(geo.nDetecU, geo.nDetecV, nangles);
    const unsigned int num_devices = gpuids.GetLength();
    size_t memfree;
    size_t memtotal;
    if (allocate){
        for (unsigned int dev = 0; dev < num_devices; dev++){
            cudaSetDevice(gpuids[dev]);
            cudaDeviceSynchronize();
            cudaCheckErrors("before cudaMalloc3DArray fail");

            //cudaArray Descriptor
            cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();
            //cuda Array
            cudaMalloc3DArray(&d_cuArrTex[dev], &channelDesc, extent);
            cudaDeviceSynchronize();
            cudaCheckErrors("cudaMalloc3DArray fail");
        }
    }
    for (unsigned int dev = 0; dev < num_devices; dev++){
        cudaSetDevice(gpuids[dev]);
        cudaMemcpy3DParms copyParams = {0};
        //Array creation
        copyParams.srcPtr   = make_cudaPitchedPtr((void *)projectiondata, extent.width*sizeof(float), extent.width, extent.height);
        copyParams.dstArray = d_cuArrTex[dev];
        copyParams.extent   = extent;
        copyParams.kind     = cudaMemcpyHostToDevice;
        cudaMemcpy3DAsync(&copyParams,stream[dev*nStreamDevice+1]);
        cudaDeviceSynchronize();
        cudaCheckErrors("cudaMemcpy3DAsync fail");
    }

    //Array creation End
    for (unsigned int dev = 0; dev < num_devices; dev++){
        
        cudaSetDevice(gpuids[dev]);
        //cudaDeviceSynchronize();
        //cudaCheckErrors("cudaCreateTextureObject init fail");
        //cudaMemGetInfo(&memfree,&memtotal);
        //printf("Free memory: %zu\n",memfree);

        cudaResourceDesc    texRes;
        memset(&texRes, 0, sizeof(cudaResourceDesc));
        texRes.resType = cudaResourceTypeArray;
        texRes.res.array.array  = d_cuArrTex[dev];
        cudaTextureDesc     texDescr;
        memset(&texDescr, 0, sizeof(cudaTextureDesc));
        texDescr.normalizedCoords = false;
        texDescr.filterMode = cudaFilterModeLinear;
        texDescr.addressMode[0] = cudaAddressModeBorder;
        texDescr.addressMode[1] = cudaAddressModeBorder;
        texDescr.addressMode[2] = cudaAddressModeBorder;
        texDescr.readMode = cudaReadModeElementType;
        cudaCreateTextureObject(&texImage[dev], &texRes, &texDescr, NULL);

        //cudaMemGetInfo(&memfree,&memtotal);
        //printf("Free memory: %zu\n",memfree);
        cudaDeviceSynchronize();
        cudaCheckErrors("cudaCreateTextureObject fail");
    }

Geometry
• 只是一个包含大小等所需元数据的结构。
GpuIds
• 是一个类，其中包含允许此代码使用哪些 GPU。与此错误无关，因为它发生在 1 GPU PC 中。
否则，这是一个标准的纹理对象创建、复制和分配代码，唯一的怪癖是它允许多 GPU 代码（但到目前为止，该错误仅在 1 个 GPU 机器上重现）并且有一个布尔值

allocate

选择是否需要分配 3D 数组。

我已经使用这个代码很多年了，它似乎运行良好。它位于分别通过 MATLAB 或 Python 中的

mex

文件或

cython

文件调用的函数中，它应该分配内存、进行计算并完全释放 GPU。最近，有人注意到，在循环运行此代码

数千次

独立调用后（对于相同的输入大小，数字始终相同，但如果输入大小发生变化，则不同），代码崩溃并显示“内存不足” ”，我将其精确到了上面代码中的最后一个 cudaCheckErrors。但是，我们调用代码的方式仅使用了 12GB 可用内存中的 1GB，并且监视所有内存我看不到任何地方的增加。这使我相信该错误不一定是全局内存的“内存不足”。我想知道是否有一些特定的数组我应该释放但我没有释放，或者我意外地填充了不同的内存（共享？常量？（我认为这两个没有意义））。当然，调试这个问题会更困难，并且如果这个问题可能超出范围，因为它是一个复杂的代码。但我现在一无所知，因为我所拥有的唯一信息对我目前的知识没有帮助（“内存不足”）[^1]。

进一步的证据表明，我以某种方式填充了

some

内存，它不是全局的，如果我在每次调用后cudaDeviceReset()，这个错误就会消失，但代价是更长的执行时间。

完全重现性

---

在两台不同的机器上进行测试，导致完全相同的迭代次数出现错误。

安装

TIGRE

。出现错误的文件是 voxel_backprojection.cu ，第 667-714 行（上面代码中的那些）。您可以使用以下 python 代码重现此错误：

import numpy as np
import tigre
from skimage.data import shepp_logan_phantom
from tqdm import tqdm

def main():
    
    gt = shepp_logan_phantom().astype(np.float32)[None, ...]
    domain = gt.shape
    NANGLES = 1000
    angles = np.linspace(0, 2 * np.pi, NANGLES)
  
    x = np.zeros(domain, dtype=np.float32)
    geo = tigre.geometry(mode="fan", nVoxel=np.array(x.shape))
    ys = tigre.Ax(gt, geo, angles)
  

    max_iterations=1000*49 # Make it longer if the error doesnt happen in your GPU. Takes ~10 minutes in my machine to crash. 
    print(ys.shape[0])
    for k in tqdm(range(max_iterations), leave=False):
        x = tigre.Atb(ys, geo, angles)
            
if __name__ == "__main__":
    main()

这在我的机器上迭代 34679 时失败

总是

。

在 RTX 4070 上使用 CUDA V12.3.103

---

[^1] 而“内存不足”是 CUDA 中最难搜索的关键字，因为这都是人们真正分配了太多内存的错误！

---

Answer 1

话虽如此，这部分：

这在我的机器上总是在迭代 34679 时失败。

让我怀疑是否发生了其他事情，例如资源句柄耗尽或某些计数器溢出或运行时内存管理器中的其他问题。为了确定起见，我绝对会将其作为错误报告给 NVIDIA。

这两种情况下的对策都是不要破坏内存管理器。

不要重复分配和释放内存。要么在代码开始时的初始化步骤中分配内存缓冲区，并在应用程序的整个生命周期中将该分配保持为运行时状态（大量证据表明 CUFFT 和 CUBLAS 正是这样做的），要么使用

stream 有序分配器

并使用运行时管理的内存池进行操作。 无论您如何做，您都应该发现内存重用 (a) 缓解了您眼前的问题并且 (b) 提高了代码的性能。