CUDA学习笔记(4) CUDA计时函数

程序员文章站 2022-04-02 10:56:11

...

精确地记录核函数运行的时间可以让我们直观地了解核函数代码的运行效率。我们需要一个可以精确测量代码运行时间的方法。一般来说，所有的电脑CPU都具有硬件计数器，可以达到微妙级的计时效果，我们可以得到核函数执行所需的时间。但是，如果我们想要知道GPU中核函数中的某几行代码的运行时间，就必须使用GPU的硬件计数器了，因为核函数完全在GPU的内部执行，CPU无法干涉。

先简单介绍下计数器工作的原理，这和《数字电路》中的计数器是同一个东西。当计数器的输入由低电平转为高电平的瞬间，计数器的值加1；当计数器的值达到它的最大值，则从0开始计数。

无论是CPU还是GPU的计数器都是开机后就会自动地开始计数，计数到到最大值又从0开始。我们通过求得这两个值之差，再除以计数的频率，就可以得到这两次计数操作之间经过的时间。
需要注意的是，假如计数器的最大值为255，那么当我们第一次读取的计数为245，第二次读取的计数为006，显然计数已经被重置过了，这时我们的计数次数实际上是第一次计数值取反加上第二次计数的值。
记第一次计数为C1，第二次计数为C2，计数器的最大值为M，计数器频率为f，两次计数的间隔时间为ΔT，伪代码：

if（C1<C2 ）
{
ΔT=(C2−C1)/f
}
else
{
ΔT=(M+C2−C1)/f
}

举个例子，假设我们需要求ai 这个数列的前128项，它的通项如下：

a i = \sum k = 1 262115 1 i 1 + sin k

其中 i∈N+,i≤128。这是一个复杂到我们无法用数学方法化简的式子，因此我们必须使用电脑计算。那么就会有至少两种方案：

使用CPU计算，在CPU中，需要计算128次 ai的等式。
使用GPU计算，在GPU中，只需要一个SM中的128个SP同时计算一次就可以了。

猜想：GPU的计算速度远快于CPU。
现在，我们分别在CPU与GPU上实现它，验证一下。

为了精确地统计我们程序运行的时间，我们有必要了解CPU上的硬件计数器使用流程：

获取CPU硬件计数器频率。
取得程序开始时的计数。
执行被测试的代码。
取得程序结束时的计数。

在使用硬件计数器时，我们需要先包含对应的头文件“windows.h”和“timer.h”。
“QueryPerformanceFrequency(&large_interger)”函数用于取得计数器的频率。
“QueryPerformanceCounter(&large_interger)”函数用于取得计数器的当前计数值。
这里，我们可以不用考虑计数器溢出的问题，因为“__int64”的最大值为+263=9.2×1018，而我们CPU的硬件计数器频率为106Hz数量级，远小于前者，所以完全可以不用考虑计数溢出的问题。

    LARGE_INTEGER  large_interger;
    double CPU_counter_frequency;
    __int64  CPU_c1_start, CPU_c2_alloc, CPU_c3_caculate, CPU_c4_end;

    // 获取频率
    QueryPerformanceFrequency(&large_interger);
    CPU_counter_frequency = large_interger.QuadPart;
    printf("CPU硬件计数器频率：\t%.2lf Hz\r\n\r\n", CPU_counter_frequency);

    // 计数
    QueryPerformanceCounter(&large_interger);
    CPU_c1_start = large_interger.QuadPart;

    /*。。。执行代码。。。*/

    // CPU 计数
    QueryPerformanceCounter(&large_interger);
    CPU_c4_end = large_interger.QuadPart;
    printf("CPU计时总时长：\t%.4lf 毫秒\r\n\r\n", \
    (double)(CPU_c4_end - CPU_c1_start) / CPU_counter_frequency * 1000);

CPU方案：
1. 初始化并分配内存。
2. CPU执行计算。
3. 释放内存。

代码实现：

    // Host变量（内存变量）
    const int ARRAY_LENGTH = 128;
    float *result_array;

#pragma region(CPU硬件计时器 开始计时)
    //用QueryPerformanceCounter()来计时  微秒
    LARGE_INTEGER  large_interger;
    double CPU_counter_frequency;
    __int64  CPU_c1_start, CPU_c2_alloc, CPU_c3_caculate, CPU_c4_end;

    QueryPerformanceFrequency(&large_interger);
    CPU_counter_frequency = large_interger.QuadPart;
    printf("CPU硬件计数器频率：\t%.2lf Hz\r\n\r\n", CPU_counter_frequency);

    // 计数
    QueryPerformanceCounter(&large_interger);
    CPU_c1_start = large_interger.QuadPart;
#pragma endregion


#pragma region(内存分配)
    result_array = (float*)malloc(ARRAY_LENGTH * sizeof(float));

    // CPU 计数
    QueryPerformanceCounter(&large_interger);
    CPU_c2_alloc = large_interger.QuadPart;
#pragma endregion


#pragma region(CPU计算)
    for (int i = 0; i < ARRAY_LENGTH; i++)
    {
        result_array[i] = 0.0;
        for (int k = 1; k < 262115; k++)
        {
            result_array[i] += 1.0 / powf((float)(i + 1), 1.0 + sin((float)k));
        }
    }

    // CPU 计数
    QueryPerformanceCounter(&large_interger);
    CPU_c3_caculate = large_interger.QuadPart;
#pragma endregion


#pragma region(释放内存、显存)
    printf("\r\n\tresult index\tresult\tcost time(ms)\r\n");
    for (int i = 0; i < ARRAY_LENGTH; i++)
    {
        printf("\t%d\t", i);
        printf("\t%f\t\r\n", result_array[i]);
    }
    cout << endl;

    // 释放内存
    free(result_array);

    // CPU 计数
    QueryPerformanceCounter(&large_interger);
    CPU_c4_end = large_interger.QuadPart;
#pragma endregion


#pragma region(CPU结束计时，输出结果)
    printf("分配内存时长：\t%.4lf 毫秒\r\n", (double)(CPU_c2_alloc - CPU_c1_start) / CPU_counter_frequency * 1000);
    printf("CPU计算时长：\t%.4lf 毫秒\r\n", (double)(CPU_c3_caculate - CPU_c2_alloc) / CPU_counter_frequency * 1000);
    printf("释放内存时长：\t%.4lf 毫秒\r\n", (double)(CPU_c4_end - CPU_c3_caculate) / CPU_counter_frequency * 1000);
    printf("CPU计时总时长：\t%.4lf 毫秒\r\n\r\n", (double)(CPU_c4_end - CPU_c1_start) / CPU_counter_frequency * 1000);
#pragma endregion

仅使用CPU做计算没有什么编程难度，但是结果如何呢，从运行结果中我们可以看到主要时间花费在计算上，耗时1.8秒。

CUDA学习笔记(4) CUDA计时函数

GPU方案：
1. 读取、编译核函数。
2. 分配内存、显存。
3. GPU计算。
4. 传出数据。
5. 释放内存、显存。

首先我们先初始化CPU的硬件计数器，用于记录各个阶段所使用的时间。

    //用QueryPerformanceCounter()来计时  微秒
    LARGE_INTEGER  large_interger;
    double CPU_counter_frequency;
    __int64  CPU_c1_start, CPU_c2_compile, CPU_c3_alloc, CPU_c4_caculate, CPU_c5_copyout, CPU_c6_end;

    QueryPerformanceFrequency(&large_interger);
    CPU_counter_frequency = large_interger.QuadPart;
    printf("CPU硬件计数器频率：\t%.2lf Hz\r\n\r\n", CPU_counter_frequency);

    // 计数
    QueryPerformanceCounter(&large_interger);
    CPU_c1_start = large_interger.QuadPart;

第一步是读取编译核函数，同CUDA学习笔记(3) NVRTC编译库的方法。记录完成时的计数为“CPU_c2_compile”。

第二步分配内存、显存。记录完成时的计数为“CPU_c3_alloc”。

第三步使用核函数在GPU进行计算。记录完成时计数为“CPU_c4_caculate”。

    // 定义网格的大小（block_rect）、块的大小（thread_rect）
    dim3 block_rect(max(1, ARRAY_LENGTH / max_thread_per_block), 1, 1);
    dim3 thread_rect(min(max_thread_per_block, ARRAY_LENGTH), 1, 1);

    void *arr[] = { (void *)&dev_thread_index_array, \
        (void *)&dev_result_array, (void *)&dev_clock_counter_array };

    // 启动核函数
    checkCudaErrors(cuLaunchKernel(kernel_addr, \
        block_rect.x, block_rect.y, block_rect.z, \
        thread_rect.x, thread_rect.y, thread_rect.z, \
        0, 0, \
        &arr[0], 0));

    // 同步，检查核函数计算过程
    checkCudaErrors(cuCtxSynchronize());

    // CPU 计数
    QueryPerformanceCounter(&large_interger);
    CPU_c4_caculate = large_interger.QuadPart;

核函数中，我们使用“clock()”函数，来读取当前GPU的计数值。为了观察“计算”部分几行代码的执行时间，我们需要将两次计数记录下来，传回内存后再输出带命令行窗口。需要注意的是用“clock()”函数计时一般不要超过2秒，一般也不允许一个核函数占用GPU太长时间（1秒以上）。在windows下，核函数占用GPU超过2秒会导致触发TDR，系统自动重置显卡。
下面是我核函数部分的代码：

// CUDA kernel fucntion file !

extern "C" __global__ void kernel_func(int *thread_index_array, float *result_array, int *clock_counter_array)
{
    int block_index = blockIdx.x + blockIdx.y * gridDim.x + blockIdx.z * gridDim.x * gridDim.y;
    int thread_index = block_index * blockDim.x * blockDim.y * blockDim.z + \
        threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.x * blockDim.y;
    // 总线程数
    int thread_num = gridDim.x * gridDim.y * gridDim.z * blockDim.x * blockDim.y * blockDim.z;
    // 记录线程号
    thread_index_array[thread_index] = thread_index;

    // 计时开始
    clock_counter_array[thread_index] = clock();

    // 计算
    result_array[thread_index] = 0.0;
    for (int k = 1; k < 262115; k++)
        result_array[thread_index] += 1.0 / powf(thread_index + 1, 1.0 + sin((float)k));

    // 计时结束
    clock_counter_array[thread_index + thread_num] = clock();
}

CUDA学习笔记(4) CUDA计时函数

通过输出的结果我们可以看到一个规律，每32个线程（Thread）的执行这个for循环的时间是相同的，也就是一个线程束（Warp）是一起作为一个整体被执行操作的。

第四步，传出数据并输出结果（注意，在对比CPU与GPU运行效率时，计时期间不能使用“printf”或“cout”输出数据，这会大大影响速度！）

    checkCudaErrors(cuMemcpyDtoH(thread_index_array, dev_thread_index_array, \
        ARRAY_LENGTH * sizeof(int)));
    checkCudaErrors(cuMemcpyDtoH(result_array, dev_result_array, \
        ARRAY_LENGTH * sizeof(float)));
    checkCudaErrors(cuMemcpyDtoH(clock_counter_array, dev_clock_counter_array, \
        2 * ARRAY_LENGTH * sizeof(unsigned int)));

    // CPU 计数
    QueryPerformanceCounter(&large_interger);
    CPU_c5_copyout = large_interger.QuadPart;

这时是需要考虑计数溢出的。因为“unsigned_int”的最大值为+232=4.2×109，GPU硬件计数器的频率就是GPU时钟的频率，这里我的GPU计数器频率是1124000kHz≈1.1×109Hz，最多可以计时3820毫秒。就可能存在这种情况：第一次计数时是一个接近最大值的数，然后计数达到最大值，计数器归零，第二次计数是一个比较小的数。我们就需要用下面这种方式来计算时间：

if（C1<C2 ）
{
ΔT=(C2−C1)/f
}
else
{
ΔT=(M+C2−C1)/f
}

for (int i = 0; i < ARRAY_LENGTH; i++)
{
    if (clock_counter_array[i + ARRAY_LENGTH] > clock_counter_array[i])
    {
        temp_float = ((float)((clock_counter_array[i + ARRAY_LENGTH] - clock_counter_array[i]))) \
            / (float)(GPU_clock_rate / 1000);
    }
    else
    {
        temp_float = ((float)((MAXUINT32 + clock_counter_array[i + ARRAY_LENGTH] \
            - clock_counter_array[i]))) \
            / (float)(GPU_clock_rate / 1000);
    }
    cost_time_array[i] = temp_float;
}

第五步，释放内存、显存，结束CPU计数。

    // 释放显存
    checkCudaErrors(cuMemFree(dev_thread_index_array));
    checkCudaErrors(cuMemFree(dev_result_array));
    checkCudaErrors(cuMemFree(dev_clock_counter_array));
    // 释放内存
    free(thread_index_array);
    free(result_array);
    free(clock_counter_array);
    free(cost_time_array);

    // CPU 计数
    QueryPerformanceCounter(&large_interger);
    CPU_c6_end = large_interger.QuadPart;

CUDA学习笔记(4) CUDA计时函数

GPU做计算仅仅用了342毫秒，总时间为697毫秒。

CUDA学习笔记(4) CUDA计时函数

可以看到利用GPU做并行计算相较于仅用CPU做计算快的多。但是相对的，编程难度也会上升，而且GPU不擅长做非线性、多重判断的算法。

总结：

CUDA中“clock()”函数可以为我们提供微秒级，乃至于纳秒级的计数。
通过结合CPU的硬件定时器，我们可以准确、清晰地统计到程序执行过程中，时间主要花费在哪一步，为我们做进一步的优化提供根据。
无论是从编程规范还是硬件支持上讲,一般不允许一个核函数占用GPU太长时间（1秒以上）。在windows下，核函数占用GPU超过2秒会导致触发TDR，系统自动重置显卡。

我的“learn_CUDA_04.cpp”完整代码：

// C/C++ IO
#include <stdio.h>
#include <iostream>
using namespace std;

// For the CUDA runtime routines (prefixed with "cuda_")
#include <cuda_runtime.h>
#include <cuda.h>

// helper functions and utilities to work with CUDA
#include <helper_cuda.h>
#include <helper_functions.h>
// NVRTC
#include <nvrtc_helper.h>

// timer
#include <timer.h>
#include <windows.h>


// GPU核心时钟频率
int GPU_clock_rate;
// Block最大线程数
int max_thread_per_block;
// 检查显卡硬件属性
void check_Cuda_information(int main_argc, char ** main_argv);
// CPU对照组
void cpu_contrast();


int main(int argc, char **argv)
{
    // Host变量（内存变量）
    const int ARRAY_LENGTH = 128;
    int *thread_index_array;
    unsigned int *clock_counter_array;
    float *result_array;
    float *cost_time_array, temp_float;

    // Device变量（显存变量）
    CUdeviceptr dev_thread_index_array, dev_result_array, dev_clock_counter_array;

    // 检查显卡硬件属性
    check_Cuda_information(argc, &argv[0]);

#pragma region(CPU硬件计时器 开始计时)
    //用QueryPerformanceCounter()来计时  微秒
    LARGE_INTEGER  large_interger;
    double CPU_counter_frequency;
    __int64
        CPU_c1_start, CPU_c2_compile, CPU_c3_alloc, CPU_c4_caculate, CPU_c5_copyout, CPU_c6_end;

    QueryPerformanceFrequency(&large_interger);
    CPU_counter_frequency = large_interger.QuadPart;
    printf("CPU硬件计数器频率：\t%.2lf Hz\r\n\r\n", CPU_counter_frequency);

    // 计数
    QueryPerformanceCounter(&large_interger);
    CPU_c1_start = large_interger.QuadPart;
#pragma endregion


#pragma region(读取、编译CUDA核函数)
    // 读取、编译、加载CUDA核函数
    char *ptx, *kernel_file;
    size_t ptxSize;

    // 打开核函数的 .cu 文件 并编译为PTX
    kernel_file = sdkFindFilePath("kernel.cu", argv[0]);
    // 如果找不到 kernel_file 会报错：error: unable to open ./xxxx.cu for reading!
    compileFileToPTX(kernel_file, NULL, NULL, &ptx, &ptxSize, 0);

    // 选择GPU设备，加载核函数到GPU设备
    CUmodule module = loadPTX(ptx, argc, argv);
    CUfunction kernel_addr;
    checkCudaErrors(cuModuleGetFunction(&kernel_addr, module, "kernel_func"));

    // CPU 计数
    QueryPerformanceCounter(&large_interger);
    CPU_c2_compile = large_interger.QuadPart;
#pragma endregion


#pragma region(内存、显存分配)
    thread_index_array = (int*)malloc(ARRAY_LENGTH * sizeof(int));
    result_array = (float*)malloc(ARRAY_LENGTH * sizeof(float));
    clock_counter_array = (unsigned int*)malloc(2 * ARRAY_LENGTH * sizeof(unsigned int));
    cost_time_array = (float*)malloc(ARRAY_LENGTH * sizeof(float));

    checkCudaErrors(cuMemAlloc(&dev_thread_index_array, ARRAY_LENGTH * sizeof(int)));
    checkCudaErrors(cuMemAlloc(&dev_result_array, ARRAY_LENGTH * sizeof(float)));
    checkCudaErrors(cuMemAlloc(&dev_clock_counter_array, 2 * ARRAY_LENGTH * sizeof(unsigned int)));

    // CPU 计数
    QueryPerformanceCounter(&large_interger);
    CPU_c3_alloc = large_interger.QuadPart;
#pragma endregion


#pragma region(执行核函数，并在核函数完成时检查错误报告)
    // 定义网格的大小（block_rect）、块的大小（thread_rect）
    dim3 block_rect(max(1, ARRAY_LENGTH / max_thread_per_block), 1, 1);
    dim3 thread_rect(min(max_thread_per_block, ARRAY_LENGTH), 1, 1);

    cout << "block_rect :\t" << block_rect.x << "\t" << block_rect.y << "\t" << block_rect.z << "\t" << endl;
    cout << "thread_rect :\t" << thread_rect.x << "\t" << thread_rect.y << "\t" << thread_rect.z << "\t" << endl;

    void *arr[] = { (void *)&dev_thread_index_array, (void *)&dev_result_array, (void *)&dev_clock_counter_array };

    // 启动核函数
    checkCudaErrors(cuLaunchKernel(kernel_addr, \
        block_rect.x, block_rect.y, block_rect.z, \
        thread_rect.x, thread_rect.y, thread_rect.z, \
        0, 0, \
        &arr[0], 0));

    // 同步，检查核函数计算过程
    checkCudaErrors(cuCtxSynchronize());

    // CPU 计数
    QueryPerformanceCounter(&large_interger);
    CPU_c4_caculate = large_interger.QuadPart;
#pragma endregion


#pragma region(传出数据，输出结果)
    checkCudaErrors(cuMemcpyDtoH(thread_index_array, dev_thread_index_array, ARRAY_LENGTH * sizeof(int)));
    checkCudaErrors(cuMemcpyDtoH(result_array, dev_result_array, ARRAY_LENGTH * sizeof(float)));
    checkCudaErrors(cuMemcpyDtoH(clock_counter_array, dev_clock_counter_array, 2 * ARRAY_LENGTH * sizeof(unsigned int)));

    /*
    // 输出结果
    printf("\r\n\tthread index\tresult\t\tclock counter\tcost time(ms)\r\n");
    for (int i = 0; i < ARRAY_LENGTH; i++)
    {
        printf("\t%d\t", thread_index_array[i]);
        printf("\t%f\t", result_array[i]);
        printf("%u", \
            ((clock_counter_array[i + ARRAY_LENGTH] - clock_counter_array[i])));

        if (clock_counter_array[i + ARRAY_LENGTH] > clock_counter_array[i])
        {
            temp_float = ((float)((clock_counter_array[i + ARRAY_LENGTH] - clock_counter_array[i])))\
                / (float)(GPU_clock_rate / 1000);
        }
        else
        {
            temp_float = ((float)((MAXUINT32 + clock_counter_array[i + ARRAY_LENGTH] - clock_counter_array[i])))\
                / (float)(GPU_clock_rate / 1000);
        }

        cost_time_array[i] = temp_float;
        printf("\t%f\r\n", cost_time_array[i]);
    }
    cout << endl;
    for (int i = 0; i < ARRAY_LENGTH; i++)
        printf("\t%u\r\n", clock_counter_array[i]);

    cout << endl << endl;
    for (int i = ARRAY_LENGTH; i < 2 * ARRAY_LENGTH; i++)
        printf("\t%u\r\n", clock_counter_array[i]);
    */

    // CPU 计数
    QueryPerformanceCounter(&large_interger);
    CPU_c5_copyout = large_interger.QuadPart;
#pragma endregion


#pragma region(释放内存、显存)
    // 释放显存
    checkCudaErrors(cuMemFree(dev_thread_index_array));
    checkCudaErrors(cuMemFree(dev_result_array));
    checkCudaErrors(cuMemFree(dev_clock_counter_array));
    // 释放内存
    free(thread_index_array);
    free(result_array);
    free(clock_counter_array);
    free(cost_time_array);

    // CPU 计数
    QueryPerformanceCounter(&large_interger);
    CPU_c6_end = large_interger.QuadPart;
#pragma endregion


#pragma region(CPU结束计时，输出结果)

    printf("编译.cu时长：\t%.4lf 毫秒\r\n", (double)(CPU_c2_compile - CPU_c1_start) / CPU_counter_frequency * 1000);
    printf("分配显存时长：\t%.4lf 毫秒\r\n", (double)(CPU_c3_alloc - CPU_c2_compile) / CPU_counter_frequency * 1000);
    printf("GPU计算时长：\t%.4lf 毫秒\r\n", (double)(CPU_c4_caculate - CPU_c3_alloc) / CPU_counter_frequency * 1000);
    printf("传出数据时长：\t%.4lf 毫秒\r\n", (double)(CPU_c5_copyout - CPU_c4_caculate) / CPU_counter_frequency * 1000);
    printf("释放内存时长：\t%.4lf 毫秒\r\n", (double)(CPU_c6_end - CPU_c5_copyout) / CPU_counter_frequency * 1000);
    printf("CPU计时总时长：\t%.4lf 毫秒\r\n\r\n", (double)(CPU_c6_end - CPU_c1_start) / CPU_counter_frequency * 1000);
#pragma endregion


    // 执行 CPU 对照组
    cpu_contrast();

    return 0;
}



// 检查显卡硬件属性
void check_Cuda_information(int main_argc, char ** main_argv)
{
    // 设备ID
    int devID;
    // 设备属性
    cudaDeviceProp deviceProps;

    //
    cout << "argc = " << main_argc << endl;
    for (int i = 0; i < main_argc; i++)
    {
        printf("argv[%d] = %s\r\n", i, main_argv[i]);
    }
    cout << endl;

    // 获取设备ID
    devID = findCudaDevice(main_argc, (const char **)main_argv);

    // 获取GPU信息
    checkCudaErrors((CUresult)cudaGetDeviceProperties(&deviceProps, devID));
    cout << "devID = " << devID << endl;
    // 显卡名称
    cout << "CUDA device is \t\t\t" << deviceProps.name << endl;
    // 每个 线程块（Block）中的最大线程数
    cout << "CUDA max Thread per Block is \t" << deviceProps.maxThreadsPerBlock << endl;
    max_thread_per_block = deviceProps.maxThreadsPerBlock;
    // 每个 多处理器组（MultiProcessor）中的最大线程数
    cout << "CUDA max Thread per SM is \t" << deviceProps.maxThreadsPerMultiProcessor << endl;
    // GPU 中 SM 的数量
    cout << "CUDA SM counter\t\t\t" << deviceProps.multiProcessorCount << endl;
    // 线程束大小
    cout << "CUDA Warp size is \t\t" << deviceProps.warpSize << endl;
    // 每个SM*享内存的大小
    cout << "CUDA shared memorize is \t" << deviceProps.sharedMemPerMultiprocessor << "\tbyte" << endl;
    // GPU时钟频率
    //  int    clockRate;       /**< Clock frequency in kilohertz */
    cout << "GPU clock frequency is \t\t" << deviceProps.clockRate << "\tkHz" << endl;
    GPU_clock_rate = deviceProps.clockRate * 1000;


    cout << endl;
}


void cpu_contrast()
{
    // Host变量（内存变量）
    const int ARRAY_LENGTH = 128;
    float *result_array;

#pragma region(CPU硬件计时器 开始计时)
    //用QueryPerformanceCounter()来计时  微秒
    LARGE_INTEGER  large_interger;
    double CPU_counter_frequency;
    __int64  CPU_c1_start, CPU_c2_alloc, CPU_c3_caculate, CPU_c4_end;

    QueryPerformanceFrequency(&large_interger);
    CPU_counter_frequency = large_interger.QuadPart;
    printf("CPU硬件计数器频率：\t%.2lf Hz\r\n\r\n", CPU_counter_frequency);

    // 计数
    QueryPerformanceCounter(&large_interger);
    CPU_c1_start = large_interger.QuadPart;
#pragma endregion


#pragma region(内存分配)
    result_array = (float*)malloc(ARRAY_LENGTH * sizeof(float));

    // CPU 计数
    QueryPerformanceCounter(&large_interger);
    CPU_c2_alloc = large_interger.QuadPart;
#pragma endregion


#pragma region(CPU计算)
    for (int i = 0; i < ARRAY_LENGTH; i++)
    {
        result_array[i] = 0.0;
        for (int k = 1; k < 262115; k++)
        {
            result_array[i] += 1.0 / powf((float)(i + 1), 1.0 + sin((float)k));
        }
    }

    // CPU 计数
    QueryPerformanceCounter(&large_interger);
    CPU_c3_caculate = large_interger.QuadPart;
#pragma endregion


#pragma region(释放内存、显存)
    /*
    // 输出结果
    printf("\r\n\tresult index\tresult\tcost time(ms)\r\n");
    for (int i = 0; i < ARRAY_LENGTH; i++)
    {
        printf("\t%d\t", i);
        printf("\t%f\t\r\n", result_array[i]);
    }
    cout << endl;
    */

    // 释放内存
    free(result_array);

    // CPU 计数
    QueryPerformanceCounter(&large_interger);
    CPU_c4_end = large_interger.QuadPart;
#pragma endregion


#pragma region(CPU结束计时)

    printf("分配内存时长：\t%.4lf 毫秒\r\n", (double)(CPU_c2_alloc - CPU_c1_start) / CPU_counter_frequency * 1000);
    printf("CPU计算时长：\t%.4lf 毫秒\r\n", (double)(CPU_c3_caculate - CPU_c2_alloc) / CPU_counter_frequency * 1000);
    printf("释放内存时长：\t%.4lf 毫秒\r\n", (double)(CPU_c4_end - CPU_c3_caculate) / CPU_counter_frequency * 1000);
    printf("CPU计时总时长：\t%.4lf 毫秒\r\n\r\n", (double)(CPU_c4_end - CPU_c1_start) / CPU_counter_frequency * 1000);  
#pragma endregion
}

参考：

1.《CUDA并行程序设计》机械工业出版社

2.NVIDIA官方NVRTC说明文档

3.CUDA Toolkit Documation

4.NVIDIA官方CUDA核函数常用数学函数

相关标签： cuda cpu 计时 gpu clock

上一篇： L1-027 出租 (20分)

下一篇： Python中countdown倒计时实现

CUDA学习笔记(4) CUDA计时函数

tensorflow学习笔记(北京大学) tf4_5.py 完全解析设损失函数 loss=(w+1)^2

CUDA 学习笔记 (一) [Chapter3]

javascript学习全过程-----js学习笔记（4）------函数（上）

CUDA从入门到精通到精通_笔记4:GPU设备属性查询的代码

Flutter学习笔记（4）--Dart函数

Cocos2d-x学习笔记（14）（更新函数scheduleUpdate、进度计时器C

Python学习笔记（4）～Python基础练习之常用内置函数（1-10）

Oracle学习笔记4--单行函数

体验笔记本MX350显卡配置深度学习环境（CUDA+tensorflow）

2021年研究生深度学习 30系显卡笔记本 Cuda CuDnn Pytorch 等环境配置日志

CUDA学习笔记(4) CUDA计时函数

tensorflow学习笔记(北京大学) tf4_5.py 完全解析 设损失函数 loss=(w+1)^2

CUDA 学习笔记 (一) [Chapter3]

javascript学习全过程-----js学习笔记（4）------函数（上）

CUDA从入门到精通到精通_笔记4:GPU设备属性查询的代码

Flutter学习笔记（4）--Dart函数

Cocos2d-x学习笔记（14）（更新函数scheduleUpdate、进度计时器C

Python学习笔记（4）～Python基础练习之常用内置函数（1-10）

Oracle学习笔记4--单行函数

体验笔记本MX350显卡配置深度学习环境（CUDA+tensorflow）

2021年研究生 深度学习 30系显卡笔记本 Cuda CuDnn Pytorch 等环境配置日志

tensorflow学习笔记(北京大学) tf4_5.py 完全解析设损失函数 loss=(w+1)^2

2021年研究生深度学习 30系显卡笔记本 Cuda CuDnn Pytorch 等环境配置日志