Merge branch 'main' into fix/container-memory

Author: loiht2

.github/workflows/style.yaml

@@ -12,4 +12,5 @@ jobs:
     - uses: actions/checkout@master
     - uses: reviewdog/action-cpplint@master
       with:
-        github_token: ${{ secrets.github_token }}
+        github_token: ${{ secrets.github_token }}
+        args: --linelength=120

CMakeLists.txt

@@ -70,7 +70,7 @@ if(GIT_FOUND)
     #string(REPLACE "." "_" GIT_BRANCH ${GIT_BRANCH})
     string(REPLACE "-" "_" GIT_BRANCH ${GIT_BRANCH})
     string(REPLACE "/" "_" GIT_BRANCH ${GIT_BRANCH})
-    message(STATUS "Git formated branch is ${GIT_BRANCH}")
+    message(STATUS "Git formatted branch is ${GIT_BRANCH}")
 endif()
 
 # Generate the static config header file

README.md

@@ -4,14 +4,14 @@ English | [中文](README_CN.md)
 
 ## Introduction
 
-HAMi-core is the in-container gpu resource controller, it has beed adopted by [HAMi](https://github.com/Project-HAMi/HAMi), [volcano](https://github.com/volcano-sh/devices)
+HAMi-core is the in-container gpu resource controller, it has been adopted by [HAMi](https://github.com/Project-HAMi/HAMi), [volcano](https://github.com/volcano-sh/devices)
 
 <img src="./docs/images/hami-arch.png" width = "600" /> 
 
 ## Features
 
 HAMi-core has the following features:
-1. Virtualize device meory
+1. Virtualize device memory
 2. Limit device utilization by self-implemented time shard
 3. Real-time device utilization monitor 
 

src/cuda/hook.c

@@ -10,7 +10,7 @@ extern fp_dlsym real_dlsym;
 cuda_entry_t cuda_library_entry[] = {
     /* Init Part    */ 
     {.name = "cuInit"},
-    /* Deivce Part */
+    /* Device Part */
     {.name = "cuDeviceGetAttribute"},
     {.name = "cuDeviceGet"},
     {.name = "cuDeviceGetCount"},
@@ -111,6 +111,7 @@ cuda_entry_t cuda_library_entry[] = {
     {.name = "cuFuncGetAttribute"},
     {.name = "cuFuncSetAttribute"},
     {.name = "cuLaunchKernel"},
+    {.name = "cuLaunchKernelEx"},
     {.name = "cuLaunchCooperativeKernel"},
     /* cuEvent Part */
     {.name = "cuEventCreate"},
@@ -134,7 +135,7 @@ cuda_entry_t cuda_library_entry[] = {
     {.name = "cuMemCreate"},
     {.name = "cuMemRelease"},
     {.name = "cuMemMap"},
-    {.name = "cucuMemImportFromShareableHandle"},
+    {.name = "cuMemImportFromShareableHandle"},
     {.name = "cuMemAllocAsync"},
     {.name = "cuMemFreeAsync"},
     /* cuda11.7 new api memory part */
@@ -330,7 +331,7 @@ void *find_symbols_in_table_by_cudaversion(const char *symbol,int  cudaVersion)
   const char *real_symbol;
   real_symbol = get_real_func_name(symbol,cudaVersion);
   if (real_symbol == NULL) {
-    // if not find in mulit func version def, use origin logic
+    // if not find in multi func version def, use origin logic
     pfn = find_symbols_in_table(symbol);
   } else {
     pfn = find_real_symbols_in_table(real_symbol);

src/cuda/memory.c

@@ -68,7 +68,7 @@ uint64_t compute_3d_array_alloc_bytes(const CUDA_ARRAY3D_DESCRIPTOR* desc) {
     }
     bytes *= cuarray_format_bytes[desc->Format];
 
-    // TODO: take acount of alignment and etc
+    // TODO: take account of alignment and etc
     // bytes ++ ???
     return bytes;
 }
@@ -87,7 +87,7 @@ uint64_t compute_array_alloc_bytes(const CUDA_ARRAY_DESCRIPTOR* desc) {
     }
     bytes *= cuarray_format_bytes[desc->Format];
 
-    // TODO: take acount of alignment and etc
+    // TODO: take account of alignment and etc
     // bytes ++ ???
     return bytes;
 }
@@ -118,7 +118,7 @@ CUresult cuArrayCreate_v2(CUarray* arr, const CUDA_ARRAY_DESCRIPTOR* desc) {
 
 CUresult cuArrayDestroy(CUarray arr) {
     CUDA_ARRAY3D_DESCRIPTOR desc;
-    LOG_DEBUG("cuArrayDestory");
+    LOG_DEBUG("cuArrayDestroy");
     CHECK_DRV_API(cuArray3DGetDescriptor(&desc, arr));
     /*uint64_t bytes*/
     compute_3d_array_alloc_bytes(&desc);
@@ -281,14 +281,14 @@ CUresult cuMemcpy(CUdeviceptr dst, CUdeviceptr src, size_t ByteCount ){
 }
 
 CUresult cuPointerGetAttribute ( void* data, CUpointer_attribute attribute, CUdeviceptr ptr ){
-    LOG_DEBUG("cuPointGetAttribue data=%p attribute=%d ptr=%llx",data,(int)attribute,ptr);
+    LOG_DEBUG("cuPointGetAttribute data=%p attribute=%d ptr=%llx", data, (int)attribute,ptr);
     ENSURE_RUNNING();
     CUresult res = CUDA_OVERRIDE_CALL(cuda_library_entry,cuPointerGetAttribute,data,attribute,ptr);
     return res;
 }
 
 CUresult cuPointerGetAttributes ( unsigned int  numAttributes, CUpointer_attribute* attributes, void** data, CUdeviceptr ptr ) {
-    LOG_DEBUG("cuPointGetAttribue data=%p ptr=%llx",data,ptr);
+    LOG_DEBUG("cuPointGetAttribute data=%p ptr=%llx", data, ptr);
     ENSURE_RUNNING();
     CUresult res = CUDA_OVERRIDE_CALL(cuda_library_entry,cuPointerGetAttributes,numAttributes,attributes,data,ptr);
     int cur=0;
@@ -307,7 +307,7 @@ CUresult cuPointerGetAttributes ( unsigned int  numAttributes, CUpointer_attribu
 }
 
 CUresult cuPointerSetAttribute ( const void* value, CUpointer_attribute attribute, CUdeviceptr ptr ){
-    LOG_DEBUG("cuPointSetAttribue value=%p attribute=%d ptr=%llx",value,(int)attribute,ptr);
+    LOG_DEBUG("cuPointSetAttribute value=%p attribute=%d ptr=%llx", value, (int)attribute, ptr);
     ENSURE_RUNNING();
     CUresult res = CUDA_OVERRIDE_CALL(cuda_library_entry,cuPointerSetAttribute,value,attribute,ptr);
     return res;
@@ -542,7 +542,7 @@ CUresult cuMipmappedArrayCreate(CUmipmappedArray* pHandle,
 
 CUresult cuMipmappedArrayDestroy(CUmipmappedArray hMipmappedArray) {
     // TODO: compute bytesize
-    LOG_DEBUG("cuMipmappedArrayDestory\n");
+    LOG_DEBUG("cuMipmappedArrayDestroy\n");
     CUresult res = CUDA_OVERRIDE_CALL(cuda_library_entry,cuMipmappedArrayDestroy, hMipmappedArray);
     return res;
 }
@@ -558,6 +558,16 @@ CUresult cuLaunchKernel ( CUfunction f, unsigned int  gridDimX, unsigned int  gr
     return res;
 }
 
+CUresult cuLaunchKernelEx(const CUlaunchConfig *config, CUfunction f, void **kernelParams, void **extra) {
+    ENSURE_RUNNING();
+    pre_launch_kernel();
+    if (pidfound==1){
+        rate_limiter(config->gridDimX * config->gridDimY * config->gridDimZ,
+                   config->blockDimX * config->blockDimY * config->blockDimZ);
+    }
+    CUresult res = CUDA_OVERRIDE_CALL(cuda_library_entry,cuLaunchKernelEx,config,f,kernelParams,extra);
+    return res;
+}
 
 CUresult cuLaunchCooperativeKernel ( CUfunction f, unsigned int  gridDimX, unsigned int  gridDimY, unsigned int  gridDimZ, unsigned int  blockDimX, unsigned int  blockDimY, unsigned int  blockDimZ, unsigned int  sharedMemBytes, CUstream hStream, void** kernelParams ){
     ENSURE_RUNNING();

src/cuda/stream.c

@@ -7,7 +7,7 @@ CUresult cuStreamCreate(CUstream *phstream, unsigned int flags){
 }
 
 CUresult cuStreamDestroy_v2 ( CUstream hStream ){
-    LOG_DEBUG("cuStreamDestory_v2 %p",hStream);
+    LOG_DEBUG("cuStreamDestroy_v2 %p",hStream);
     return CUDA_OVERRIDE_CALL(cuda_library_entry,cuStreamDestroy_v2,hStream);
 }
 

src/include/libcuda_hook.h

@@ -34,13 +34,16 @@ typedef CUresult (*cuda_sym_t)();
   ({    \
     LOG_DEBUG("Hijacking %s", #sym);                                           \
     cuda_sym_t _entry = (cuda_sym_t)CUDA_FIND_ENTRY(table, sym);               \
+    if (_entry == NULL) {                                                      \
+      LOG_ERROR("Hijack failed: %s is NULL", #sym);                            \
+    }                                                                          \
     _entry(__VA_ARGS__);                                                       \
   })
 
 typedef enum {
     /* cuInit Part */
     CUDA_OVERRIDE_ENUM(cuInit),
-    /* cuDeivce Part */
+    /* cuDevice Part */
     CUDA_OVERRIDE_ENUM(cuDeviceGetAttribute),
     CUDA_OVERRIDE_ENUM(cuDeviceGet),
     CUDA_OVERRIDE_ENUM(cuDeviceGetCount),
@@ -142,6 +145,7 @@ typedef enum {
     CUDA_OVERRIDE_ENUM(cuFuncGetAttribute),
     CUDA_OVERRIDE_ENUM(cuFuncSetAttribute),
     CUDA_OVERRIDE_ENUM(cuLaunchKernel),
+    CUDA_OVERRIDE_ENUM(cuLaunchKernelEx),
     CUDA_OVERRIDE_ENUM(cuLaunchCooperativeKernel),
     /* cuEvent Part */
     CUDA_OVERRIDE_ENUM(cuEventCreate),

src/include/nvml-subset.h

@@ -655,7 +655,7 @@ typedef enum nvmlGpuP2PCapsIndex_enum {
 /* P2P Capability Index Status*/
 typedef enum nvmlGpuP2PStatus_enum {
   NVML_P2P_STATUS_OK = 0,
-  NVML_P2P_STATUS_CHIPSET_NOT_SUPPORED,
+  NVML_P2P_STATUS_CHIPSET_NOT_SUPPORTED,
   NVML_P2P_STATUS_GPU_NOT_SUPPORTED,
   NVML_P2P_STATUS_IOH_TOPOLOGY_NOT_SUPPORTED,
   NVML_P2P_STATUS_DISABLED_BY_REGKEY,
@@ -808,7 +808,7 @@ typedef struct nvmlVgpuPgpuMetadata_st {
   char hostDriverVersion
   [NVML_SYSTEM_DRIVER_VERSION_BUFFER_SIZE];  //!< Host driver version
   unsigned int
-      pgpuVirtualizationCaps;   //!< Pgpu virtualizaion capabilities bitfileld
+      pgpuVirtualizationCaps;   //!< Pgpu virtualization capabilities bitfield
   unsigned int reserved[7];     //!< Reserved for internal use
   unsigned int opaqueDataSize;  //!< Size of opaque data field in bytes
   char opaqueData[4];           //!< Opaque data
@@ -966,7 +966,7 @@ typedef enum nvmlVgpuVmCompatibility_enum {
   NVML_VGPU_VM_COMPATIBILITY_HIBERNATE =
   0x2,  //!< vGPU is runnable from a hibernated state (ACPI S4)
   NVML_VGPU_VM_COMPATIBILITY_SLEEP =
-  0x4,  //!< vGPU is runnable from a sleeped state (ACPI S3)
+  0x4,  //!< vGPU is runnable from a slept state (ACPI S3)
   NVML_VGPU_VM_COMPATIBILITY_LIVE =
   0x8,  //!< vGPU is runnable from a live/paused (ACPI S0)
 } nvmlVgpuVmCompatibility_t;
@@ -1093,7 +1093,7 @@ typedef enum nvmlVgpuVmIdType {
  * Represents frame buffer capture session type
  */
 typedef enum nvmlFBCSessionType_enum {
-  NVML_FBC_SESSION_TYPE_UNKNOWN = 0,  //!< Unknwon
+  NVML_FBC_SESSION_TYPE_UNKNOWN = 0,  //!< Unknown
   NVML_FBC_SESSION_TYPE_TOSYS,        //!< ToSys
   NVML_FBC_SESSION_TYPE_CUDA,         //!< Cuda
   NVML_FBC_SESSION_TYPE_VID,          //!< Vid

src/libvgpu.c

@@ -74,7 +74,23 @@ FUNC_ATTR_VISIBLE void* dlsym(void* handle, const char* symbol) {
     LOG_DEBUG("into dlsym %s",symbol);
     pthread_once(&dlsym_init_flag,init_dlsym);
     if (real_dlsym == NULL) {
-        real_dlsym = dlvsym(RTLD_NEXT,"dlsym","GLIBC_2.2.5");
+        const char* glibc_versions[] = {
+                "GLIBC_2.2.5",  // for amd64
+                "GLIBC_2.17",   // for arm64
+                "GLIBC_2.3",
+                "GLIBC_2.4",
+                "GLIBC_2.10",
+                "GLIBC_2.18",
+                "GLIBC_2.22",
+                NULL
+        };
+        for (int i = 0; glibc_versions[i] != NULL; i++) {
+            real_dlsym = dlvsym(RTLD_NEXT, "dlsym", glibc_versions[i]);
+            if (real_dlsym != NULL) {
+                LOG_DEBUG("found dlsym with version: %s", glibc_versions[i]);
+                break;
+            }
+        }
         char *path_search=getenv("CUDA_REDIRECT");
         if ((path_search!=NULL) && (strlen(path_search)>0)){
             vgpulib = dlopen(path_search,RTLD_LAZY);
@@ -203,6 +219,7 @@ void* __dlsym_hook_section(void* handle, const char* symbol) {
     DLSYM_HOOK_FUNC(cuFuncGetAttribute);
     DLSYM_HOOK_FUNC(cuFuncSetAttribute);
     DLSYM_HOOK_FUNC(cuLaunchKernel);
+    DLSYM_HOOK_FUNC(cuLaunchKernelEx);
     DLSYM_HOOK_FUNC(cuLaunchCooperativeKernel);
     DLSYM_HOOK_FUNC(cuIpcOpenMemHandle_v2);
     DLSYM_HOOK_FUNC(cuIpcGetMemHandle);
@@ -838,7 +855,13 @@ void* __dlsym_hook_section_nvml(void* handle, const char* symbol) {
 void preInit(){
     LOG_MSG("Initializing.....");
     if (real_dlsym == NULL) {
-        real_dlsym = _dl_sym(RTLD_NEXT, "dlsym", dlsym);
+        real_dlsym = dlvsym(RTLD_NEXT,"dlsym","GLIBC_2.2.5");
+        if (real_dlsym == NULL) {
+            LOG_ERROR("real dlsym not found");
+            real_dlsym = _dl_sym(RTLD_NEXT, "dlsym", dlsym);
+            if (real_dlsym == NULL)
+                LOG_ERROR("real dlsym not found");
+        }
     }
     real_realpath = NULL;
     load_cuda_libraries();
@@ -848,7 +871,7 @@ void preInit(){
 
 void postInit(){
     allocator_init();
-
+    map_cuda_visible_devices();
     try_lock_unified_lock();
     nvmlReturn_t res = set_task_pid();
     try_unlock_unified_lock();
@@ -860,7 +883,6 @@ void postInit(){
         pidfound=1;
     }
 
-    map_cuda_visible_devices();
     //add_gpu_device_memory_usage(getpid(),0,context_size,0);
     env_utilization_switch = set_env_utilization_switch();
     init_utilization_watcher();

src/multiprocess/multiprocess_memory_limit.c

@@ -436,7 +436,7 @@ int fix_lock_shrreg() {
     if (current_owner != 0) {
         int flag = 0;
         if (current_owner == region_info.pid) {
-            LOG_INFO("Detect onwer pid = self pid (%d), "
+            LOG_INFO("Detect owner pid = self pid (%d), "
                 "indicates pid loopback or race condition", current_owner);
             flag = 1;
         } else {
@@ -589,7 +589,7 @@ void init_proc_slot_withlock() {
     }
     signal(SIGUSR2,sig_swap_stub);
     signal(SIGUSR1,sig_restore_stub);
-    // If, by any means a pid of itself is found in region->proces, then it is probably caused by crashloop
+    // If, by any means a pid of itself is found in region->process, then it is probably caused by crashloop
     // we need to reset it.
     int i,found=0;
     for (i=0; i<region->proc_num; i++) {