Skip to content

Commit

Permalink
add support for cuModuleLoadData
Browse files Browse the repository at this point in the history 
Signed-off-by: Niklas Eiling <[email protected]>
  • Loading branch information
@n-eiling
n-eiling committed Jul 13, 2023
1 parent ce21d8a commit 481dec9
Show file tree
Hide file tree
Showing 5 changed files with 107 additions and 9 deletions.
46 changes: 45 additions & 1 deletion cpu/cpu-client-driver.c
View file
Original file line number Diff line number Diff line change
Expand Up @@ -6,6 +6,7 @@
#include <cudaEGL.h>
#include <vdpau/vdpau.h>
#include <cudaVDPAU.h>
#include <elf.h>

#include <driver_types.h>
#include <string.h>
Expand Down Expand Up @@ -438,7 +439,50 @@ CUresult cuModuleLoad(CUmodule* module, const char* fname)
}
return result.err;
}
DEF_FN(CUresult, cuModuleLoadData, CUmodule*, module, const void*, image)


CUresult cuModuleLoadData(CUmodule* module, const void* image)
{
enum clnt_stat retval;
ptr_result result;
mem_data mem;

if (image == NULL) {
LOGE(LOG_ERROR, "image is NULL!");
return CUDA_ERROR_INVALID_IMAGE;
}
Elf64_Ehdr *ehdr = (Elf64_Ehdr*)image;

if (ehdr->e_ident[EI_MAG0] != ELFMAG0 ||
ehdr->e_ident[EI_MAG1] != ELFMAG1 ||
ehdr->e_ident[EI_MAG2] != ELFMAG2 ||
ehdr->e_ident[EI_MAG3] != ELFMAG3) {
LOGE(LOG_ERROR, "image is not an ELF!");
return CUDA_ERROR_INVALID_IMAGE;
}

mem.mem_data_len = ehdr->e_shoff + ehdr->e_shnum * ehdr->e_shentsize;
mem.mem_data_val = (uint8_t*)image;

LOGE(LOG_DEBUG, "image_size = %#0zx", mem.mem_data_len);

if (elf2_parameter_info(&kernel_infos, mem.mem_data_val, mem.mem_data_len) != 0) {
LOGE(LOG_ERROR, "could not get kernel infos from memory");
return CUDA_ERROR_INVALID_IMAGE;
}

retval = rpc_cumoduleloaddata_1(mem, &result, clnt);
printf("[rpc] %s(%p) = %d, result %p\n", __FUNCTION__, image, result.err, (void*)result.ptr_result_u.ptr);
if (retval != RPC_SUCCESS) {
fprintf(stderr, "[rpc] %s failed.", __FUNCTION__);
return CUDA_ERROR_UNKNOWN;
}
if (module != NULL) {
*module = (CUmodule)result.ptr_result_u.ptr;
}
return result.err;
}

DEF_FN(CUresult, cuModuleLoadDataEx, CUmodule*, module, const void*, image, unsigned int, numOptions, CUjit_option*, options, void**, optionValues)
DEF_FN(CUresult, cuModuleLoadFatBinary, CUmodule*, module, const void*, fatCubin)
CUresult cuModuleUnload(CUmodule hmod)
Expand Down
2 changes: 0 additions & 2 deletions cpu/cpu-elf2.c
View file
Original file line number Diff line number Diff line change
Expand Up @@ -895,8 +895,6 @@ int elf2_parameter_info(list *kernel_infos, void* memory, size_t memsize)
return -1;
}

hexdump(memory, 0x10);

#define ELF_DUMP_TO_FILE 1

#ifdef ELF_DUMP_TO_FILE
Expand Down
28 changes: 25 additions & 3 deletions cpu/cpu-server-driver.c
View file
Original file line number Diff line number Diff line change
Expand Up @@ -299,6 +299,26 @@ bool_t rpc_cumodulegetfunction_1_svc(uint64_t module, char *name, ptr_result *re
return 1;
}

bool_t rpc_cumoduleloaddata_1_svc(mem_data mem, ptr_result *result,
struct svc_req *rqstp)
{
RECORD_API(mem_data);
RECORD_SINGLE_ARG(mem);
LOG(LOG_DEBUG, "%s(%p, %#0zx)", __FUNCTION__, mem.mem_data_val, mem.mem_data_len);
GSCHED_RETAIN;
result->err = cuModuleLoadData((CUmodule*)&result->ptr_result_u.ptr, mem.mem_data_val);
GSCHED_RELEASE;
if (resource_mg_create(&rm_modules, (void*)result->ptr_result_u.ptr) != 0) {
LOGE(LOG_ERROR, "error in resource manager");
}
if (result->err != 0) {
char *err_str = NULL;
cuGetErrorName(result->err, &err_str);
LOGE(LOG_DEBUG, "cuModuleLoadData result: %s", err_str);
}
RECORD_RESULT(ptr_result_u, *result);
return 1;
}
bool_t rpc_cumoduleload_1_svc(char* path, ptr_result *result,
struct svc_req *rqstp)
{
Expand All @@ -311,9 +331,11 @@ bool_t rpc_cumoduleload_1_svc(char* path, ptr_result *result,
if (resource_mg_create(&rm_modules, (void*)result->ptr_result_u.ptr) != 0) {
LOGE(LOG_ERROR, "error in resource manager");
}
char *err_str = NULL;
cuGetErrorName(result->err, &err_str);
LOGE(LOG_DEBUG, "cuModuleLoad result: %s", err_str);
if (result->err != 0) {
char *err_str = NULL;
cuGetErrorName(result->err, &err_str);
LOGE(LOG_DEBUG, "cuModuleLoad result: %s", err_str);
}
RECORD_RESULT(ptr_result_u, *result);
return 1;
}
Expand Down
5 changes: 3 additions & 2 deletions cpu/cpu_rpc_prot.x
View file
Original file line number Diff line number Diff line change
Expand Up @@ -424,6 +424,7 @@ program RPC_CD_PROG {
mem_result rpc_cuDeviceGetProperties(int) = 1023;
dint_result rpc_cuDeviceComputeCapability(int) = 1024;
int_result rpc_cuDeviceGetP2PAttribute(int, ptr, ptr) = 1025;
ptr_result rpc_cuModuleLoadData(mem_data mem) = 1026;

/* HIDDEN DRIVER API */
/* ptr_result rpc_hidden_get_device_ctx(int) = 1101;
Expand Down Expand Up @@ -562,12 +563,12 @@ program RPC_CD_PROG {
int rpc_cudnnBackendSetAttribute(ptr descriptor,
int attributeName,
int attributeType,
int64_t elementCount,
hyper elementCount,
mem_data arrayOfElements) = 5313;
mem_result rpc_cudnnBackendGetAttribute(ptr descriptor,
int attributeName,
int attributeType,
int64_t requestedElementCount) = 5314;
hyper requestedElementCount) = 5314;
int rpc_cudnnBackendExecute(ptr handle, ptr executionPlan, ptr variantPack) = 5315;
} = 1;
} = 99;
35 changes: 34 additions & 1 deletion tests/cpu/cubin/main.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -3,6 +3,9 @@
#include <cuda_runtime.h>
#include <cuda.h>
#include <unistd.h>
#include <sys/mman.h>
#include <sys/stat.h>
#include <fcntl.h>


#define printCudaErrors(err) __printCudaErrors (err, __FILE__, __LINE__)
Expand Down Expand Up @@ -76,6 +79,32 @@ int getModuleFromCubin(CUmodule *module, const char *cubin)
return 0;
}

int getModuleFromCubinInMemory(CUmodule *module, const char *cubin)
{
int fd = open(cubin, O_RDONLY);
if (fd < 0) {
printf("error\n");
return 1;
}
struct stat st;
if (fstat(fd, &st) < 0) {
printf("error\n");
return 1;
}
printf("size: %#0zx\n", (int)st.st_size);
void *buf = mmap(NULL, st.st_size, PROT_READ, MAP_PRIVATE, fd, 0);
if (buf == MAP_FAILED) {
printf("error\n");
return 1;
}
CUresult err;
if ((err = cuModuleLoadData(module, buf)) != CUDA_SUCCESS) {
printCudaErrors(err);
return 1;
}
return 0;
}

int getModuleFromShared(CUmodule **module, const char *cubin)
{
return 0;
Expand All @@ -97,10 +126,14 @@ int main(int argc, char** argv)
CUmodule module;
CUfunction func;
printf("testing cubin...\n");
if (getModuleFromCubin(&module, "kernel.cubin") != 0) {
if (getModuleFromCubinInMemory(&module, "kernel.cubin") != 0) {
printf("error\n");
return 1;
}
// if (getModuleFromCubin(&module, "kernel.cubin") != 0) {
// printf("error\n");
// return 1;
// }
// if ((err = getModuleFromShared(&module, "kernel.so")) != 0) {
// printf("error\n");
// return 1;
Expand Down

0 comments on commit 481dec9

Please sign in to comment.