如何在CUDA中分配可执行内存并执行SASS/CUBIN代码？

核心限制说明 #

RTX 4070基于Ampere架构（sm_89），CUDA默认cudaMalloc分配的内存不可执行，且cuMemSetAccess无法直接设置可执行权限——这是因为NVIDIA设备内存的权限模型与CPU不同，可执行内存需从特定内存池分配，而非通用显存。

可行方案：使用底层CUDA API分配并配置可执行内存（需sudo权限） #

1. 分配可执行的托管内存 #

普通cudaMalloc无法获取可执行权限，需用cuMemAllocManaged分配托管内存，再通过属性配置标记为可执行：

CUdeviceptr exec_mem;
size_t sass_size = 256; // 对应你的sm_89 SASS代码实际大小
cuMemAllocManaged(&exec_mem, sass_size, CU_MEM_ATTACH_GLOBAL);

2. 设置可执行权限 #

使用cuMemSetAttribute而非cuMemSetAccess配置可执行属性，Ampere架构支持通过CU_MEM_ATTRIBUTE_EXECUTABLE标记：

unsigned int exec_flag = 1;
cuMemSetAttribute(exec_mem, CU_MEM_ATTRIBUTE_EXECUTABLE, &exec_flag, sizeof(exec_flag));

注意：此操作必须以sudo权限运行，否则会返回权限错误。

写入并执行SASS代码 #

1. 拷贝sm_89兼容的SASS指令 #

先通过cuobjdump -sass your_kernel.cubin提取对应sm_89的SASS指令，再用cuMemcpyHtoD写入可执行内存：

uint8_t sm89_sass[] = { /* 提取的SASS字节码 */ };
cuMemcpyHtoD(exec_mem, sm89_sass, sizeof(sm89_sass));

2. 调用可执行内存中的SASS代码 #

将设备指针转换为全局内核函数指针，直接发起调用：

typedef void(*custom_kernel)(int*);
custom_kernel run_sass = (custom_kernel)exec_mem;

int* d_result;
cudaMalloc(&d_result, sizeof(int));
run_sass<<<1, 1>>>(d_result);
cudaDeviceSynchronize();

关键注意事项 #

• SASS兼容性：必须确保SASS代码完全匹配sm_89架构，包括寄存器分配、指令格式，否则会触发设备异常。
• 权限验证：分配后可通过cuMemGetAttribute检查CU_MEM_ATTRIBUTE_EXECUTABLE是否为1，确认权限设置成功。
• 内存生命周期：可执行内存需在使用完成后用cuMemFree释放，避免显存泄漏。

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