分配可执行内存并在CUDA中执行

Question

我想在CUDA中分配可执行内存，在那里编写SASS/CUBIN代码，然后执行该代码。 在 Linux 系统的 CPU 上，这非常简单并且有详细记录 - 只需组合

mprotect

和

mmap

即可完成内存分配工作，并且您可以分配可执行的内存。

我尝试在 RTX 4070 上执行以下操作，显示内存默认情况下不可执行（通过

nvcc -arch=sm_89 FILE.cu -lcuda

编译）：

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

typedef void (* funptr)(int *);

__global__ void globalfunc(int * a, void * fun)
{
    funptr ptr = (funptr) fun;
    ptr(a);
}

int main(void)
{
    int h_a[1];
    int * d_a;
    uint64_t h_ins[32] =
    {
        // This is the SASS code for sm_89 with function signature
        // __device__ void myfunc(int * a)
        // {
        //     *a = 1337;
        // }
        0x0000053900037802,
        0x000fe20000000f00,
        0x0000460000047ab9,
        0x000fc80000000a00,
        0x0000000304007985,
        0x0001e4000c101904,
        0x0000000014007950,
        0x001fea0003e00000,
        0xfffffff000007947,
        0x000fc0000383ffff,
        0x0000000000007918,
        0x000fc00000000000,
        0x0000000000007918,
        0x000fc00000000000,
        0x0000000000007918,
        0x000fc00000000000,
        0x0000000000007918,
        0x000fc00000000000,
        0x0000000000007918,
        0x000fc00000000000,
        0x0000000000007918,
        0x000fc00000000000,
        0x0000000000007918,
        0x000fc00000000000,
        0x0000000000007918,
        0x000fc00000000000,
        0x0000000000007918,
        0x000fc00000000000,
        0x0000000000007918,
        0x000fc00000000000,
        0x0000000000007918,
        0x000fc00000000000
    };
    void * d_ins;

    cudaMalloc((void **) &d_a, 1 * sizeof(int));
    cudaMalloc((void **) &d_ins, 32 * sizeof(uint64_t));
    cudaMemcpy(d_ins, h_ins, 32 * sizeof(uint64_t), cudaMemcpyHostToDevice);

    // Executable code seems to require 128 byte alignments, at least on Ada architecture.
    // cudaMalloc allegedly allocate on 256 byte alignments, so we assert that this indeed
    // is the case.
    assert(((uint64_t) d_ins) % 256 == 0);

    // Launch the kernel with one block and 1 thread
    globalfunc<<<1, 1>>>(d_a, d_ins);

    // Copy the result back to the host
    cudaMemcpy(h_a, d_a, sizeof(int), cudaMemcpyDeviceToHost);

    // Print the result
    printf("*h_a = %d\n", *h_a);

    // Free device memory
    cudaFree(d_a);
    cudaFree(d_ins);

    return 0;
}

也就是说，使用实际的

__device__ void myfunc(int * a)

运行代码可以按预期工作，但将 SASS 指令推入内存只会产生

*h_a = 0

。

我还尝试使用

cuMemSetAccess

，通过使用这个答案中提供的代码并更改行

accessDesc.flags = CU_MEM_ACCESS_FLAGS_PROT_READWRITE;

到

accessDesc.flags = (CUmemAccess_flags) 0x7;

因为这似乎对应于 NVIDIA Linux 开放内核模块（内部？）标头中的可执行、可读和可执行内存访问

nvport/memory.h

。 然而，这样的更改会产生 CUDA 错误。

我知道 NVIDIA 的 nvJitLink，但我对此处涉及此的答案不感兴趣。

那么，如何为 NVIDIA 卡分配和使用可执行内存？

在回答问题时，您可能会假设我使用的是最新的 Ubuntu 系统，具有 sudo 访问权限、x86 CPU 和 RTX 4070 GPU。

Answer 1

我想在 CUDA 中分配可执行内存......

不存在用户可分配的“可执行”内存之类的东西。我所看到的所有经验证据以及 NVIDIA 多年来发布的架构白皮书都表明 GPU 具有可编程的 MMU，并且 NVIDIA 选择将 GPU DRAM 逻辑上划分为不同功能的区域（全局内存、常量内存、本地内存） 、代码页）。后者在设计上似乎完全无法从用户代码中访问。

在那里编写SASS/CUBIN代码，然后执行该代码。

我也不明白这是如何运作的。 CUDA 执行模型需要在链接阶段静态分配全局符号、寄存器、本地内存和常量内存，这必须在代码加载到 GPU 并执行之前执行。此链接阶段可以在编译时或运行时完成，但必须完成。这就是您在问题中拒绝的 nvjitlink API 的目的。据我所知，您不可能运行资源需求未知的代码先验。

最后，我认为绕过 NVIDIA 在其驱动程序和运行时中实施的所有保护并在 GPU 上注入和运行任意代码的能力是一个潜在的安全缺陷，并期望 NVIDIA 消除它，如果这样的向量是有记录存在。