HAMi源码解析——HAMi-Core-3

原创

DifficultWork

修改于 2025-07-09 14:53:45

4710

4 GPU Core 限制

GPU 的利用率取决于提交到 GPU 的计算任务，当遇到数据并行处理时，CUDA 会将程序编译成 GPU 能处理的程序并交给 GPU 处理。这种程序在 CUDA 中被称为 Kernel（注意与通常情况下和我们理解的不一样，这里可以把它当作一种程序的别称）。CUDA 将Kernel 交予 GPU 处理的过程为 Launch Kernel。

HAMi 是通过限制 Kernel 的提交来实现 Core limit，具体算法类似令牌桶，初始化时会生成一批 Token，每次提交 Kernel 都会消耗 Token，当提交 Kernel 发现没有 Token 时就会直接 sleep，当时间片结束后又会分配一批 Token。

下面从 CUDA 的提交开始分析。

4.1 Kernel 提交

cuLaunchKernel 是 CUDA 的一个 API，用于提交处理一个 Kernel：

// src/cuda/memory.c
CUresult cuLaunchKernel ( 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, void** extra ){
    ENSURE_RUNNING();
    pre_launch_kernel();
    if (pidfound==1){ 
        // rate_limiter 用来实现 core limit
        rate_limiter(gridDimX * gridDimY * gridDimZ,
                   blockDimX * blockDimY * blockDimZ);
    }
    CUresult res = CUDA_OVERRIDE_CALL(cuda_library_entry,cuLaunchKernel,f,gridDimX,gridDimY,gridDimZ,blockDimX,blockDimY,blockDimZ,sharedMemBytes,hStream,kernelParams,extra);
    return res;
}

4.2 Core 限制实现

rate_limiter 会比较当前的使用情况和从环境变量中获取到的 limit 限值：

// src/multiprocess/multiprocess_utilization_watcher.c
void rate_limiter(int grids, int blocks) {
  long before_cuda_cores = 0;
  long after_cuda_cores = 0;
  long kernel_size = grids;

  while (get_recent_kernel()<0) {
    sleep(1);
  }
  set_recent_kernel(2);
  // SM 没有限制
  if ((get_current_device_sm_limit(0)>=100) || (get_current_device_sm_limit(0)==0))
    	return;
  // 限制关闭
  if (get_utilization_switch()==0)
      return;
  LOG_DEBUG("grid: %d, blocks: %d", grids, blocks);
  LOG_DEBUG("launch kernel %ld, curr core: %ld", kernel_size, g_cur_cuda_cores);
  //if (g_vcuda_config.enable) {
    do {
CHECK:
      // 每次提交成功都会消耗 g_cur_cuda_cores，这里的 g_cur_cuda_cores 可以认为是 Token
      before_cuda_cores = g_cur_cuda_cores;
      LOG_DEBUG("current core: %ld", g_cur_cuda_cores);
      // 当 g_cur_cuda_cores 小于0时表示没有 Token 了，就会触发 block 强制 sleep 等待时间片过后释放 Token 再试
      if (before_cuda_cores < 0) {
        nanosleep(&g_cycle, NULL); // 大概10ms
        goto CHECK;
      }
      after_cuda_cores = before_cuda_cores - kernel_size;
      // 最后则通过 CAS 把 after_cuda_cores 的值赋给 g_cur_cuda_cores
      // 如果 g_cur_cuda_cores 和 before_cuda_cores 值一样说明就没有其他地方更新 g_cur_cuda_cores，就放心把 after_cuda_cores 赋值给 g_cur_cuda_cores，CAS 会返回 true
      // 否则需要重新申请
    } while (!CAS(&g_cur_cuda_cores, before_cuda_cores, after_cuda_cores));
  //}
}

4.2.1 SM 概念

SM（Streaming Multiprocessor，流多处理器）是 NVIDIA GPU 架构的核心计算单元。

一个现代NVIDIA GPU包含多个SM（从几个到几十个不等），每个SM包含：

多个CUDA核心（执行整数和浮点运算）
共享内存
寄存器
纹理单元
调度器
其他专用计算单元（如张量核心） SM以"线程束"(warp)为单位执行指令，一个线程束通常包含32个线程，同时执行相同的指令。多个线程束并发执行在一个SM上。

对于 SM 限制，是 do_init_device_sm_limits 从环境变量中读取的：

// src/multiprocess/multiprocess_memory_limit.c
void do_init_device_sm_limits(uint64_t *arr, int len) {
    // 从 CUDA_DEVICE_SM_LIMIT 环境变量读取
    size_t fallback_limit = get_limit_from_env(CUDA_DEVICE_SM_LIMIT);
    if (fallback_limit == 0) fallback_limit = 100;
    int i;
    for (i = 0; i < len; ++i) {
        char env_name[CUDA_DEVICE_SM_LIMIT_KEY_LENGTH] = CUDA_DEVICE_SM_LIMIT;
        char index_name[8];
        snprintf(index_name, 8, "_%d", i);
        strcat(env_name, index_name);
        size_t cur_limit = get_limit_from_env(env_name);
        if (cur_limit > 0) {
            arr[i] = cur_limit;
        } else if (fallback_limit > 0) {
            arr[i] = fallback_limit;
        } else {
            arr[i] = 0;
        }
    }
}

4.3 Token 恢复

相当于有一个“守护线程”不停的监控 GPU 的使用情况——utilization_watcher：

// src/multiprocess/multiprocess_utilization_watcher.c
void* utilization_watcher() {
    nvmlInit();
    int userutil[CUDA_DEVICE_MAX_COUNT];
    int sysprocnum;
    long share = 0;
    // 初始化并获取限制
    int upper_limit = get_current_device_sm_limit(0);
    ensure_initialized();
    LOG_DEBUG("upper_limit=%d\n",upper_limit);
    while (1){
        nanosleep(&g_wait, NULL);
        if (pidfound==0) {
          update_host_pid();
          if (pidfound==0)
            continue;
        }
        init_gpu_device_utilization();
        // 获取当前 GPU 使用率
        get_used_gpu_utilization(userutil,&sysprocnum);
        //if (sysprocnum == 1 &&
        //    userutil < upper_limit / 10) {
        //    g_cur_cuda_cores =
        //        delta(upper_limit, userutil, share);
        //    continue;
        //}
        if ((share==g_total_cuda_cores) && (g_cur_cuda_cores<0)) {
          g_total_cuda_cores *= 2;
          share = g_total_cuda_cores;
        }
        if ((userutil[0]<=100) && (userutil[0]>=0)){
          // 根据限制值和当前使用情况计算本轮可以增加的 Token
          share = delta(upper_limit, userutil[0], share);
          // 调用 change_token 增加 Token
          change_token(share);
        }
        LOG_INFO("userutil1=%d currentcores=%ld total=%ld limit=%d share=%ld\n",userutil[0],g_cur_cuda_cores,g_total_cuda_cores,upper_limit,share);
    }
}

change_token 给 g_cur_cuda_cores 赋值：

// src/multiprocess/multiprocess_utilization_watcher.c
static void change_token(long delta) {
  int cuda_cores_before = 0, cuda_cores_after = 0;

  LOG_DEBUG("delta: %ld, curr: %ld", delta, g_cur_cuda_cores);
  do {
    cuda_cores_before = g_cur_cuda_cores;
    cuda_cores_after = cuda_cores_before + delta;

    if (cuda_cores_after > g_total_cuda_cores) {
      cuda_cores_after = g_total_cuda_cores;
    }
  } while (!CAS(&g_cur_cuda_cores, cuda_cores_before, cuda_cores_after));
}

原创声明：本文系作者授权腾讯云开发者社区发表，未经许可，不得转载。

如有侵权，请联系 cloudcommunity@tencent.com 删除。

kubernetes