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 构建时出现分段错误，而运行在计算能力 7.0 的 GPU 上构建它？

Answer 1

根据评论中的讨论，我遇到这个问题是因为我必须在构建期间使用 与我的 GPU 完全相同的计算能力。我没有收到任何错误的原因是我应该手动检查它们 (What is the canonical way to check for errors using the CUDA runtime API?)。

如果我像这样扩展这段代码：

#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

CUDA 托管内存不适用于不同的计算能力

CUDA managed memory doesn't work with different compute capability

cuda

gpu-managed-memory