每个线程了解CUDA堆内存的限制(Understanding cuda heap memory li

这个问题是关于CUDA堆大小限制。我曾参观过有关这个话题的一些问题，包括这一个：在内核中新的运营商..奇怪的行为，我做了一些测试。给定一个内核如下：

#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;
}

我想测试堆大小限制。

动态地分配一个阵列（温度）与一种螺纹，其大小大约是超过96万*的sizeof（双）（接近8MB，这是堆大小的缺省极限）给出一个错误：确定。 900000幅作品。（没有人知道如何计算真正的限制吗？）
上升堆的大小限制允许分配更多的内存：正常，确定。
回到一个8MB堆大小，分配每个线程一个阵列与两个线程（因此，如果替换（I == 0）如果由（ⅰ== 0 ||我== 1），每一个900000 *的sizeof（双）失败但45万*的sizeof（双）各自的作品。还行。
这里说到我的问题：有一个线程（因此，温度和TEMP2线程0）分配两个阵列，每个阵列900000 *的sizeof（双）工作过，但它不应该？事实上，当我尝试在两个数组来写，它失败。但是，任何人都使用同一个线程，而不是两个阵列，两个线程两个数组时有一个想法，为什么在这个分配不同的行为？

编辑：另一个测试，我觉得这对于那些谁像我一样，会学习堆的使用有趣：5.执行内核的两倍，大小为90万*的sizeof（双）由单个线程0分配的一个阵列，如果有删除的作品。如果删除被省略，它会失败，第二次，但在第一次调用将被执行。

编辑2：如何通过所有线程一次，但可写分配设备范围内的变量（不是从主机上，使用在设备代码动态分配）？

也许你不是测试对在返回的空指针new操作，这是一个在C ++操作员报告故障的有效方法。

当我修改您的代码如下，我得到“第二个新的失败”的消息：

#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可用。如果删除分配，则第二个会成功。参考文档，我们可以看到以这种方式动态分配内存活的CUDA上下文的整个生命周期，或直到执行相应的删除操作（即不只是一个单一的内核调用。内核的完成并不一定免费分配。）

关于你的第二编辑的问题，你已经证明的方法，使用__device__ double *temp; 指针，由其中一个线程可以分配存储所有线程都可以访问。您将有跨越块有问题，但是，由于没有同步为了之中块或块之间的执行顺序的保障，因此，如果您从线程0 0块分配，如果块0以外块之前执行，这只是有用的。你能想出一个复杂的方案，以检查是否可变的分配已经完成（也许通过测试指针NULL，也可能使用原子能），但它创建脆弱的代码。最好是提前计划你的全球分配和从主机相应地分配。