自学内容网 自学内容网
自学内容网 > 正文

CUDA Context学习及实验

🕗 发布于 2024-10-11 15:12 GPU  CUDA  GPGPU  

CUDA Context学习及实验

CUDA上下文（CUDA Context）是一个核心概念，负责在GPU设备上管理和维护CUDA程序执行所需的所有状态和资源。

一、CUDA上下文的原理

1. 什么是CUDA上下文

  • 定义：CUDA上下文是CUDA程序在GPU设备上的执行环境或工作空间，类似于CPU编程中的进程或线程上下文。
  • 包含内容：设备内存分配、内核（kernel）代码、流（stream）、事件（event）等所有GPU执行所需的状态信息和资源。

2. 上下文的创建与绑定

  • 创建：当主机线程首次与GPU设备交互（如调用CUDA API函数）时，CUDA驱动程序会为该线程创建一个CUDA上下文。
  • 绑定：创建的CUDA上下文绑定到调用线程，一个线程只能有一个当前的CUDA上下文，对GPU的所有操作都在该上下文中执行。

3. 上下文的作用范围

  • 线程独立性：CUDA上下文的作用范围仅限于创建它的主机线程，不同线程有各自的CUDA上下文，彼此独立。
  • 资源隔离：GPU资源和状态在上下文之间隔离，避免并发访问时的资源冲突和数据竞争。

二、CUDA上下文的功能

1. 管理GPU资源

  • 内存管理：负责设备内存的分配和释放，包括全局内存、共享内存、常量内存等。
  • 内核加载与执行：管理CUDA内核代码的加载和执行，编译PTX代码，管理内核参数。
  • 流和事件：包含用于并发执行和同步的流（streams）和事件（events）。

2. 维护执行状态

  • 配置状态：包括当前活动的设备、线程块配置、访存配置等。
  • 错误状态：维护CUDA API调用的错误状态，供开发者调试和处理异常。

3. 资源和状态的隔离

  • 独立性：每个上下文有自己独立的资源，不会与其他上下文冲突。
  • 安全性：防止多个线程同时访问同一资源导致的不确定行为。

4. 支持并行和异步执行

  • 流（Streams）：允许在同一设备上并发地执行多个内核或内存传输操作，提高设备利用率。
  • 事件（Events）：用于标记和测量GPU执行的特定点，帮助实现主机和设备之间的同步。

三、CUDA上下文的工作机制

1. 上下文的创建与销毁

  • 自动创建：主机线程首次调用CUDA API与设备交互时，CUDA驱动自动创建上下文。
  • 手动管理：高级用户可以使用CUDA的上下文管理API手动创建、绑定和销毁上下文，获得更细粒度的控制。
  • 销毁：上下文不再需要时，可以显式销毁以释放GPU资源；否则，当主机线程结束时，CUDA驱动会自动清理。

2. 上下文的切换

  • 单线程多上下文：一个主机线程可以创建多个CUDA上下文，但同一时刻只能有一个当前上下文。
  • 切换代价：上下文切换需要保存和恢复大量状态，有一定的性能开销，应尽量避免频繁切换。

3. 多线程与多上下文

  • 多线程环境：每个线程都有自己的CUDA上下文，彼此独立。
  • 共享上下文：可以通过线程同步和上下文共享机制，让多个线程访问同一个上下文，但需要谨慎管理以避免资源冲突。

四、CUDA上下文的内存可见性与隔离

同一进程内的不同上下文

  • 设备内存的可见性：在同一进程中，不同的CUDA上下文分配的设备内存物理上位于同一GPU的全局内存中，可以相互拷贝数据。
  • 数据拷贝：可使用cudaMemcpy、cudaMemcpyAsync、cudaMemcpyPeer等CUDA API，在不同上下文的设备内存之间执行数据拷贝。

不同进程之间的设备内存隔离

  • 默认隔离：不同进程中的CUDA上下文及其设备内存默认是隔离的，彼此不可见，不能直接访问或拷贝数据。
  • MMU与进程关联：GPU的内存管理单元在进程级别上管理设备内存的虚拟地址空间，确保进程间的内存隔离。

进程间通信的特殊机制

  1. CUDA进程间通信（IPC）
    • 共享内存句柄：一个进程可以通过cudaIpcGetMemHandle获取设备内存的句柄，然后通过进程间通信方式（如socket、管道、共享内存等）将句柄传递给另一个进程。
    • 打开句柄访问内存：接收方进程使用cudaIpcOpenMemHandle打开内存句柄，从而访问同一块设备内存。
    • 限制条件：要求两个进程在同一物理GPU上运行，且GPU和驱动程序需要支持CUDA IPC功能。
  2. 统一虚拟地址空间（UVA）
    • 跨进程地址一致性：在支持UVA的系统上，GPU的虚拟地址空间在一定程度上是一致的，但仍需使用CUDA提供的IPC机制进行跨进程内存共享。
    • 简化内存管理：UVA简化了设备和主机以及不同设备间的内存指针管理，但不改变进程间的内存隔离特性。

五、资源的上下文绑定与限制

流、模块和函数的上下文绑定

  • 上下文依赖性：流（stream）、模块（module）和函数（function）都是在特定的CUDA上下文中创建的，操作和生命周期与创建的上下文密切相关。
  • 不可跨上下文使用：不能在一个上下文中使用在另一个上下文中创建的流、模块或函数，否则会导致未定义的行为或运行时错误。
  • 符号解析：模块内的符号（如全局变量、常量）只能在其所属的上下文中解析和访问。

上下文隔离的设计目的

  • 资源管理：上下文提供对GPU资源的逻辑隔离，防止不同任务或库之间的干扰。
  • 稳定性和安全性：隔离不同上下文的资源，提高系统稳定性，防止一个上下文中的错误影响其他上下文。

六、实践中的注意事项

  • 一致性使用：确保在哪个上下文中创建的资源，就在该上下文中使用，不要混用。
  • 上下文切换：如果需要在多个上下文中工作，确保正确地切换到相应的上下文。
  • 减少上下文数量：如有可能，尽量在同一个上下文中完成所有操作，避免不必要的上下文切换。
  • 减少上下文开销：避免频繁创建和销毁上下文，尽可能重用已有的上下文。
  • 流的利用：充分利用流和异步操作，提高GPU利用率和程序的并发性。
  • 显式释放资源：及时释放不再需要的设备内存和资源，防止内存泄漏。
  • 上下文销毁：程序结束前，显式销毁上下文，确保所有资源被正确释放。
  • 错误检查：在CUDA API调用后，检查返回状态，处理可能的错误。
  • 上下文复位：发生不可恢复的错误时，可以重置上下文，使GPU处于干净的状态。

七.CUDA driver API Context Management介绍

1. 创建新的 CUDA 上下文

  • 使用 cuCtxCreate() 函数创建新的 CUDA 上下文，并将其与调用线程关联。
  • 上下文创建时可以指定 flags 参数来控制 CPU 线程在等待 GPU 结果时的调度行为。
  • 如果线程已经有一个当前上下文，新创建的上下文将取代之前的，可以通过 cuCtxPopCurrent() 恢复。
  • 上下文创建后，使用计数为 1，使用完毕后需要调用 cuCtxDestroy() 销毁上下文。

2. 执行关联性和资源限制

  • 上下文可以通过 paramsArray 和 numParams 来限制其可以使用的执行资源，例如限制可用的 SM 数量。
  • 通过指定 CUexecAffinityParam 数组，可以设置执行关联性的参数。

3. 控制 CPU 线程调度行为的标志

创建上下文时，可以使用 flags 参数的最低三位来设置 CPU 线程在等待 GPU 结果时的行为：

  • CU_CTX_SCHED_SPIN：主动轮询等待 GPU 结果，降低延迟，但可能影响其他 CPU 线程性能。
  • CU_CTX_SCHED_YIELD：让出 CPU 线程，等待 GPU 结果时增加延迟，但提高其他 CPU 线程性能。
  • CU_CTX_SCHED_BLOCKING_SYNC：在等待 GPU 完成工作时，阻塞 CPU 线程。
  • CU_CTX_SCHED_AUTO：默认值，根据进程中的活动上下文数量和系统的逻辑处理器数量，自动选择策略。

4. 其他上下文创建标志

  • CU_CTX_MAP_HOST：支持映射固定的主机内存，使其可被 GPU 访问。
  • CU_CTX_LMEM_RESIZE_TO_MAX（已弃用）：指示 CUDA 不要在调整内核的本地内存后减少其大小。
  • CU_CTX_COREDUMP_ENABLE：如果未全局启用 GPU 核心转储，设置此标志可在上下文执行期间发生异常时创建核心转储。
  • CU_CTX_USER_COREDUMP_ENABLE：允许用户触发 GPU 核心转储，需要先全局设置管道名称。

5. 上下文销毁

  • 使用 cuCtxDestroy() 销毁 CUDA 上下文，无论其被多少线程作为当前上下文。
  • 销毁上下文会清理所有与之关联的资源，包括 CUDA 模块、函数、流、事件、数组、纹理对象等。
  • 如果要销毁的上下文是调用线程的当前上下文，它也会从线程的上下文堆栈中弹出。

6. 获取和设置上下文属性

  • cuCtxGetApiVersion()：获取上下文的 API 版本号。
  • cuCtxGetCacheConfig() 和 cuCtxSetCacheConfig()：获取和设置上下文的缓存配置。
  • cuCtxGetCurrent() 和 cuCtxSetCurrent()：获取和设置当前线程的 CUDA 上下文。
  • cuCtxGetDevice()：获取当前上下文的设备编号。
  • cuCtxGetFlags()：获取当前上下文的标志位。
  • cuCtxGetLimit() 和 cuCtxSetLimit()：获取和设置上下文的限制参数，如线程堆栈大小、设备运行时同步深度等。

7. 上下文同步和事件

  • cuCtxSynchronize()：阻塞直到当前上下文完成所有先前请求的任务。
  • cuCtxWaitEvent()：使上下文中的所有未来工作等待指定的事件完成。
  • cuCtxRecordEvent()：在上下文中捕获当前所有活动到事件对象。

8. 注意事项

  • 如果设备的计算模式为 CU_COMPUTEMODE_PROHIBITED，上下文创建将失败，返回 CUDA_ERROR_UNKNOWN。
  • 某些标志和功能在特定的 CUDA 版本或计算能力下可能已弃用或不支持，需要参考相应的文档。
  • 在创建上下文前，确保已正确设置所需的全局属性（如核心转储的管道名称）。

9. 错误处理

  • 上下文创建失败可能返回的错误包括 CUDA_ERROR_INVALID_VALUE（参数无效）和 CUDA_ERROR_NOT_SUPPORTED（不支持的功能）。
  • 调用函数前，确保上下文和资源处于有效状态，避免未定义的行为。

八.复现过程

tee cuda_contexts.cu<<-'EOF'
#include <cuda_runtime.h>
#include <iostream>

__global__ void Kernel_v1(float *data) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    float val=data[idx];
    for(int i=0;i<10240;i++)
    {
        val+=__sinf((float)i);
    }
    data[idx]=val;
}
EOF

/usr/local/cuda/bin/nvcc -std=c++17 -dc -arch=sm_86 -ptx cuda_contexts.cu -o cuda_contexts.ptx
/usr/local/cuda/bin/nvcc -arch=sm_86 cuda_contexts.ptx -cubin -o cuda_contexts.cubin
/usr/local/cuda/bin/nvcc -arch=sm_86 cuda_contexts.cubin -fatbin -o cuda_contexts.fatbin

tee cuda_contexts_main.cpp<<-'EOF'
#include <stdio.h>
#include <string.h>
#include <cuda_runtime.h>
#include <cuda.h>
#include <thread>
#include <chrono>
#include <assert.h>
#include <cstdio>
#include <cuda.h>
#include <iostream>

#define CHECK_CUDA(call)                      \
  do {                              \
    cudaError_t err = call;                  \
    if (err != cudaSuccess) {                 \
      std::cerr << "CUDA error at " << __FILE__ << ":" << __LINE__ << " "<< #call; \
      std::cerr << " code=" << err << " (" << cudaGetErrorString(err) << ")" << std::endl; \
    }                                      \
  } while (0)

#define CHECK_CUDA_DRV_API(call)                      \
  do {                              \
    CUresult err = call;                  \
    if (err != CUDA_SUCCESS) {                 \
        char *error_str=new char[1024];  \
        cuGetErrorString(err,(const char**)&error_str); \
        printf("[%s:%d] %s Error :%s!\n",__FILE__,__LINE__,#call,error_str); \
    }                                      \
  } while (0)

const char* module_file = "cuda_contexts.fatbin";
const char *kernel_name = "_Z9Kernel_v1Pf";

int block_count=10000;
int block_size=1024;
int thread_size=block_count*block_size;
int data_size=sizeof(float)*thread_size;

CUmodule module[128]={0};
CUfunction function[128]={0};
CUcontext cuContext[128]={0};
CUstream hStream[128]={0};
float *ptrs[128]={0};
float *uva_ptrs[128]={0};

int create_resource(int count)
{
    CUresult error;
    cuInit(0);
    CUdevice cuDevice;
    int deviceCount = 0;
    CHECK_CUDA_DRV_API(cuDeviceGetCount(&deviceCount));
    CHECK_CUDA_DRV_API(cuDeviceGet(&cuDevice, 0));
    for(int i=0;i<count;i++)
    {
        CHECK_CUDA_DRV_API(cuCtxCreate(&cuContext[i], 0, cuDevice));
        CHECK_CUDA_DRV_API(cuCtxSetCurrent(cuContext[i]));
        CHECK_CUDA(cudaMalloc((void**)&ptrs[i],data_size));

        CHECK_CUDA(cudaMallocManaged(&uva_ptrs[i], data_size));
        CHECK_CUDA(cudaMemPrefetchAsync(uva_ptrs[i],data_size,0));
        
        CHECK_CUDA_DRV_API(cuStreamCreate(&hStream[i],CU_STREAM_NON_BLOCKING));
        CHECK_CUDA_DRV_API(cuModuleLoad(&module[i], module_file));
        CHECK_CUDA_DRV_API(cuModuleGetFunction(&function[i], module[i], kernel_name));
        void *kernelParams[]= {(void*)&ptrs[i]};
        CHECK_CUDA_DRV_API(cuLaunchKernel(function[i],block_count,1,1,block_size,1,1,0,hStream[i],kernelParams, 0));
        CHECK_CUDA_DRV_API(cuStreamSynchronize(hStream[i]));
    }
    return 0;
}

int invoke(int fun_idx,int stream_idx,int mem_idx)
{
    printf("invoke fun_idx:%d stream_idx:%d mem_idx:%d\n",fun_idx,stream_idx,mem_idx);
    auto start = std::chrono::high_resolution_clock::now();
    void *kernelParams[]= {(void*)&ptrs[mem_idx]};
    CHECK_CUDA_DRV_API(cuLaunchKernel(function[fun_idx],block_count,1,1,block_size,1,1,0,hStream[stream_idx],kernelParams, 0));
    CHECK_CUDA_DRV_API(cuStreamSynchronize(hStream[stream_idx]));
    auto end = std::chrono::high_resolution_clock::now(); 
    std::chrono::duration<double> diff = end - start; 
    printf("E2E:%7.2fms\n",diff.count()*1000);
    return 0;
}

int copy(int first_idx,int second_idx)
{
    printf("copy first_idx:%d second_idx:%d\n",first_idx,second_idx);
    CHECK_CUDA(cudaMemcpy(ptrs[first_idx],ptrs[second_idx],data_size,cudaMemcpyDeviceToDevice));
    return 0;
}

int uva_copy(int first_idx,int second_idx)
{
    printf("uva_copy first_idx:%d second_idx:%d\n",first_idx,second_idx);
    CHECK_CUDA(cudaMemcpy(uva_ptrs[first_idx],uva_ptrs[second_idx],data_size,cudaMemcpyDeviceToDevice));
    return 0;
}

int del_resource(int count)
{
    for(int i=0;i<count;i++)
    {
        if(hStream[i]==0) break;        
        CHECK_CUDA(cudaFree(ptrs[i]));
        CHECK_CUDA(cudaFree(uva_ptrs[i]));
        CHECK_CUDA_DRV_API(cuModuleUnload(module[i]));
        CHECK_CUDA_DRV_API(cuStreamDestroy(hStream[i]));
        CHECK_CUDA_DRV_API(cuCtxDestroy(cuContext[i]));
        hStream[i]=0;
    }
}

int del_resource_test(int count)
{
    for(int i=0;i<count;i++)
    {
        if(hStream[i]==0) break;
        printf("release index:%d\n",i);
        CHECK_CUDA_DRV_API(cuCtxDestroy(cuContext[i]));
        CHECK_CUDA(cudaFree(ptrs[i]));
        CHECK_CUDA(cudaFree(uva_ptrs[i]));
        CHECK_CUDA_DRV_API(cuStreamDestroy(hStream[i]));
        CHECK_CUDA_DRV_API(cuModuleUnload(module[i]));
        
    }
}

int main(int argc,char *argv[])
{
    int mode=atoi(argv[1]);
    int fun_idx=atoi(argv[2]);
    int stream_idx=atoi(argv[3]);
    int mem_idx=atoi(argv[4]);
    create_resource(7);
    if(mode==0)
    {
       invoke(fun_idx,stream_idx,mem_idx);
    }
    if(mode==1)
    {
       copy(fun_idx,stream_idx);
    }
    if(mode==2)
    {
       uva_copy(fun_idx,stream_idx);
    }   
    if(mode==3)
    {
       del_resource_test(7);
    }      
    del_resource(7);
    return 0;
}
EOF
g++ cuda_contexts_main.cpp -g -o cuda_contexts_main \
    -I /usr/local/cuda/include -L /usr/local/cuda/lib64 -lcudart -lcuda -lpthread
./cuda_contexts_main 0 0 0 0 #正常
./cuda_contexts_main 0 1 0 0 #cuLaunchKernel失败
./cuda_contexts_main 0 0 1 0 #cuLaunchKernel失败
./cuda_contexts_main 0 0 0 1 #正常
./cuda_contexts_main 1 0 0 0 #正常
./cuda_contexts_main 1 1 0 0 #正常
./cuda_contexts_main 1 0 1 0 #正常
./cuda_contexts_main 2 0 0 0 #正常
./cuda_contexts_main 2 1 0 0 #正常
./cuda_contexts_main 2 0 1 0 #正常
./cuda_contexts_main 3 0 0 0 #如果先销毁context,再销毁由它创建的stream,mem,model会失败(销毁context时,会销毁它包含的所有资源)

原文地址：https://blog.csdn.net/m0_61864577/article/details/142819709

免责声明：本站文章内容转载自网络资源，如本站内容侵犯了原著者的合法权益，可联系本站删除。更多内容请关注自学内容网（zxcms.com）！

相关文章

自学内容网
Copyright © 2015-2024