了解每个线程的 cuda 堆内存限制

这个问题是关于cuda中堆大小限制的。 访问了有关该主题的一些问题，包括这个：内核中的新运算符..奇怪的行为 https://stackoverflow.com/questions/13072624/new-operator-in-kernel-strange-behaviour/13073677#13073677我做了一些测试。给定一个内核如下：

#include <cuda.h>
#include <cuda_runtime.h>
#define CUDA_CHECK( err ) __cudaSafeCall( err, __FILE__, __LINE__ )
#define CUDA_CHECK_ERROR()    __cudaCheckError( __FILE__, __LINE__ )
inline void __cudaSafeCall( cudaError err, const char *file, const int line )
{
    if ( cudaSuccess != err )
    {
        fprintf( stderr, "cudaSafeCall() failed at %s:%i : %s\n",
                file, line, cudaGetErrorString( err ) );
        exit( -1 );
    }
    return;
}

inline void __cudaCheckError( const char *file, const int line )
{
    cudaError err = cudaGetLastError();
    if ( cudaSuccess != err )
    {
        fprintf( stderr, "cudaCheckError() failed at %s:%i : %s\n",
                file, line, cudaGetErrorString( err ) );
        exit( -1 );
    }
    return;
}
#include <stdio>
#define NP 900000
__device__ double *temp;
__device__ double *temp2;
__global__
void test(){
int i = blockDim.x*blockIdx.x + threadIdx.x;
if(i==0){
    temp = new double[NP];
    //temp2 = new double[NP];
}

if(i==0){
    for(int k=0;k<NP;k++){
        temp[i] = 1.;
        if(k%1000 == 0){
            printf("%d : %g\n", k, temp[i]);
        }
    }
}
if(i==0){
    delete(temp);
    //delete(temp2);
}
}
int main(){
    //cudaDeviceSetLimit(cudaLimitMallocHeapSize, 32*1024*1024);
    //for(int k=0;k<2;k++){
        test<<<ceil((float)NP/512), 512>>>();
        CUDA_CHECK_ERROR();
    //}
    return 0;
}

我想测试堆大小限制。

1. 为一个线程动态分配一个数组（temp），其大小为 大约超过 960,000*sizeof(double) （接近 8MB，这是 堆大小的默认限制）给出错误：好的。 900,000 幅作品。 （有人知道如何计算真实极限吗？）
2. 提高堆大小限制允许分配更多内存：正常，好的。
3. 回到 8MB 堆大小，为两个线程的每个线程分配一个数组（因此，将 if (i==0) 替换为 if(i==0 || i==1)，每个 900,000 * sizeof(double) 失败.但是每个 450,000*sizeof(double) 都可以.仍然可以.
4. 我的问题来了：用一个线程分配两个数组（因此，线程 0 的 temp 和 temp2），每个 900,000 * sizeof(double) 也可以工作，但它不应该吗？事实上，当我尝试写入两个数组时，它失败了。但是任何人都知道为什么在使用两个数组和一个线程而不是两个数组和两个线程时分配时会出现这种不同的行为？

编辑：另一个测试，我发现对于像我一样学习堆用法的人来说很有趣： 5. 执行内核两次，单线程0分配一个大小为900,000 * sizeof(double)的数组，如果有删除，就可以工作。如果省略delete，第二次会失败，但第一次调用会执行。

编辑2：如何分配一次设备范围的变量，但可由所有线程写入（不是来自主机，在设备代码中使用动态分配）？

可能您没有测试返回的空指针new操作，这是一个C++ 中操作员报告失败的有效方法 http://www.cplusplus.com/reference/new/operator%20new%5b%5d/.

当我按如下方式修改代码时，我收到消息“第二个新失败”：

#include <stdio.h>

 #define NP 900000
__device__ double *temp;
__device__ double *temp2;
__global__
void test(){
int i = blockDim.x*blockIdx.x + threadIdx.x;
if(i==0){
    temp = new double[NP];
    if (temp == 0) {printf("first new failed\n"); return;}
    temp2 = new double[NP];
    if (temp2 == 0) {printf("second new failed\n"); return;}
}

if(i==0){
    for(int k=0;k<NP;k++){
        temp[i] = 1.;
        if(k%1000 == 0){
            printf("%d : %g\n", k, temp[i]);
        }
    }
}
if(i==0){
    delete(temp);
    delete(temp2);
}
}

int main() {

  test<<<1,1>>>();
  cudaDeviceSynchronize();
  return 0;
}

如果您提供完整的、可编译的代码供其他人使用，就像我一样，那就很方便了。

对于您的第一个编辑问题，如果删除第一个问题，第二个新问题将起作用并不奇怪。第一个分配几乎所有可用的 8MB。如果删除该分配，则第二个分配将会成功。参考文档 http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#dynamic-global-memory-allocation-and-operations，我们看到以这种方式动态分配的内存在 cuda 上下文的整个生命周期内都存在，或者直到执行相应的删除操作为止（即不仅仅是单个内核调用。内核完成并不一定会释放分配。）

对于你的第二个编辑问题，你已经在演示一种方法，使用你的__device__ double *temp;指针，一个线程可以通过该指针分配所有线程都可以访问的存储空间。但是，跨块会遇到问题，因为无法保证块之间的同步顺序或块之间的执行顺序，因此如果从块 0 中的线程 0 进行分配，则仅当块 0 在其他块之前执行时才有用。您可以想出一个复杂的方案来检查变量分配是否已经完成（也许通过测试 NULL 指针，也可能使用原子），但它会创建脆弱的代码。最好提前计划全局分配并从主机进行相应分配。