nvcc 编译过程

nvcc 编译分为两个阶段
首先将 .cu 编译为面向 虚拟架构 的 .ptx 代码 (stage1)
然后将 .ptx 编译为面向 实际架构 的 二进制代码 (stage2)

- 阅读剩余部分 -

cupti : CUDA profiling tools interface
搭建面向 CUDA 应用的 profiling 和 tracing 工具

CUDA dynamic parallelism : CDP

CUPTI 提供四种 API :

• activity api
• callback api
• event api
• metric api

CUPTI 在第一次调用 CUPTI 函数时懒惰初始化
cuptiSubscribe() : 最先调用，防止多个 CUPTI client 互相干扰

目前理解：CUPTI 分为 用户端 和 服务端， 服务端记录 CUDA 设备和 CPU 上产生的事件 CUpti_Activity，储存在 用户端提供的 Activity Buffer 上
CUPTI 不保证 activity 在 activity buffer 中的顺序
用户端 调用 cuptiActivityFlushPeriod 和 cuptiActivityFlushAll

CUPTI 创建一个 worker thread，以减少对 application thread 的干扰

activate api

activate record : 记录事件，使用基类 CUpti_Activity
activity buffer : 将 activity record 从 CUPTI 转移到 client
使用 cuptiActivityEnable 或 cuptiActivityEnableContext 初始化
activity kind

这是一个大坑，早晚要来填。。。

PCI BAR : base address register

一篇关于 GPU 虚拟化 的 survey 文章，发表于 ACM Computing Surveys

Background 里面一段比较有意思的话：

On the contrary, the design of conventional processors is optimized
for reducing the execution time of sequential code on each core, thus adding complexity
to each core at the cost of offering fewer cores in the processor package. Conventional
processors typically use sophisticated control logic and large cache memories to efficiently
deal with conditional branches, pipeline stalls, and poor data locality.

传统处理器的目标是尽可能加速每个核心的串行执行时间，因此每个核心有大量处理分支预测，流水线延迟，data locality 相关的资源。代价是核心数量较少

之前做过劫持 cuda runtime 动态链接库的事情，在这里记录一下：

cudaError_t cudaMalloc ( void** devPtr, size_t size )
{
    if(!fp) {
        fp = fopen("/home/leenldk/sc/race/ASC2021-RACE/mem.out", "w");
    }
    cudaError_t (*lcudaMalloc) (void**, size_t) = (cudaError_t (*) (void**, size_t))dlsym(RTLD_NEXT, "cudaMalloc");
    cudaError_t ret = lcudaMalloc(devPtr, size);
    printf("cudaMalloc size : %#lx %#lx\n", size, (size_t)(*devPtr));
    fprintf(fp, "cudaMalloc size : %#lx %#lx\n", size, (size_t)(*devPtr));
    fflush(fp);
    return ret;
}

编译为动态链接库，使用 LD_PRELOAD 预加载

update 2022.11.2

C++ 动态库 加载，卸载 时调用函数 ：

static void init() __attribute__((constructor));
void init() {}
static void fini() __attribute__((destructor));
void fini() {}