Управляемая память CUDA не работает с другими вычислительными возможностями

Question

Этот код не работает одинаково при компиляции с различными вычислительными возможностями:

#include <cuda.h>
#include <stdio.h>

__managed__ int m;

int main() {
    printf("hi 1\n");
    m = -123;
    printf("hi 2\n");
}

Устройство с вычислительными возможностями 6.0:

$ nvcc main.cu -gencode arch=compute_60,code=sm_60 -rdc=true && ./a.out
hi 1
hi 2

Устройство с вычислительными возможностями 7.0:

$ nvcc main.cu -gencode arch=compute_60,code=sm_60 -rdc=true && ./a.out
hi 1
Segmentation fault

Устройство с возможностью вычислений 7.0:

$ nvcc main.cu -gencode arch=compute_70,code=sm_70 -rdc=true && ./a.out
hi 1
hi 2

Почему у меня возникает ошибка сегментации при построении с возможностью вычислений 6.0 и при запуске его на GPU с возможностью вычислений 7.0?

Answer 1

Согласно обсуждению в комментариях, у меня есть эта проблема, потому что я должен использовать точно такую ​​же вычислительную возможность во время сборки, как у моего GPU . Причина, по которой я не получил никаких ошибок, заключается в том, что я должен проверить их вручную ( Что такое канонический способ проверки на ошибки с использованием API времени выполнения CUDA? ).

Если я расширяюфрагмент кода, подобный следующему:

#include <cuda.h>
#include <stdio.h>

__managed__ int m;

__global__ void foo() {
    printf("from foo: %d %d\n", blockIdx.x, threadIdx.x);
}

int main() {
    foo<<<2,2>>>();
    printf("001\n");
    if (cudaPeekAtLastError() != cudaSuccess) abort();
    printf("002\n");
    if (cudaDeviceSynchronize() != cudaSuccess) abort();
    printf("hi 1\n");
    m = -123;
    printf("hi 2\n");
}

Устройство с вычислительными возможностями 7.0:

$ nvcc main.cu -gencode arch=compute_70,code=sm_70 -rdc=true && ./a.out
001
002
from foo: 0 0
from foo: 0 1
from foo: 1 0
from foo: 1 1
hi 1
hi 2

Устройство с вычислительными возможностями 7.0:

$ nvcc main.cu -gencode arch=compute_60,code=sm_60 -rdc=true && ./a.out
001
Aborted