Fermi 一代 GPU 的单精度计算应该比 double 快 2 倍。 然而,尽管我将所有声明“double”重写为“float”,但我并没有加速。 有什么错误吗?编译选项等..?
显卡:特斯拉C2075 操作系统:win7专业版 编译器:VS2013(nvcc) CUDA:v.7.5 命令行:nvcc test.cu
我写了测试代码:
#include<stdio.h>
#include<stdlib.h>
#include<math.h>
#include<time.h>
#include<conio.h>
#include<cuda_runtime.h>
#include<cuda_profiler_api.h>
#include<device_functions.h>
#include<device_launch_parameters.h>
#define DOUBLE 1
#define MAXI 10
__global__ void Kernel_double(double*a,int nthreadx)
{
double b=1.e0;
int i;
i = blockIdx.x * nthreadx + threadIdx.x + 0;
a[i] *= b;
}
__global__ void Kernel_float(float*a,int nthreadx)
{
float b=1.0F;
int i;
i = blockIdx.x * nthreadx + threadIdx.x + 0;
a[i] *= b;
}
int main()
{
#if DOUBLE
double a[10];
for(int i=0;i<MAXI;++i){
a[i]=1.e0;
}
double*d_a;
cudaMalloc((void**)&d_a, sizeof(double)*(MAXI));
cudaMemcpy(d_a, a, sizeof(double)*(MAXI), cudaMemcpyHostToDevice);
#else
float a[10];
for(int i=0;i<MAXI;++i){
a[i]=1.0F;
}
float*d_a;
cudaMalloc((void**)&d_a, sizeof(float)*(MAXI));
cudaMemcpy(d_a, a, sizeof(float)*(MAXI), cudaMemcpyHostToDevice);
#endif
dim3 grid(2, 2, 1);
dim3 block(2, 2, 1);
clock_t start_clock, end_clock;
double sec_clock;
printf("[%d] start\n", __LINE__);
start_clock = clock();
for (int i = 1; i <= 100000; ++i){
#if DOUBLE
Kernel_double << < grid, block >> > (d_a, 2);
cudaMemcpy(a, d_a, sizeof(double)*(MAXI), cudaMemcpyDeviceToHost);
#else
Kernel_float << < grid, block >> > (d_a, 2);
cudaMemcpy(a, d_a, sizeof(float)*(MAXI), cudaMemcpyDeviceToHost);
#endif
}
end_clock = clock();
sec_clock = (end_clock - start_clock) / (double)CLOCKS_PER_SEC;
printf("[%d] %f[s]\n", __LINE__, sec_clock);
printf("[%d] end\n", __LINE__);
return 0;
}
最佳答案
好吧,经过一些调查,那是因为你只是乘以常数 1,它在二进制文件中被优化为“什么都不做”:
如果改为对数组进行平方(以防止这种微不足道的优化),则会得到以下程序集:
下面(简化的)代码恢复了性能提升,我在其中做了一些更改:
- 更大的数组 (100M)
- 使用 blockDim.x 而不是实参参数
- 为我的机器使用更好的内核配置 (GTX 980)
- 在堆上而不是堆栈上分配输入数组(允许超过 1M)
代码如下:
#include<stdio.h>
#include<stdlib.h>
#include<math.h>
#include<time.h>
#include<conio.h>
#include<cuda_runtime.h>
#include<cuda_profiler_api.h>
#include<device_functions.h>
#include<device_launch_parameters.h>
#define DOUBLE float
#define ITER 10
#define MAXI 100000000
__global__ void kernel(DOUBLE*a)
{
for(int i = blockIdx.x * blockDim.x + threadIdx.x ; i < MAXI; i += blockDim.x * gridDim.x)
{
a[i] *= a[i];
}
}
int main()
{
DOUBLE* a = (DOUBLE*) malloc(MAXI*sizeof(DOUBLE));
for(int i=0;i<MAXI;++i)
{
a[i]=(DOUBLE)1.0;
}
DOUBLE* d_a;
cudaMalloc((void**)&d_a, sizeof(DOUBLE)*(MAXI));
cudaMemcpy(d_a, a, sizeof(DOUBLE)*(MAXI), cudaMemcpyHostToDevice);
clock_t start_clock, end_clock;
double sec_clock;
printf("[%d] start\n", __LINE__);
start_clock = clock();
for (int i = 1; i <= ITER; ++i){
kernel <<< 32, 256>>> (d_a);
}
cudaDeviceSynchronize();
end_clock = clock();
cudaMemcpy(a, d_a, sizeof(DOUBLE)*(MAXI), cudaMemcpyDeviceToHost);
sec_clock = (end_clock - start_clock) / (double)CLOCKS_PER_SEC;
printf("[%d] %f/%d[s]\n", __LINE__, sec_clock, CLOCKS_PER_SEC);
printf("[%d] end\n", __LINE__);
return 0;
}
(您会注意到我分配了一个长度为 100M 的数组以获得可衡量的性能。)
关于c - 为什么我的 cuda C 代码在单精度下没有变得更快?,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/37018634/