opencl初始化速度提升：kernel编译和clFFT初始化
opencl的kernel脚本一般在应用程序运行时调用clCreateProgramWithSource、clBuildProgram进行编译，若脚本中程序较多，则编译时消耗一定的时间。
在桌面级GPU中，这个时间消耗很小，一般在10~30ms，但在嵌入式GPGPU中，编译时间会很长，达到10秒以上，严重影响系统的初始化时间。
opencl提供了clCreateProgramWithBinary接口，用于减少初始化时间。
使用时，首先通过clGetProgramInfo函数将已经编译好的程序保存下来，下次启动时，调用clCreateProgramWithBinary加载，即可大幅降低编译时间。
1 int save_bin(const char *filename) //在已经完成编译的条件下，保存成文件 2 { 3 size_t binary_size = 0; 4 clGetProgramInfo( program_, CL_PROGRAM_BINARY_SIZES, //获得本程序的大小 5 sizeof(size_t), &binary_size, NULL); 6 7 vector<u8> program_binary(binary_size); 8 u8* binary_ptrs[] = { program_binary.data() }; 9 clGetProgramInfo( program_, CL_PROGRAM_BINARIES, 10 sizeof(program_binary.data()), binary_ptrs, NULL); 11 12 ofstream ofs(filename, ios::out | ios::binary); 13 if (!ofs.is_open()) 14 { 15 return 1; 16 } 17 18 ofs.write((const char*)program_binary.data(), binary_size); 19 ofs.close(); 20 return 0; 21 } 22 int load_bin(const char *filename) //从二进制文件读取程序，失败返回非0 23 { 24 CComFile file; 25 file.open(filename, "rb"); 26 vector<u8> buf(file.len); 27 file.read(buf.data(), file.len); 28 file.close(); 29 30 cl_int err= 0; 31 cl_int status= 0; 32 cl_device_id dev_ids[1] = {device->id}; 33 size_t len=buf.size(); 34 const u8* binary_ptrs[] = { buf.data() }; 35 program_ = clCreateProgramWithBinary(device->ctx, 1, dev_ids, &len, binary_ptrs, &status, &err); 36 if (program_ == nullptr) { return 1; } 37 //编译Program对象 38 const char* opt_str=" -cl-mad-enable"; 39 //DBGL; 40 err = clBuildProgram(program_, 1, dev_ids, opt_str, NULL, NULL); 41 //DBGL; 42 if (err != CL_SUCCESS) 43 { 44 char build_log[KB(16)]={0}; 45 clGetProgramBuildInfo(program_, dev_ids[0], CL_PROGRAM_BUILD_LOG, KB(16), build_log, NULL); 46 THROW<<sFormat("fail to build program, ret=%d %s", err, build_log).c_str(); 47 } 48 return 0; 49 }
下次程序启动时，首先判断是否有保存的bin文件，若有，则加载，若没有则编译，并存储为bin文件。
保存的文件其实也不是二进制文件，而是文本指令：
1 // 2 // Generated by NVIDIA NVVM Compiler 3 // 4 // Compiler Build ID: UNKNOWN 5 // Unknown Toolkit Version 6 // Based on NVVM 7.0.1 7 // 8 9 .version 8.3 10 .target sm_86, texmode_independent 11 .address_size 64 12 13 // .globl s16_2_float 14 // transpose_$_tile has been demoted 15 // sar_1_R_$_tile has been demoted 16 .global .align 4 .b8 __cudart_i2opi_f[24] = {65, 144, 67, 60, 153, 149, 98, 219, 192, 221, 52, 245, 209, 87, 39, 252, 41, 21, 68, 78, 110, 131, 249, 162}; 17 18 .entry s16_2_float( 19 .param .u32 s16_2_float_param_0, 20 .param .u64 .ptr .global .align 2 s16_2_float_param_1, 21 .param .u64 .ptr .global .align 4 s16_2_float_param_2 22 ) 23 { 24 .reg .b16 %rs<3>; 25 .reg .f32 %f<3>; 26 .reg .b32 %r<15>; 27 .reg .b64 %rd<12>; 28 29 30 ld.param.u32 %rd1, [s16_2_float_param_0]; 31 ld.param.u64 %rd2, [s16_2_float_param_1]; 32 ld.param.u64 %rd3, [s16_2_float_param_2]; 33 mov.b32 %r1, %envreg3; 34 mov.u32 %r2, %ntid.x; 35 mov.u32 %r3, %ctaid.x; 36 mov.u32 %r4, %tid.x; 37 add.s32 %r5, %r4, %r1; 38 mad.lo.s32 %r6, %r2, %r3, %r5; 39 cvt.s64.s32 %rd4, %r6; 40 mov.b32 %r7, %envreg4; 41 mov.u32 %r8, %ntid.y; 42 mov.u32 %r9, %ctaid.y; 43 mov.u32 %r10, %tid.y; 44 add.s32 %r11, %r10, %r7; 45 mad.lo.s32 %r12, %r8, %r9, %r11; 46 mov.b32 %r13, %envreg6; 47 mul.lo.s32 %r14, %r13, %r2; 48 mul.wide.s32 %rd5, %r12, %r14; 49 add.s64 %rd6, %rd5, %rd4; 50 add.s64 %rd7, %rd6, %rd1; 51 shl.b64 %rd8, %rd7, 2; 52 add.s64 %rd9, %rd2, %rd8; 53 ld.global.u16 %rs1, [%rd9]; 54 ld.global.u16 %rs2, [%rd9+2]; 55 // begin inline asm 56 cvt.rn.f32.s16 %f1, %rs1; 57 // end inline asm 58 // begin inline asm 59 cvt.rn.f32.s16 %f2, %rs2; 60 // end inline asm 61 shl.b64 %rd10, %rd6, 3; 62 add.s64 %rd11, %rd3, %rd10; 63 st.global.f32 [%rd11], %f1; 64 st.global.f32 [%rd11+4], %f2; 65 ret; 66 67 }
经过优化后，在英伟达RTX 3060上测试，发现对于大约700行的kernel，编译时间由20~30ms，缩短到了10~20ms，比较随机。
但在某国产化GPGPU-SOC上测试时，发现clBuildProgram函数接受的有效代码（除注释）有限，大约500行就不行了（不报错，程序直接退出，显示Killed），代码数量越多，编译速度越慢。应该是有什么东西限制。如果一个程序太大，需要分多个部分编译。
编译时间12s左右，直接加载上次编译好的bin文件，时间降低到了0.7秒左右，效果明显。
对于clFFT的初始化，也采用相同的办法，只是clFFT的初始化内部带有cache机制，只需设置环境变量使能cache即可。
clFFT的初始化主要是Bake过程，如果不bake，直接使用plan，首次执行时间就会很长，总时间与先bake后执行一致。
在某国产化GPGPU-SOC上测试clFFT的初始化速度，设置环境变量：export CLFFT_CACHE_PATH=.
让clFFT初始化bake的时候，将编译结果缓存至指定目录。下次启动时，先从缓存目录读取文件，若没有文件，才进行编译。
clFFT的启动速度由原来的4秒左右，降低到了3.8秒左右，测试了4组不同的FFT配置，都是缩短200ms。所以clFFT的启动速度没有明显提升。将clFFT的初始化放在辅线程中执行，结果最后一个plan的初始化与第一个plan的计算发生了并行，导致两者执行速度都变慢了，总时间与串行一致。
