Выделить исполняемую память и выполнить ее в CUDA ⇐ Linux

[Anonymous]

Я хотел бы выделить исполняемую память в CUDA, написать туда код SASS/CUBIN, а затем выполнить этот код. На ЦП для систем Linux это довольно просто и хорошо документировано — просто комбинация mprotect и mmap выполнит работу по распределению памяти, и вы сможете выделить память, которая исполняемый файл.
Я попытался сделать следующее на RTX 4070, показав, что память по умолчанию не является исполняемой (компилировать через nvcc -arch=sm_89 FILE.cu -lcuda):

Код: Выделить всё

#include 
#include 
#include 

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(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< /code>.
Я также пробовал использовать cuMemSetAccess, используя код, приведенный в этом ответе, и изменив строку

Код: Выделить всё

accessDesc.flags = CU_MEM_ACCESS_FLAGS_PROT_READWRITE;

для

Код: Выделить всё

accessDesc.flags = (CUmemAccess_flags) 0x7;

поскольку это соответствует доступу к исполняемой, читаемой и исполняемой памяти в (внутреннем?) заголовке nvport/memory.h модулей открытого ядра NVIDIA Linux. Однако такое изменение приводит к ошибке CUDA.
Мне известен nvJitLink от NVIDIA, но меня не интересуют ответы, связанные с этим здесь.
Итак, как я могу выделить и использовать исполняемую память для карт NVIDIA?
Отвечая на вопрос, вы можете предположить, что я использую новейшую систему Ubuntu с доступом sudo, процессором x86 и РТХ 4070 графический процессор.

Подробнее здесь: https://stackoverflow.com/questions/792 ... it-in-cuda