Понимание ограничений памяти кучи cuda на поток

Этот вопрос касается ограничения размера кучи в 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;
}

Я хочу проверить ограничение размера кучи.

1. Динамическое выделение одного массива (temp) с одним потоком, размер которого примерно превышает 960 000 * sizeof (double) (около 8 МБ, что является пределом размера кучи по умолчанию), дает ошибку: ok. 900 000 работ. (кто-то знает, как рассчитать истинный предел?)
2. Увеличение предела размера кучи позволяет выделить больше памяти: нормально, хорошо.
3. Вернуться к размеру кучи 8 МБ, выделяя один массив на поток с ДВУМЯ потоками (таким образом, заменив if (i==0) на if(i==0 || i==1), каждый из 900 000 * sizeof(double) терпит неудачу , Но 450 000*sizeof(double) каждый работает.
4. Вот моя проблема: выделение ДВУХ массивов с ОДНИМ потоком (таким образом, temp и temp2 для потока 0), каждое 900 000 * sizeof(double) тоже работает, но не должно? Действительно, когда я пытаюсь записать в оба массива, это не удается. Но у кого-нибудь есть идея, почему это различное поведение при распределении при использовании двух массивов с одним потоком вместо двух массивов с двумя потоками?

РЕДАКТИРОВАТЬ: еще один тест, который я нахожу интересным для тех, кто, как и я, будет изучать использование кучи: 5. Выполнение ядра два раза с одним массивом размером 900 000 * sizeof(double), выделенным одним потоком 0, работает, если есть удаление. Если удаление не указано, во второй раз произойдет сбой, но будет выполнен первый вызов.

РЕДАКТИРОВАТЬ 2: как выделить переменную для всего устройства один раз, но доступную для записи всеми потоками (не с хоста, используя динамическое распределение в коде устройства)?