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

【问题标题】：CUDA managed memory doesn't work with different compute capabilityCUDA 托管内存不适用于不同的计算能力
【发布时间】：2019-10-25 07:28:25
【问题描述】：

当使用不同的计算能力编译时，此代码的工作方式不同：

#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 上运行时出现分段错误？

【问题讨论】：

因为不支持？并非所有计算能力都支持所有功能。事实上，这是了解您的计算能力很重要的唯一原因！
@AnderBiguri 根据 CUDA 文档，从计算能力 6.0 开始支持统一内存。我提到它实际上适用于具有计算能力 6.0 的 GPU。问题是：当我为计算能力 6.0 编译二进制文件时，它在支持计算能力 7.0 的 GPU 上不起作用。
我不确定。在我的工具中，如果运行 CUDA 代码为旧版 cc 编译，我通常会得到“没有可在设备上执行的内核映像”，即使只是以前的版本。你可能没有得到它，因为你没有内核
每个 CUDA 安装都对使用不正确的计算能力保持沉默，除非您 explicitly check for errors。
仔细阅读链接的答案，因为内核启动的错误检查有点古怪。您需要在调用cudaDeviceSynchronize() 之前检查cudaPeekAtLastError() 或cudaGetLastError() 的返回值以检查此类启动错误，否则错误会在您注意到它之前被重置。

标签： cuda

【解决方案1】：

根据 cmets 中的讨论，我遇到了这个问题，因为我必须在构建期间使用 与我的 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

【讨论】：

您不必为运行代码的 GPU 的确切计算能力进行编译。包含可以在程序加载时即时编译的 PTX 代码就足够了（如果您不指定任何内容，这是默认设置）。但是，如果您通过明确指定为真实架构生成二进制文件来禁用它，那么它必须匹配。