Cuda 性能测量 - 经过时间返回零
Posted
技术标签:
【中文标题】Cuda 性能测量 - 经过时间返回零【英文标题】：Cuda Performance measuring - Elapsed time returns zero 【发布时间】：2013-05-07 02:57:47 【问题描述】：我写了几个内核函数，想知道处理这些函数需要多少毫秒。
using namespace std;
#include <iostream>
#include <stdio.h>
#include <stdlib.h>
#define N 8000
void fillArray(int *data, int count)
for (int i = 0; i < count; i++)
data[i] = rand() % 100;
__global__ void add(int* a, int *b)
int add = 0;
int tid = threadIdx.x + blockIdx.x * blockDim.x;
if (tid < N)
add = a[tid] + b[tid];
__global__ void subtract(int* a, int *b)
int subtract = 0;
int tid = threadIdx.x + blockIdx.x * blockDim.x;
if (tid < N)
subtract = a[tid] - b[tid];
__global__ void multiply(int* a, int *b)
int multiply = 0;
int tid = threadIdx.x + blockIdx.x * blockDim.x;
if (tid < N)
multiply = a[tid] * b[tid];
__global__ void divide(int* a, int *b)
int divide = 0;
int tid = threadIdx.x + blockIdx.x * blockDim.x;
if (tid < N)
divide = a[tid] / b[tid];
__global__ void modu(int* a, int *b)
int modulus = 0;
int tid = threadIdx.x + blockIdx.x * blockDim.x;
if (tid < N)
modulus = a[tid] % b[tid];
__global__ void neg(int *data)
int tid = threadIdx.x + blockIdx.x * blockDim.x;
if (tid < N)
data[tid] = -data[tid];
float duration(int *devA, int *devB, int blocksPerGrid, int threadsPerBlock)
cudaEvent_t start, stop;
float elapsedTime;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start, 0);
add<<<blocksPerGrid, threadsPerBlock>>>(devA, devB);
subtract<<<blocksPerGrid, threadsPerBlock>>>(devA, devB);
multiply<<<blocksPerGrid, threadsPerBlock>>>(devA, devB);
divide<<<blocksPerGrid, threadsPerBlock>>>(devA, devB);
modu<<<blocksPerGrid, threadsPerBlock>>>(devA, devB);
neg<<<blocksPerGrid, threadsPerBlock>>>(devA);
neg<<<blocksPerGrid, threadsPerBlock>>>(devB);
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&elapsedTime, start, stop);
cudaEventDestroy(start);
cudaEventDestroy(stop);
return elapsedTime;
int main(void)
int a[N], b[N];
float dur = 0;
int *devA, *devB;
cudaMalloc((void**) &devA, N * sizeof(int));
cudaMalloc((void**) &devB, N * sizeof(int));
fillArray(a, N);
fillArray(b, N);
cudaMemcpy(devA, a, N * sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(devA, b, N * sizeof(int), cudaMemcpyHostToDevice);
dur = duration(a, b, N, 1);
cout << "Global memory version:\n";
cout << "Process completed in " << dur;
cout << " for a data set of " << N << " integers.";
return 0;
毫秒总是返回零。为什么？我在这里缺少什么？如果我从持续时间函数中删除否定函数。它返回 0.15687 毫秒。我认为处理这些功能的数量很少。那个程序有什么问题？
编辑后，我这样做了：
using namespace std;
#include <iostream>
#include <stdio.h>
#include <stdlib.h>
const int N = 8000;
void fillArray(int *data, int count)
for (int i = 0; i < count; i++)
data[i] = rand() % 100;
__global__ void add(int* a, int *b, int *c)
int tid = threadIdx.x + blockIdx.x * blockDim.x;
if (tid < N)
c[tid] = a[tid] + b[tid];
__global__ void subtract(int* a, int *b, int *c)
int tid = threadIdx.x + blockIdx.x * blockDim.x;
if (tid < N)
c[tid] = a[tid] - b[tid];
__global__ void multiply(int* a, int *b, int *c)
int tid = threadIdx.x + blockIdx.x * blockDim.x;
if (tid < N)
c[tid] = a[tid] * b[tid];
__global__ void divide(int* a, int *b, int *c)
int tid = threadIdx.x + blockIdx.x * blockDim.x;
if (tid < N)
c[tid] = a[tid] / b[tid];
__global__ void modu(int* a, int *b, int *c)
int tid = threadIdx.x + blockIdx.x * blockDim.x;
if (tid < N)
c[tid] = a[tid] % b[tid];
__global__ void neg(int *data, int *c)
int tid = threadIdx.x + blockIdx.x * blockDim.x;
if (tid < N)
c[tid] = -data[tid];
float duration(int *devA, int *devB, int *devC, int blocksPerGrid, int threadsPerBlock)
cudaEvent_t start, stop;
float elapsedTime;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start, 0);
double hArrayC[N];
add<<<blocksPerGrid, threadsPerBlock>>>(devA, devB,devC);
cudaMemcpy(hArrayC,devC,N*sizeof(int),cudaMemcpyDeviceToHost);
subtract<<<blocksPerGrid, threadsPerBlock>>>(devA, devB,devC);
cudaMemcpy(hArrayC,devC,N*sizeof(int),cudaMemcpyDeviceToHost);
multiply<<<blocksPerGrid, threadsPerBlock>>>(devA, devB,devC);
cudaMemcpy(hArrayC,devC,N*sizeof(int),cudaMemcpyDeviceToHost);
divide<<<blocksPerGrid, threadsPerBlock>>>(devA, devB,devC);
cudaMemcpy(hArrayC,devC,N*sizeof(int),cudaMemcpyDeviceToHost);
modu<<<blocksPerGrid, threadsPerBlock>>>(devA, devB,devC);
cudaMemcpy(hArrayC,devC,N*sizeof(int),cudaMemcpyDeviceToHost);
neg<<<blocksPerGrid, threadsPerBlock>>>(devA,devC);
cudaMemcpy(hArrayC,devC,N*sizeof(int),cudaMemcpyDeviceToHost);
neg<<<blocksPerGrid, threadsPerBlock>>>(devB,devC);
cudaMemcpy(hArrayC,devC,N*sizeof(int),cudaMemcpyDeviceToHost);
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&elapsedTime, start, stop);
cudaEventDestroy(start);
cudaEventDestroy(stop);
return elapsedTime;
int main(void)
int a[N], b[N],c[N];
float dur = 0;
int *devA, *devB,*devC;
cudaMalloc((void**) &devA, N * sizeof(int));
cudaMalloc((void**) &devB, N * sizeof(int));
cudaMalloc((void**) &devC, N * sizeof(int));
fillArray(a, N);
fillArray(b, N);
cudaMemcpy(devA, a, N * sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(devB, b, N * sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(devC, c, N * sizeof(int), cudaMemcpyHostToDevice);
dur = duration(devA, devB, devC,N, 1);
cout << "Global memory version:\n";
cout << "Process completed in " << dur;
cout << " for a data set of " << N << " integers.";
cudaFree(devA);
cudaFree(devB);
return 0;
【问题讨论】：
【参考方案1】：您的内核没有做任何事情，因为您只将结果存储在寄存器中。编译时，您会收到一些警告：
kernel.cu(13)：警告：变量“add”已设置但从未使用过
此外，如果您想查看更好的时序，请使用 NVIDIA 的分析器：nvprof (CLI) 或 nvvp (GUI)。
$ nvprof ./kernel
======== NVPROF is profiling kernel... ======== Command: kernel Global memory version: Process completed in 0 for a data set of 8000 integers. ======== Profiling result: Time(%) Time Calls Avg Min Max Name 100.00 18.46us 2 9.23us 6.02us 12.45us [CUDA memcpy HtoD] 0.00 0ns 1 0ns 0ns 0ns multiply(int*, int*) 0.00 0ns 1 0ns 0ns 0ns add(int*, int*) 0.00 0ns 1 0ns 0ns 0ns modu(int*, int*) 0.00 0ns 2 0ns 0ns 0ns neg(int*) 0.00 0ns 1 0ns 0ns 0ns subtract(int*, int*) 0.00 0ns 1 0ns 0ns 0ns divide(int*, int*)
您还在每个网格使用 N 块，每个块使用 1 个线程。你应该考虑阅读this question的答案。
更新
关于向量加法（和其他简单的操作）本身，你应该学习CUDA SDK的vectorAdd sample，或者使用Thrust。第一个选项将教您如何使用 CUDA，第二个选项将向您展示您可以使用 Thrust 执行的高级操作。如果我是你，我会两者兼得。
【讨论】：
那么如何让内核工作呢？我需要将总和值复制到主机端吗？ 您的结果与我的导师一致。我怎样才能使这些内核工作？我需要复制吗？ 真正的问题是：你想做什么？您将a 和b 提供给GPU 内核，但您不存储结果，因此您还应该提供一个包含操作结果的c 数组。计算结束后，将c 复制回主机。如果这不仅仅是 CUDA 练习，并且您想在“严肃”项目中对向量进行操作，您应该考虑使用Thrust。它更易于使用（类似于 STL），开发速度更快，并且它提供了许多对向量的操作。
能否请您再次检查问题。我编辑了问题并按照您所说的做了。对吗？
我看不出有什么区别...如果您想了解它是如何完成的，请查看 CUDA SDK 的 VectorAdd 示例。该示例中包含您需要的一切。【参考方案2】：
Cuda 任务正在设备上运行而不会阻塞 CPU 线程。因此，cuda 调用只会在您尝试从设备内存中获取计算数据且尚未准备好时才会阻塞。或者，当您使用 cudaDeviceSyncronize() 调用显式将 CPU 线程与 GPU 同步时。如果要测量计算时间，则需要在停止计时器之前进行同步。
如果您对测量内存复制时间感兴趣，您需要在计算开始之后和复制计时器开始之前进行同步，否则计算时间将显示为复制时间。
您可以使用cuda SDK 中包含的分析器来测量所有cuda 调用的时间。
【讨论】：
【参考方案3】：尝试使用float（或double）变量和数组而不是int 来存储所有算术变量和运算。有时时间间隔太小，整数值总是四舍五入为零。
【讨论】：
你能投到int 只是为了 mod 功能吗？
我将所有变量加倍，但仍返回 0以上是关于Cuda 性能测量 - 经过时间返回零的主要内容，如果未能解决你的问题，请参考以下文章