如何增加使用 CUDA 实现的 FLOPS

所以我刚刚进入 CUDA(我已经使用 c++ 好几年了，但我是处理 GPU 的新手，所以请原谅我缺乏经验(。我正在用 NVIDIA GEFORCE GTX860M 显卡在我的计算机上编程 3D n 体模拟。该卡公布的峰值理论 FP32 性能为 1,389 GFLOPS (https://www.techpowerup.com/gpu-specs/geforce-gtx-860m.c2536(。我正在使用下面的代码来大致计算我可以实现多少"有效 FLOPS"，目前我在使用全局内存时只得到 7.100 GFLOPS，在使用共享内存时获得 5.100 GFLOPS。我的印象是共享内存比全局内存快 100 倍，那么为什么我没有看到 FLOPS 的增加呢？

旁注 1：我假设"cudaFunction"中的每个线程每次内核调用执行大约 100,000 个浮点操作。因此，(512 * 128( 线程 * (100000( FP32 操作/(1.285( 秒 = 5.100 GFLOPS。

旁注 2：我意识到我可能没有正确测量 FLOPS，但我的目标是最大化单位时间内所有 CUDA 线程中完成的浮点计算数量，因此我将这个数量称为"有效 FLOPS"。

我的第二个问题是，我可以期望达到什么样的有效翻牌率，我可以实施什么样的优化才能将我的 5.1 GFLOPS 提高到更接近公布的最大值？ 0.37%(5.1 GFLOPS/1389 GFLOPS(的峰值似乎很低，所以我假设我在某个地方遇到了瓶颈？

#include <cuda_runtime.h>
#include <iostream>
#include <time.h>
#include <math.h>
#include "device_launch_parameters.h"
#include <iomanip>
#include <cuda.h>

#define numPtcls 512*128//Total number of particles
#define threadsPerBlock 128//Number of threads per block
#define BLOCKS numPtcls / threadsPerBlock//total number of blocks
using namespace std;
struct Particles {
float testVariable;
};
//USING SHARED MEMORY
__global__ void cudaFunction(Particles *particle)
{
int index = threadIdx.x + blockIdx.x * blockDim.x;
float sum = 1;
__shared__ float position;// Allocate share memory
position = particle[0].testVariable;
for (int i = 0; i < 100000; i++) {
sum *= position;
}
particle[0].testVariable = 1;
}
////USING GLOBAL MEMORY
//__global__ void cudaFunction(Particles *particle)
//{
//  int index = threadIdx.x + blockIdx.x * blockDim.x;
//  float sum = 1;
//
//  for (int i = 0; i < 100000; i++) {
//      sum *= particle[0].testVariable;
//  }
//
//  particle[0].testVariable = 1;
//}
int main()
{
Particles *particle = new Particles[numPtcls];
particle[0].testVariable = 1;
Particles *device_location;//POINTER TO MEMORY FOR CUDA
int size = numPtcls * sizeof(Particles);//SIZE OF PARTICLE DATA TO MAKE ROOM FOR IN CUDA
cudaMalloc((void**)&device_location, size);// allocate device copies
cudaMemcpy(device_location, particle, size, cudaMemcpyHostToDevice);// copy inputs to device
clock_t start, end;
double cpu_time_used;
while (true) {
start = clock();
cudaFunction << <BLOCKS, threadsPerBlock >> > (device_location);//CUDA CALL
cudaMemcpy(particle, device_location, size, cudaMemcpyDeviceToHost);
end = clock();
cpu_time_used = ((double)(end - start)) / CLOCKS_PER_SEC;
std::cout << fixed << setprecision(6) << cpu_time_used << std::endl;
}
cudaFree(device_location);//FREE DEVICE MEMORY
delete[] particle;//FREE CPU MEMORY
return 0;
}

TLDR：性能测量很难正确 - 你使用的代码，你如何编译它，以及你如何计时都很重要。

至少，您的尝试有很多错误：

除非内核循环的结果参与内存写入，否则编译器优化会将浮点计算视为死代码并删除它们
除非你为发布而不是调试而编译，否则对代码进行基准测试是没有意义的，因为它消除了所有编译器优化
的使用完全无关紧要，因为编译器无论如何都会在寄存器中的循环中缓存结果，并且在这种情况下，使用共享内存不会优化内存事务模式

clock

而不是挂钟时间，因此使用它来计时不消耗 CPU 周期的 GPU 上的异步操作是无效
是破碎的，但也包括memcpy时间，如果你的目标是测量内核中的FLOP，这是无效的。

修复上述所有问题，这让我明白了：

$ cat floppy.cu 
#include <iostream>
#include <iomanip>
#include <cmath>
#include <limits>
#define numPtcls (512*128) //Total number of particles
#define threadsPerBlock (128) //Number of threads per block
#define BLOCKS numPtcls / threadsPerBlock//total number of blocks
#define niters (10000) // FMAD iterations per thread

struct Particles {
float testVariable;
};
__global__ void cudaFunction(Particles *particle)
{
int index = threadIdx.x + blockIdx.x * blockDim.x;
float sum = 1;
float position = particle[0].testVariable;
for (int i = 0; i < niters; i++) {
sum *= position;
}
particle[0].testVariable = sum;
}
int main()
{
Particles *particle = new Particles[numPtcls];
particle[0].testVariable = 1;
Particles *device_location;//POINTER TO MEMORY FOR CUDA
int size = numPtcls * sizeof(Particles);//SIZE OF PARTICLE DATA TO MAKE ROOM FOR IN CUDA
cudaMalloc((void**)&device_location, size);// allocate device copies
cudaMemcpy(device_location, particle, size, cudaMemcpyHostToDevice);// copy inputs to device
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
float flopcount = float(niters) * float(numPtcls);
for(int i=0; i<10; i++) {
cudaEventRecord(start, 0);
cudaFunction << <BLOCKS, threadsPerBlock >> > (device_location);//CUDA CALL
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
cudaMemcpy(particle, device_location, size, cudaMemcpyDeviceToHost);
float gpu_time_used;
cudaEventElapsedTime(&gpu_time_used, start, stop);
std::cout << std::fixed << std::setprecision(6) << 1e-6 * (flopcount / gpu_time_used) << std::endl;
}
cudaFree(device_location);//FREE DEVICE MEMORY
delete[] particle;//FREE CPU MEMORY
return 0;
}

这只是对你所拥有的非常适度的修改(基本上将内核的结果存储到内存中以击败死代码删除，使用 CUDA 事件对内核进行计时(

$ nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2019 NVIDIA Corporation
Built on Fri_Feb__8_19:08:17_PST_2019
Cuda compilation tools, release 10.1, V10.1.105
$ nvcc -arch=sm_52 -std=c++11 -Xptxas="-v" -o floppy floppy.cu 
floppy.cu(18): warning: variable "index" was declared but never referenced
ptxas info    : 0 bytes gmem
ptxas info    : Compiling entry function '_Z12cudaFunctionP9Particles' for 'sm_52'
ptxas info    : Function properties for _Z12cudaFunctionP9Particles
0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 6 registers, 328 bytes cmem[0]

$ ./floppy 
1557.296000
1534.312192
1575.505792
1547.762944
1541.820288
1555.521792
1561.808896
1545.193856
1545.543680
1581.345152

这个相当幼稚的代码在大约 4 毫秒内运行，在我的 GTX970 上让我大约 1550 GFLOP/s，这大约是我用来运行它的设备上大约 4000 GFLOP/s 峰值的 40%。编译器发出的代码值得检查：

.version 6.4
.target sm_52
.address_size 64
// .globl   _Z12cudaFunctionP9Particles
.visible .entry _Z12cudaFunctionP9Particles(
.param .u64 _Z12cudaFunctionP9Particles_param_0
)
{
.reg .pred  %p<2>;
.reg .f32   %f<55>;
.reg .b32   %r<5>;
.reg .b64   %rd<3>;

ld.param.u64    %rd2, [_Z12cudaFunctionP9Particles_param_0];
cvta.to.global.u64  %rd1, %rd2;
ld.global.f32   %f1, [%rd1];
mov.f32     %f54, 0f3F800000;
mov.u32     %r4, -10000;
BB0_1:
mul.f32     %f5, %f1, %f54;
mul.f32     %f6, %f1, %f5;
mul.f32     %f7, %f1, %f6;
mul.f32     %f8, %f1, %f7;
mul.f32     %f9, %f1, %f8;
mul.f32     %f10, %f1, %f9;
mul.f32     %f11, %f1, %f10;
mul.f32     %f12, %f1, %f11;
mul.f32     %f13, %f1, %f12;
mul.f32     %f14, %f1, %f13;
mul.f32     %f15, %f1, %f14;
mul.f32     %f16, %f1, %f15;
mul.f32     %f17, %f1, %f16;
mul.f32     %f18, %f1, %f17;
mul.f32     %f19, %f1, %f18;
mul.f32     %f20, %f1, %f19;
mul.f32     %f21, %f1, %f20;
mul.f32     %f22, %f1, %f21;
mul.f32     %f23, %f1, %f22;
mul.f32     %f24, %f1, %f23;
mul.f32     %f25, %f1, %f24;
mul.f32     %f26, %f1, %f25;
mul.f32     %f27, %f1, %f26;
mul.f32     %f28, %f1, %f27;
mul.f32     %f29, %f1, %f28;
mul.f32     %f30, %f1, %f29;
mul.f32     %f31, %f1, %f30;
mul.f32     %f32, %f1, %f31;
mul.f32     %f33, %f1, %f32;
mul.f32     %f34, %f1, %f33;
mul.f32     %f35, %f1, %f34;
mul.f32     %f36, %f1, %f35;
mul.f32     %f37, %f1, %f36;
mul.f32     %f38, %f1, %f37;
mul.f32     %f39, %f1, %f38;
mul.f32     %f40, %f1, %f39;
mul.f32     %f41, %f1, %f40;
mul.f32     %f42, %f1, %f41;
mul.f32     %f43, %f1, %f42;
mul.f32     %f44, %f1, %f43;
mul.f32     %f45, %f1, %f44;
mul.f32     %f46, %f1, %f45;
mul.f32     %f47, %f1, %f46;
mul.f32     %f48, %f1, %f47;
mul.f32     %f49, %f1, %f48;
mul.f32     %f50, %f1, %f49;
mul.f32     %f51, %f1, %f50;
mul.f32     %f52, %f1, %f51;
mul.f32     %f53, %f1, %f52;
mul.f32     %f54, %f1, %f53;
add.s32     %r4, %r4, 50;
setp.ne.s32 %p1, %r4, 0;
@%p1 bra    BB0_1;
st.global.f32   [%rd1], %f54;
ret;
}

您可以看到编译器已将循环展开为一长串单精度mul指令，这些指令以每个时钟周期 1 个或每个内核每个时钟周期 1 个 FLOP 的速率停用。请注意，如果您将内核更改为以下内容：

__global__ void cudaFunction(Particles *particle)
{
int index = threadIdx.x + blockIdx.x * blockDim.x;
float sum = 1;
float position = particle[0].testVariable;
for (int i = 0; i < niters; i++) {
sum += sum * position;
}
particle[0].testVariable = sum;
}

编译将发出以下内容：

.version 6.4
.target sm_52
.address_size 64
// .globl   _Z12cudaFunctionP9Particles
.visible .entry _Z12cudaFunctionP9Particles(
.param .u64 _Z12cudaFunctionP9Particles_param_0
)
{
.reg .pred  %p<2>;
.reg .f32   %f<45>;
.reg .b32   %r<5>;
.reg .b64   %rd<5>;

ld.param.u64    %rd2, [_Z12cudaFunctionP9Particles_param_0];
cvta.to.global.u64  %rd1, %rd2;
ld.global.f32   %f1, [%rd1];
mov.f32     %f44, 0f3F800000;
mov.u32     %r4, -10000;
BB0_1:
fma.rn.f32  %f5, %f1, %f44, %f44;
fma.rn.f32  %f6, %f1, %f5, %f5;
fma.rn.f32  %f7, %f1, %f6, %f6;
fma.rn.f32  %f8, %f1, %f7, %f7;
fma.rn.f32  %f9, %f1, %f8, %f8;
fma.rn.f32  %f10, %f1, %f9, %f9;
fma.rn.f32  %f11, %f1, %f10, %f10;
fma.rn.f32  %f12, %f1, %f11, %f11;
fma.rn.f32  %f13, %f1, %f12, %f12;
fma.rn.f32  %f14, %f1, %f13, %f13;
fma.rn.f32  %f15, %f1, %f14, %f14;
fma.rn.f32  %f16, %f1, %f15, %f15;
fma.rn.f32  %f17, %f1, %f16, %f16;
fma.rn.f32  %f18, %f1, %f17, %f17;
fma.rn.f32  %f19, %f1, %f18, %f18;
fma.rn.f32  %f20, %f1, %f19, %f19;
fma.rn.f32  %f21, %f1, %f20, %f20;
fma.rn.f32  %f22, %f1, %f21, %f21;
fma.rn.f32  %f23, %f1, %f22, %f22;
fma.rn.f32  %f24, %f1, %f23, %f23;
fma.rn.f32  %f25, %f1, %f24, %f24;
fma.rn.f32  %f26, %f1, %f25, %f25;
fma.rn.f32  %f27, %f1, %f26, %f26;
fma.rn.f32  %f28, %f1, %f27, %f27;
fma.rn.f32  %f29, %f1, %f28, %f28;
fma.rn.f32  %f30, %f1, %f29, %f29;
fma.rn.f32  %f31, %f1, %f30, %f30;
fma.rn.f32  %f32, %f1, %f31, %f31;
fma.rn.f32  %f33, %f1, %f32, %f32;
fma.rn.f32  %f34, %f1, %f33, %f33;
fma.rn.f32  %f35, %f1, %f34, %f34;
fma.rn.f32  %f36, %f1, %f35, %f35;
fma.rn.f32  %f37, %f1, %f36, %f36;
fma.rn.f32  %f38, %f1, %f37, %f37;
fma.rn.f32  %f39, %f1, %f38, %f38;
fma.rn.f32  %f40, %f1, %f39, %f39;
fma.rn.f32  %f41, %f1, %f40, %f40;
fma.rn.f32  %f42, %f1, %f41, %f41;
fma.rn.f32  %f43, %f1, %f42, %f42;
fma.rn.f32  %f44, %f1, %f43, %f43;
add.s32     %r4, %r4, 40;
setp.ne.s32 %p1, %r4, 0;
@%p1 bra    BB0_1;
ld.param.u64    %rd4, [_Z12cudaFunctionP9Particles_param_0];
cvta.to.global.u64  %rd3, %rd4;
st.global.f32   [%rd3], %f44;
ret;
}

请注意，mul指令现在已被替换为fma(融合乘加(，后者仍然以每个时钟周期 1 次的速率退役，但每个周期每个内核执行 2 次 FLOP(即每单位时间将浮点运算加倍(。在这种情况下，上述代码中的操作计数更改为：

float flopcount = 2.0f * float(niters) * float(numPtcls);

此版本的代码与原始代码运行的时间相同，但现在执行的 FLOP 数量增加了一倍：

$ ./floppy 
3158.544128
3134.614016
3083.408640
3098.570240
3100.915968
3089.688576
3182.842368
3108.682496
3139.659520
3098.570240

这代表了我的器件理论峰值的 80%(这也是基于融合的单精度乘加指令(。

最后，为了进行比较，以下是为设备调试编译的最佳性能代码：

$ nvcc -arch=sm_52 -std=c++11 -G -o floppy floppy.cu
$ ./floppy 
66.823832
69.371288
67.816480
69.234680
68.168728
76.703976
79.013264
78.954016
79.187560
77.139656

即性能从峰值的约80%下降到峰值的约2%。编译器发出的代码具有指导意义：

.visible .entry _Z12cudaFunctionP9Particles(
.param .u64 _Z12cudaFunctionP9Particles_param_0
)
{
.reg .pred  %p<3>;
.reg .f32   %f<9>;
.reg .b32   %r<12>;
.reg .b64   %rd<2>;

.loc 1 16 1
func_begin6:
.loc    1 0 0
.loc 1 16 1
ld.param.u64    %rd1, [_Z12cudaFunctionP9Particles_param_0];
func_exec_begin6:
.loc    1 18 15
tmp12:
mov.u32     %r4, %tid.x;
mov.u32     %r5, %ctaid.x;
mov.u32     %r6, %ntid.x;
mul.lo.s32  %r7, %r5, %r6;
add.s32     %r8, %r4, %r7;
mov.b32     %r9, %r8;
tmp13:
mov.f32     %f5, 0f3F800000;
.loc    1 19 15
mov.f32     %f1, %f5;
tmp14:
.loc    1 21 20
ld.f32  %f6, [%rd1];
mov.f32     %f2, %f6;
tmp15:
.loc    1 23 16
mov.u32     %r10, 0;
mov.b32     %r1, %r10;
tmp16:
mov.f32     %f8, %f1;
tmp17:
mov.u32     %r11, %r1;
tmp18:
BB6_1:
.loc    1 23 5
mov.u32     %r2, %r11;
mov.f32     %f3, %f8;
tmp19:
setp.lt.s32 %p1, %r2, 10000;
not.pred    %p2, %p1;
@%p2 bra    BB6_4;
bra.uni     BB6_2;
BB6_2:
.loc    1 24 9
tmp20:
mul.f32     %f7, %f3, %f2;
add.f32     %f4, %f3, %f7;
tmp21:
.loc    1 23 34
add.s32     %r3, %r2, 1;
tmp22:
mov.f32     %f8, %f4;
tmp23:
mov.u32     %r11, %r3;
tmp24:
bra.uni     BB6_1;
tmp25:
BB6_4:
.loc    1 27 5
st.f32  [%rd1], %f3;
.loc    1 28 1
ret;
tmp26:
func_end6:
}

循环展开被抑制，融合的乘加指令被替换为单独的mul和add。永远不要低估编译器优化的强大功能 - 在这里编译器免费为您提供大约 40 倍的性能提升。忽视这一点，后果自负。

相关内容

最新更新

热门标签：