CUDA编程（一）

目的

目前xid使用的是全比对的方式，也就是每张图片的特征值会和底库中每张图片的特征值进行比较，然后打出一个相似度的分数；

目前比较过程是cpu并行完成的，本文尝试使用在相同过程的提速效应；

使用nvdia的cuda并行计算框架；

环境

host : gpu001.hogpu.cc

cpu的物理指标：

?

processor       : 0 vendor_id       : GenuineIntel cpu family      : 6 model           : 79 model name      : Intel(R) Xeon(R) CPU E5-2630 v4 @ 2.20GHz stepping        : 1 microcode       : 0xb00002e cpu MHz         : 1200.842 cache size      : 25600 KB physical id     : 0 siblings        : 20 core id         : 0 cpu cores       : 10 apicid          : 0 initial apicid  : 0 fpu             : yes fpu_exception   : yes cpuid level     : 20 wp              : yes flags           : fpu vme de pse tsc msr pae mce cx8 apic sep mtrr pge mca cmov pat pse36 clflush dts acpi mmx fxsr sse sse2 ss ht tm pbe syscall nx pdpe1gb rdtscp lm constant_tsc arch_perfmon pebs bts rep_good nopl xtopology nonstop_tsc aperfmperf eagerfpu pni pclmulqdq dtes64 monitor ds_cpl vmx smx est tm2 ssse3 sdbg fma cx16 xtpr pdcm pcid dca sse4_1 sse4_2 x2apic movbe popcnt tsc_deadline_timer aes xsave avx f16c rdrand lahf_lm abm 3dnowprefetch epb cat_l3 cdp_l3 intel_ppin intel_pt ssbd ibrs ibpb stibp tpr_shadow vnmi flexpriority ept vpid fsgsbase tsc_adjust bmi1 hle avx2 smep bmi2 erms invpcid rtm cqm rdt_a rdseed adx smap xsaveopt cqm_llc cqm_occup_llc cqm_mbm_total cqm_mbm_local dtherm ida arat pln pts spec_ctrl intel_stibp flush_l1d bogomips        : 4399.64 clflush size    : 64 cache_alignment : 64 address sizes   : 46 bits physical, 48 bits virtual power management:

gpu物理指标：

?

total gpu number is 4 general info for device 0 name: GeForce GTX 1080 Ti compute capability: 6.1 clock rate: 1582000 total global memory: 11721506816 total const memory: 65536 multiprocessor count: 28 shared memory per block: 49152 max thread per block: 1024 is integrated: 0 map host memory: 1

比较过程

考虑到比较过程是每条请求最为耗时的地方，因此本次测试只对该处进行了测试，摘录之前一个版本的源代码精简如下:

?

std::vector<DistNode> ComputeAllDistance(const float *all_center, // 底库特征的平均值                                          const float *all_variance, // 每张底库的偏差                                          size_t feat_size, // 特征的维度，目前模型中该值固定为256                                          size_t num, // 底库的数量                                          const feat_t &feat // 进行比对的特征值) {   std::vector<DistNode> all_dis(num);   for (size_t i = 0; i < num; ++i) {     float dis = 0.0f;     for (size_t j=0 ; j < feat_size ; ++j) {         dis += all_center[feat_size * i +j] * feat[j];     }     all_dis[i].data = 2 - 2*dis;   // 计算的分数     all_dis[i].index = i;   }   return all_dis; }

可以看出cpu版本的时间复杂度为O(feat_size*num)；

gpu的代码和cpu稍有不同

?

__global__ void  // “__global__”表示由主机到gpu的函数调用，也就是下面这个函数是运行在gpu上的 scoreCalc(const float *features, // 底库特征值，指向gpu的内存空间           const float *item, // 进行比对的特征值，指向gpu的内存空间           float *score, // 输出分数           int total_num,  // 总的底库数量           int size) // 特征的维度，目前模型中该值固定为256 {     int i = blockDim.x * blockIdx.x + threadIdx.x;     if (i < total_num)     {         score[i] = 0.0;         double dis = 0.0;         for(int j=0; j!=size; j++) {             dis += features[i*size + j] * item[j];         }         score[i] = 2 - 2*dis;     } }

实验

实验一. 单cpu core和单块gpu的性能对比实验

底库数量	cpu（毫秒）	gpu（毫秒）
50万	521	11
500万	3901	94

小结：gpu有40x的提速；

 

threadPerBlock	gpu（毫秒）
32	6.48
64	9.40
128	9.37
256	9.69
512	9.55

注：该时间未考虑将结果从device拷贝回host的步骤

实验二. 多gpu的性能测试

通常服务器上会挂载多个gpu，因此想看一下不同数量gpu对于整体性能的提升，还是选用50万底库

	单卡（微秒）	双卡（微秒）	四卡（微秒）
耗时	10955	5789	2960

小结：gpu的并行性能还是很好的，基本同个数成线性增长

实验三. 使用常量内存对于性能的影响

threadPerBlock	default pu（毫秒）	const memory gpu（毫秒）
32	6.48	5.61
64	9.40	9.56
128	9.37	9.82
256	9.69	10.46
512	9.55	10.48

实验四. 纹理内存对于性能的影响

threadPerBlock	default pu（毫秒）	const memory gpu（毫秒）	texture memory gpu（毫秒）	texture + const gpu（毫秒）
32	6.48	5.61	4.60	6.20
64	9.40	9.56	5.94	12.89
128	9.37	9.82	6.00	12.90
256	9.69	10.46	6.09	13.97
512	9.55	10.48	6.16	14.84

分析

1. 在选定的服务器上，单个gpu能获取到40倍左右的提速；
2. gpu的性能基本随着gpu的个数线性增长；
3. 本服务器上的cpu物理性能一般，参考开发机，其主频只有开发机的一半。因此具体的提速倍数是和硬件的性能相关的；
4. 测试的数据还同模型相关，目前的选用的模型只是其中的一个版本；
5. 无论是gpu，还是cpu编程，应该都有优化的空间；