之前做过劫持 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() {}

linux 将外设抽象为 /dev 下的文件，通过统一的文件读写接口访问外设

linux与外设通信：

• I/O 端口：通过 I/O 读写访问设备
• I/O 内存映射：对 I/O 端口进行内存映射，将外设地址映射到内存地址，PCI 总线寻址通过内存映射完成
• 中断

- 阅读剩余部分 -

memory

managed memory : 使用 linux heterogeneous memory management (HMM)， device 和 host 端可以以相同指针访问同一块内存

coherent memory : 可以在 kernel 运行时执行对 host 和其他 peer 可见的原子操作，通过不 cache 内存实现
non-coherent memory : device 端 cache 的内存，修改不实时可见

调度

direct dispatch : runtime 直接将操作发送至 AQL 队列
device side malloc

编译与链接

HIP 支持两种 static lib：

• 只包含 host 函数，可以使用 gcc 等非 hipcc 编译器链接
• 包含 device 函数，只能使用 hipcc 链接

driver api v.s. runtime api

driver api 更为细粒度，例如 runtime 中所有 kernel 初始化时自动 load 且程序运行时保持 load ，driver api 可以只保持当前需要 load 的 module

从用户接口上：
driver api 通常返回 CUresult
runtime api 通常返回 cudaError_t

- 阅读剩余部分 -

代码版本：linux-4.9.327
（最老的一个 long term 版）
多年来第一次鼓起勇气开 linux 源码
感觉也没有那么吓人（？）
不得不感慨 open source 以及 free software 真的是好文明

- 阅读剩余部分 -