NVCC 8.0.61编译Accelerator.cu时出现__managed__属性适用错误

报错原因分析 #

你遇到的error: attribute "managed" does not apply here错误，在CUDA 8.0（NVCC 8.0.61）中通常是因为违反了该版本对统一内存__managed__属性的使用限制，主要有以下几种常见情况：

• __managed__用于类/结构体成员变量：CUDA 8.0并不支持将__managed__属性应用于类或结构体的成员变量，这个特性是在CUDA 9.0及以后的版本才引入的。如果你的第8行代码是在类定义中给成员变量加上了__managed__，这就会触发该错误。
• __managed__用于局部变量（非静态）：CUDA 8.0中，__managed__只能用于全局作用域变量或静态局部变量，普通的函数局部变量不能使用该属性。
• __managed__应用于非POD类型的全局变量：CUDA 8.0对全局作用域的__managed__变量类型有限制，仅支持POD（Plain Old Data）类型（比如基本数据类型、无自定义构造/析构函数的简单结构体等）。如果你的全局__managed__变量是带有复杂逻辑的自定义类型，也会报错。

针对不同场景的解决方案 #

场景1：类成员变量误用__managed__ #

如果是想让类的成员在统一内存中，你可以改为在类中保存指向统一内存的指针，然后在运行时通过cudaMallocManaged动态分配内存：

#include <cuda_runtime_api.h>

class MyAccelerator {
public:
    float* device_buffer; // 用指针替代直接的__managed__成员

    MyAccelerator() {
        // 动态分配统一内存
        cudaMallocManaged(&device_buffer, sizeof(float) * 1024);
    }

    ~MyAccelerator() {
        // 释放内存
        cudaFree(device_buffer);
    }
};

场景2：局部变量误用__managed__ #

将局部变量改为静态局部变量，或者改为全局作用域的__managed__变量，也可以用cudaMallocManaged动态分配：

// 正确用法1：全局作用域的__managed__变量
__managed__ int global_managed_data[256];

void kernel_wrapper() {
    // 正确用法2：静态局部的__managed__变量
    static __managed__ int static_managed_val;

    // 正确用法3：动态分配统一内存
    double* dynamic_managed_ptr;
    cudaMallocManaged(&dynamic_managed_ptr, sizeof(double) * 512);
    
    // 使用内存...
    
    cudaFree(dynamic_managed_ptr);
}

场景3：非POD类型的全局变量 #

如果你的自定义类型需要放在统一内存中，在CUDA 8.0中只能通过动态分配的方式实现：

#include <string>

struct ComplexData {
    int id;
    std::string name; // 非POD类型
};

void init_complex_data() {
    ComplexData* managed_data;
    cudaMallocManaged(&managed_data, sizeof(ComplexData));
    
    // 初始化数据
    managed_data->id = 100;
    managed_data->name = "cuda_managed";
    
    // 使用数据...
    
    cudaFree(managed_data);
}

额外注意事项 #

• 如果你有条件升级到CUDA 9.0或更高版本，很多__managed__的使用限制会被移除，比如支持类成员变量、非POD类型的全局__managed__变量等，会大幅简化代码。
• 编译时可以加上对应GPU架构的参数（比如-arch=sm_35，根据你的GPU计算能力调整），确保统一内存特性被正确启用。

内容的提问来源于stack exchange，提问作者Jake Long