CUDA时长统计
技术背景前面的一篇文章中介绍了在CUDA中使用宏来监测CUDA C函数或者Kernel函数的运行报错问题。同样的思路，我们可用写一个用于统计函数运行时长的宏，这样不需要使用额外的工具来对函数体的性能进行测试。
文件准备
因为这里的宏改动，主要涉及CUDA头文件和CUDA文件的修改，所以Cython文件和Python文件还有异常捕获宏我们还是复用这篇文章里面用到的。测试内容是，定义一个原始数组和一个索引数组，输出索引的结果数组。
wrapper.pyx
# cythonize -i -f wrapper.pyximport numpy as npcimport numpy as npcimport cythoncdef extern from "<dlfcn.h>" nogil: void *dlopen(const char *, int) char *dlerror() void *dlsym(void *, const char *) int dlclose(void *) enum: RTLD_LAZYctypedef int (*GatherFunc)(float *source, int *index, float *res, int N, int M) noexcept nogilcdef 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 int M = x.shape float[:] res = np.zeros((N, ), dtype=np.float32) Gather = <GatherFunc>dlsym(handle, "Gather") success = Gather(&x, &idx, &res, N, M) return reswhile not True: dlclose(handle)test_gather.py
import numpy as npnp.random.seed(0)from wrapper import cuda_gatherM = 1024 * 1024 * 128N = 1024 * 1024x = 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).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 = source]; }}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 = source]; }}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.pyximport numpy as npcimport numpy as npcimport cythoncdef extern from "<dlfcn.h>" nogil: void *dlopen(const char *, int) char *dlerror() void *dlsym(void *, const char *) int dlclose(void *) enum: RTLD_LAZYctypedef float (*GatherFunc)(float *source, int *index, float *res, int N, int M) noexcept nogilcdef 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 int M = x.shape float[:] res = np.zeros((N, ), dtype=np.float32) Gather = <GatherFunc>dlsym(handle, "Gather") timeTaken = Gather(&x, &idx, &res, N, M) print (timeTaken) return reswhile 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
页:
[1]