技术背景

前面的一篇文章中介绍了在CUDA中使用宏来监测CUDA C函数或者Kernel函数的运行报错问题。同样的思路，我们可用写一个用于统计函数运行时长的宏，这样不需要使用额外的工具来对函数体的性能进行测试。

文件准备

因为这里的宏改动，主要涉及CUDA头文件和CUDA文件的修改，所以Cython文件和Python文件还有异常捕获宏我们还是复用这篇文章里面用到的。测试内容是，定义一个原始数组和一个索引数组，输出索引的结果数组。

wrapper.pyx

# cythonize -i -f wrapper.pyx

import numpy as np
cimport numpy as np
cimport cython

cdef extern from "<dlfcn.h>" nogil:
    void *dlopen(const char *, int)
    char *dlerror()
    void *dlsym(void *, const char *)
    int dlclose(void *)
    enum:
        RTLD_LAZY

ctypedef int (*GatherFunc)(float *source, int *index, float *res, int N, int M) noexcept nogil

cdef void* handle = dlopen('/path/to/libcuindex.so', RTLD_LAZY)

@cython.boundscheck(False)
@cython.wraparound(False)
cpdef float[:] cuda_gather(float[:] x, int[:] idx):
    cdef:
        GatherFunc Gather
        int success
        int N = idx.shape[0]
        int M = x.shape[0]
        float[:] res = np.zeros((N, ), dtype=np.float32)
    Gather = <GatherFunc>dlsym(handle, "Gather")
    success = Gather(&x[0], &idx[0], &res[0], N, M)
    return res

while not True:
    dlclose(handle)

test_gather.py

import numpy as np
np.random.seed(0)
from wrapper import cuda_gather

M = 1024 * 1024 * 128
N = 1024 * 1024
x = np.random.random((M,)).astype(np.float32)
idx = np.random.randint(0, M, (N,)).astype(np.int32)
res = np.asarray(cuda_gather(x, idx))
print (res.shape)
print ((res==x[idx]).sum())

error.cuh

#pragma once
#include <stdio.h>

#define CHECK(call) do{const cudaError_t error_code = call; if (error_code != cudaSuccess){printf("CUDA Error:\n"); printf("    File:   %s\n", __FILE__); printf("    Line:   %d\n", __LINE__); printf("    Error code: %d\n", error_code); printf("    Error text: %s\n", cudaGetErrorString(error_code)); exit(1);}} while (0)

计时宏

这里增加一个用于计时的record.cuh头文件，里面写一个TIME_CUDA_FUNCTION宏，然后在CUDA中需要统计的函数前调用，就可以输出CUDA函数的运行时长了。

#pragma once
#include <stdio.h>
#include <cuda_runtime.h>

// 宏定义，用于测量CUDA函数的执行时间
#define TIME_CUDA_FUNCTION(func) \
    do { \
        cudaEvent_t start, stop; \
        float elapsedTime; \
        cudaEventCreate(&start); \
        cudaEventCreate(&stop); \
        cudaEventRecord(start, NULL); \
        \
        func; \
        \
        cudaEventRecord(stop, NULL); \
        cudaEventSynchronize(stop); \
        cudaEventElapsedTime(&elapsedTime, start, stop); \
        printf("Time taken by function %s is: %f ms\n", #func, elapsedTime); \
        \
        cudaEventDestroy(start); \
        cudaEventDestroy(stop); \
    } while (0)

计时宏的使用

我们在CUDA文件cuda_index.cu中调用record.cuh里面的计时宏，这里用来统计一个CUDA核函数的执行时间：

// nvcc -shared ./cuda_index.cu -Xcompiler -fPIC -o ./libcuindex.so
#include <stdio.h>
#include "cuda_index.cuh"
#include "error.cuh"
#include "record.cuh"

void __global__ GatherKernel(float *source, int *index, float *res, int N){
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N){
        res[idx] = source[index[idx]];
    }
}

extern "C" int Gather(float *source, int *index, float *res, int N, int M){
    float *souce_device, *res_device;
    int *index_device;
    CHECK(cudaMalloc((void **)&souce_device, M * sizeof(float)));
    CHECK(cudaMalloc((void **)&res_device, N * sizeof(float)));
    CHECK(cudaMalloc((void **)&index_device, N * sizeof(int)));
    CHECK(cudaMemcpy(souce_device, source, M * sizeof(float), cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(res_device, res, N * sizeof(float), cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(index_device, index, N * sizeof(int), cudaMemcpyHostToDevice));
    int block_size = 1024;
    int grid_size = (N + block_size - 1) / block_size;
    TIME_CUDA_FUNCTION((GatherKernel<<<grid_size, block_size>>>(souce_device, index_device, res_device, N)));
    CHECK(cudaGetLastError());
    CHECK(cudaDeviceSynchronize());
    CHECK(cudaMemcpy(res, res_device, N * sizeof(float), cudaMemcpyDeviceToHost));
    CHECK(cudaFree(souce_device));
    CHECK(cudaFree(index_device));
    CHECK(cudaDeviceSynchronize());
    CHECK(cudaFree(res_device));
    CHECK(cudaDeviceReset());
    return 1;
}

需要注意的是，TIME_CUDA_FUNCTION宏只能有一个输入，但是使用CUDA核函数的时候实际上会被当作是两个输入，因此我们需要将CUDA核函数用括号再封装起来。

输出结果

最终按照这篇文章中的运行流程，可以得到这样的输出结果：

Time taken by function (GatherKernel<<<grid_size, block_size>>>(souce_device, index_device, res_device, N)) is: 0.584224 ms
(1048576,)
1048576

这里CUDA核函数的运行时长被正确的格式化输出了。

返回耗时数值

除了在CUDA中直接打印耗时的数值，我们还可以修改record.cuh中的宏，让其返回耗时数值：

#pragma once
#include <stdio.h>
#include <cuda_runtime.h>

// 宏定义，用于测量CUDA函数的执行时间
#define TIME_CUDA_FUNCTION(func) \
    do { \
        cudaEvent_t start, stop; \
        float elapsedTime; \
        cudaEventCreate(&start); \
        cudaEventCreate(&stop); \
        cudaEventRecord(start, NULL); \
        \
        func; \
        \
        cudaEventRecord(stop, NULL); \
        cudaEventSynchronize(stop); \
        cudaEventElapsedTime(&elapsedTime, start, stop); \
        printf("Time taken by function %s is: %f ms\n", #func, elapsedTime); \
        \
        cudaEventDestroy(start); \
        cudaEventDestroy(stop); \
    } while (0)

// 宏定义，用于测量CUDA函数的执行时间并返回该时间
#define GET_CUDA_TIME(func) \
    ({ \
        cudaEvent_t start, stop; \
        float elapsedTime = 0.0f; \
        cudaEventCreate(&start); \
        cudaEventCreate(&stop); \
        cudaEventRecord(start, NULL); \
        \
        func; \
        \
        cudaEventRecord(stop, NULL); \
        cudaEventSynchronize(stop); \
        cudaEventElapsedTime(&elapsedTime, start, stop); \
        \
        cudaEventDestroy(start); \
        cudaEventDestroy(stop); \
        \
        elapsedTime; \
    })

修改头文件cuda_index.cuh，因为这里我们需要返回一个运行时长的float数值，不再是int类型了：

#include <stdio.h>

extern "C" float Gather(float *source, int *index, float *res, int N, int M);

最后再对应修改下cuda_index.cu中的内容：

// nvcc -shared ./cuda_index.cu -Xcompiler -fPIC -o ./libcuindex.so
#include <stdio.h>
#include "cuda_index.cuh"
#include "error.cuh"
#include "record.cuh"

void __global__ GatherKernel(float *source, int *index, float *res, int N){
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N){
        res[idx] = source[index[idx]];
    }
}

extern "C" float Gather(float *source, int *index, float *res, int N, int M){
    float *souce_device, *res_device;
    int *index_device;
    CHECK(cudaMalloc((void **)&souce_device, M * sizeof(float)));
    CHECK(cudaMalloc((void **)&res_device, N * sizeof(float)));
    CHECK(cudaMalloc((void **)&index_device, N * sizeof(int)));
    CHECK(cudaMemcpy(souce_device, source, M * sizeof(float), cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(res_device, res, N * sizeof(float), cudaMemcpyHostToDevice));
    CHECK(cudaMemcpy(index_device, index, N * sizeof(int), cudaMemcpyHostToDevice));
    int block_size = 1024;
    int grid_size = (N + block_size - 1) / block_size;
    float timeTaken = GET_CUDA_TIME((GatherKernel<<<grid_size, block_size>>>(souce_device, index_device, res_device, N)));
    CHECK(cudaGetLastError());
    CHECK(cudaDeviceSynchronize());
    CHECK(cudaMemcpy(res, res_device, N * sizeof(float), cudaMemcpyDeviceToHost));
    CHECK(cudaFree(souce_device));
    CHECK(cudaFree(index_device));
    CHECK(cudaDeviceSynchronize());
    CHECK(cudaFree(res_device));
    CHECK(cudaDeviceReset());
    return timeTaken;
}

这样就可以把函数运行耗时的数值返回给Cython文件，然后在Cython文件wrapper.pyx中打印耗时：

# cythonize -i -f wrapper.pyx

import numpy as np
cimport numpy as np
cimport cython

cdef extern from "<dlfcn.h>" nogil:
    void *dlopen(const char *, int)
    char *dlerror()
    void *dlsym(void *, const char *)
    int dlclose(void *)
    enum:
        RTLD_LAZY

ctypedef float (*GatherFunc)(float *source, int *index, float *res, int N, int M) noexcept nogil

cdef void* handle = dlopen('/home/dechin/projects/gitee/dechin/tests/cuda/libcuindex.so', RTLD_LAZY)

@cython.boundscheck(False)
@cython.wraparound(False)
cpdef float[:] cuda_gather(float[:] x, int[:] idx):
    cdef:
        GatherFunc Gather
        float timeTaken
        int N = idx.shape[0]
        int M = x.shape[0]
        float[:] res = np.zeros((N, ), dtype=np.float32)
    Gather = <GatherFunc>dlsym(handle, "Gather")
    timeTaken = Gather(&x[0], &idx[0], &res[0], N, M)
    print (timeTaken)
    return res

while not True:
    dlclose(handle)

最后再通过Python模块调用（无需改动），输出结果为：

0.6107839941978455
(1048576,)
1048576

这里的单位是ms。

总结概要

这篇文章主要介绍了一个CUDA入门的技术：使用CUDA头文件写一个专门用于CUDA函数运行时长统计的宏，这样就可以统计目标Kernel函数的运行时长。可以直接在CUDA中打印相应的数值，也可以回传到Cython或者Python中进行打印。

版权声明

本文首发链接为：https://www.cnblogs.com/dechinphy/p/cuda-time-record.html

作者ID：DechinPhy

更多原著文章：https://www.cnblogs.com/dechinphy/

请博主喝咖啡：https://www.cnblogs.com/dechinphy/gallery/image/379634.html