You can not select more than 25 topics Topics must start with a chinese character,a letter or number, can include dashes ('-') and can be up to 35 characters long.

mat.cpp 14 kB

adreno image shader + fp16 + fp16a (#1714) * wip * wip * fix * image and imageview can not be destroyed until command execution ends * fast copy path for tightly packed data * wip * texture load works * 1d 3d image * record clone image, multiple commands share one image reference * upload download image * layer forward accept vkimagemat * vkimagemat graph works * staging vkimagemat for passing dynamic parameters, macro for fp32+image shader, padding image shader * vkimagemat elemsize * convolution test pass * conv1x1s1 image shader * fast staging image allocator from host memory, pooling image shader * convolutiondepthwise image shader * innerproduct image shader * packing image shader * crop deconvolution image shader * resolve spirv binding types * image fp16 and fp16a, cast image shader * eltwise image shader * wip * absval image shader * deconvolutiondepthwise image shader * concat image shader, squeezenet works * noop split image shader * uniform precision hint * layer support_image_storage * wip * vulkan device utility operator * command is storage and packing option aware * fallback to cpu on image allocation failed, mobilenetssd works * flatten image shader, enable more test * ci test * check imgfp32 imgfp16 imgfp16a features * fix ci test * fix ci test * upgrade swiftshader * wip * opt aggressive * imgfp16p * opt none * convolution winograd image shader * fix flush range, fast copy path for continous buffer * minor fix * fix innerproduct * wip ... * wip * cast fix * packing test * wip * image fp16p is fp16p * wip * silence * more line info * code clean * softmax image shader
6 years ago
7 years ago
7 years ago
7 years ago
7 years ago
7 years ago
123456789101112131415161718192021222324252627282930313233343536373839404142434445464748495051525354555657585960616263646566676869707172737475767778798081828384858687888990919293949596979899100101102103104105106107108109110111112113114115116117118119120121122123124125126127128129130131132133134135136137138139140141142143144145146147148149150151152153154155156157158159160161162163164165166167168169170171172173174175176177178179180181182183184185186187188189190191192193194195196197198199200201202203204205206207208209210211212213214215216217218219220221222223224225226227228229230231232233234235236237238239240241242243244245246247248249250251252253254255256257258259260261262263264265266267268269270271272273274275276277278279280281282283284285286287288289290291292293294295296297298299300301302303304305306307308309310311312313314315316317318319320321322323324325326327328329330331332333334335336337338339340341342343344345346347348349350351352353354355356357358359360361362363364365366367368369370371372373374375376377378379380381382383384385386387388389390391392393394395396397398399400401402403404405406407408409410411412413414415416417418419420421422423424425426427428429430431432433434435436437438439440441442443444445446447448449450451452453454455456457458459460461462463464465466467468469470471472473474475476477478479480481482483484485486487488489490491492493494495496497498499500501502503504505506507508509510511512513514515516517518519520521522523524525526527528529530531532533534535536537538539540541542543544545546547548549550551552553554555556557558559560561562563564565566567568569570571572573574575576577578579580581582583584585586587588589590591592593594595596597598599600601602603604605606607608609610611612613614615616617618619620621622
  1. // Tencent is pleased to support the open source community by making ncnn available.
  2. //
  3. // Copyright (C) 2017 THL A29 Limited, a Tencent company. All rights reserved.
  4. //
  5. // Licensed under the BSD 3-Clause License (the "License"); you may not use this file except
  6. // in compliance with the License. You may obtain a copy of the License at
  7. //
  8. // https://opensource.org/licenses/BSD-3-Clause
  9. //
  10. // Unless required by applicable law or agreed to in writing, software distributed
  11. // under the License is distributed on an "AS IS" BASIS, WITHOUT WARRANTIES OR
  12. // CONDITIONS OF ANY KIND, either express or implied. See the License for the
  13. // specific language governing permissions and limitations under the License.
  14. #include "mat.h"
  15. #if __ARM_NEON
  16. #include <arm_neon.h>
  17. #endif // __ARM_NEON
  18. #include "cpu.h"
  19. #include "layer.h"
  20. #include "layer_type.h"
  21. #include <math.h>
  22. #if NCNN_VULKAN
  23. #if NCNN_PLATFORM_API
  24. #if __ANDROID_API__ >= 26
  25. #include <android/hardware_buffer.h>
  26. #endif // __ANDROID_API__ >= 26
  27. #endif // NCNN_PLATFORM_API
  28. #endif // NCNN_VULKAN
  29. namespace ncnn {
  30. void Mat::substract_mean_normalize(const float* mean_vals, const float* norm_vals)
  31. {
  32. Layer* op;
  33. if (mean_vals && !norm_vals)
  34. {
  35. // substract mean only
  36. op = create_layer(LayerType::Bias);
  37. ParamDict pd;
  38. pd.set(0, c);
  39. op->load_param(pd);
  40. Mat weights[1];
  41. weights[0] = Mat(c);
  42. for (int q = 0; q < c; q++)
  43. {
  44. weights[0][q] = -mean_vals[q];
  45. }
  46. op->load_model(ModelBinFromMatArray(weights));
  47. }
  48. else if (!mean_vals && norm_vals)
  49. {
  50. // normalize only
  51. op = create_layer(LayerType::Scale);
  52. ParamDict pd;
  53. pd.set(0, c);
  54. op->load_param(pd);
  55. Mat weights[1];
  56. weights[0] = Mat(c);
  57. for (int q = 0; q < c; q++)
  58. {
  59. weights[0][q] = norm_vals[q];
  60. }
  61. op->load_model(ModelBinFromMatArray(weights));
  62. }
  63. else if (mean_vals && norm_vals)
  64. {
  65. // substract mean and normalize
  66. op = create_layer(LayerType::Scale);
  67. ParamDict pd;
  68. pd.set(0, c);
  69. pd.set(1, 1);
  70. op->load_param(pd);
  71. Mat weights[2];
  72. weights[0] = Mat(c);
  73. weights[1] = Mat(c);
  74. for (int q = 0; q < c; q++)
  75. {
  76. weights[0][q] = norm_vals[q];
  77. weights[1][q] = -mean_vals[q] * norm_vals[q];
  78. }
  79. op->load_model(ModelBinFromMatArray(weights));
  80. }
  81. else // if (!mean_vals && !norm_vals)
  82. {
  83. return;
  84. }
  85. Option opt;
  86. opt.num_threads = 1; // TODO
  87. op->create_pipeline(opt);
  88. op->forward_inplace(*this, opt);
  89. op->destroy_pipeline(opt);
  90. delete op;
  91. }
  92. Mat Mat::from_float16(const unsigned short* data, int size)
  93. {
  94. Mat m(size);
  95. if (m.empty())
  96. return m;
  97. float* ptr = m; //.data;
  98. #if __ARM_NEON && (__ARM_FP & 2)
  99. int nn = cpu_support_arm_vfpv4() ? size >> 2 : 0;
  100. int remain = size - (nn << 2);
  101. #else
  102. int remain = size;
  103. #endif // __ARM_NEON
  104. #if __ARM_NEON && (__ARM_FP & 2)
  105. #if __aarch64__
  106. if (nn > 0)
  107. {
  108. asm volatile(
  109. "0: \n"
  110. "ld1 {v0.4h}, [%1], #8 \n"
  111. "fcvtl v1.4s, v0.4h \n"
  112. "subs %w0, %w0, #1 \n"
  113. "st1 {v1.4s}, [%2], #16 \n"
  114. "bne 0b \n"
  115. : "=r"(nn), // %0
  116. "=r"(data), // %1
  117. "=r"(ptr) // %2
  118. : "0"(nn),
  119. "1"(data),
  120. "2"(ptr)
  121. : "cc", "memory", "v0", "v1");
  122. }
  123. #else
  124. if (nn > 0)
  125. {
  126. asm volatile(
  127. "0: \n"
  128. "pld [%1, #64] \n"
  129. "vld1.s16 {d0}, [%1 :64]! \n"
  130. "vcvt.f32.f16 q1, d0 \n"
  131. "subs %0, #1 \n"
  132. "vst1.f32 {d2-d3}, [%2 :128]! \n"
  133. "bne 0b \n"
  134. : "=r"(nn), // %0
  135. "=r"(data), // %1
  136. "=r"(ptr) // %2
  137. : "0"(nn),
  138. "1"(data),
  139. "2"(ptr)
  140. : "cc", "memory", "q0", "q1");
  141. }
  142. #endif // __aarch64__
  143. #endif // __ARM_NEON
  144. for (; remain > 0; remain--)
  145. {
  146. *ptr = float16_to_float32(*data);
  147. data++;
  148. ptr++;
  149. }
  150. return m;
  151. }
  152. #if NCNN_VULKAN
  153. #if NCNN_PLATFORM_API
  154. #if __ANDROID_API__ >= 26
  155. VkImageMat VkImageMat::from_android_hardware_buffer(VkAndroidHardwareBufferImageAllocator* allocator)
  156. {
  157. int width = allocator->width();
  158. int height = allocator->height();
  159. return VkImageMat(width, height, allocator);
  160. }
  161. #endif // __ANDROID_API__ >= 26
  162. #endif // NCNN_PLATFORM_API
  163. #endif // NCNN_VULKAN
  164. unsigned short float32_to_float16(float value)
  165. {
  166. // 1 : 8 : 23
  167. union
  168. {
  169. unsigned int u;
  170. float f;
  171. } tmp;
  172. tmp.f = value;
  173. // 1 : 8 : 23
  174. unsigned short sign = (tmp.u & 0x80000000) >> 31;
  175. unsigned short exponent = (tmp.u & 0x7F800000) >> 23;
  176. unsigned int significand = tmp.u & 0x7FFFFF;
  177. // NCNN_LOGE("%d %d %d", sign, exponent, significand);
  178. // 1 : 5 : 10
  179. unsigned short fp16;
  180. if (exponent == 0)
  181. {
  182. // zero or denormal, always underflow
  183. fp16 = (sign << 15) | (0x00 << 10) | 0x00;
  184. }
  185. else if (exponent == 0xFF)
  186. {
  187. // infinity or NaN
  188. fp16 = (sign << 15) | (0x1F << 10) | (significand ? 0x200 : 0x00);
  189. }
  190. else
  191. {
  192. // normalized
  193. short newexp = exponent + (-127 + 15);
  194. if (newexp >= 31)
  195. {
  196. // overflow, return infinity
  197. fp16 = (sign << 15) | (0x1F << 10) | 0x00;
  198. }
  199. else if (newexp <= 0)
  200. {
  201. // underflow
  202. if (newexp >= -10)
  203. {
  204. // denormal half-precision
  205. unsigned short sig = (significand | 0x800000) >> (14 - newexp);
  206. fp16 = (sign << 15) | (0x00 << 10) | sig;
  207. }
  208. else
  209. {
  210. // underflow
  211. fp16 = (sign << 15) | (0x00 << 10) | 0x00;
  212. }
  213. }
  214. else
  215. {
  216. fp16 = (sign << 15) | (newexp << 10) | (significand >> 13);
  217. }
  218. }
  219. return fp16;
  220. }
  221. float float16_to_float32(unsigned short value)
  222. {
  223. // 1 : 5 : 10
  224. unsigned short sign = (value & 0x8000) >> 15;
  225. unsigned short exponent = (value & 0x7c00) >> 10;
  226. unsigned short significand = value & 0x03FF;
  227. // NCNN_LOGE("%d %d %d", sign, exponent, significand);
  228. // 1 : 8 : 23
  229. union
  230. {
  231. unsigned int u;
  232. float f;
  233. } tmp;
  234. if (exponent == 0)
  235. {
  236. if (significand == 0)
  237. {
  238. // zero
  239. tmp.u = (sign << 31);
  240. }
  241. else
  242. {
  243. // denormal
  244. exponent = 0;
  245. // find non-zero bit
  246. while ((significand & 0x200) == 0)
  247. {
  248. significand <<= 1;
  249. exponent++;
  250. }
  251. significand <<= 1;
  252. significand &= 0x3FF;
  253. tmp.u = (sign << 31) | ((-exponent + (-15 + 127)) << 23) | (significand << 13);
  254. }
  255. }
  256. else if (exponent == 0x1F)
  257. {
  258. // infinity or NaN
  259. tmp.u = (sign << 31) | (0xFF << 23) | (significand << 13);
  260. }
  261. else
  262. {
  263. // normalized
  264. tmp.u = (sign << 31) | ((exponent + (-15 + 127)) << 23) | (significand << 13);
  265. }
  266. return tmp.f;
  267. }
  268. void copy_make_border(const Mat& src, Mat& dst, int top, int bottom, int left, int right, int type, float v, const Option& opt)
  269. {
  270. Layer* padding = create_layer(LayerType::Padding);
  271. ParamDict pd;
  272. pd.set(0, top);
  273. pd.set(1, bottom);
  274. pd.set(2, left);
  275. pd.set(3, right);
  276. pd.set(4, type);
  277. pd.set(5, v);
  278. padding->load_param(pd);
  279. padding->create_pipeline(opt);
  280. padding->forward(src, dst, opt);
  281. padding->destroy_pipeline(opt);
  282. delete padding;
  283. }
  284. void copy_cut_border(const Mat& src, Mat& dst, int top, int bottom, int left, int right, const Option& opt)
  285. {
  286. if (left + right > src.w || top + bottom > src.h)
  287. {
  288. NCNN_LOGE("copy_cut_border parameter error, top: %d, bottom: %d, left: %d, right: %d, src.w: %d, src.h: %d", top, bottom, left, right, src.w, src.h);
  289. return;
  290. }
  291. Layer* crop = create_layer(LayerType::Crop);
  292. ParamDict pd;
  293. pd.set(0, left);
  294. pd.set(1, top);
  295. pd.set(2, 0);
  296. pd.set(3, src.w - left - right);
  297. pd.set(4, src.h - top - bottom);
  298. pd.set(5, -233);
  299. crop->load_param(pd);
  300. crop->create_pipeline(opt);
  301. crop->forward(src, dst, opt);
  302. crop->destroy_pipeline(opt);
  303. delete crop;
  304. }
  305. void resize_nearest(const Mat& src, Mat& dst, int w, int h, const Option& opt)
  306. {
  307. Layer* interp = create_layer(LayerType::Interp);
  308. ParamDict pd;
  309. pd.set(0, 1);
  310. pd.set(3, h);
  311. pd.set(4, w);
  312. interp->load_param(pd);
  313. interp->create_pipeline(opt);
  314. interp->forward(src, dst, opt);
  315. interp->destroy_pipeline(opt);
  316. delete interp;
  317. }
  318. void resize_bilinear(const Mat& src, Mat& dst, int w, int h, const Option& opt)
  319. {
  320. Layer* interp = create_layer(LayerType::Interp);
  321. ParamDict pd;
  322. pd.set(0, 2);
  323. pd.set(3, h);
  324. pd.set(4, w);
  325. interp->load_param(pd);
  326. interp->create_pipeline(opt);
  327. interp->forward(src, dst, opt);
  328. interp->destroy_pipeline(opt);
  329. delete interp;
  330. }
  331. void resize_bicubic(const Mat& src, Mat& dst, int w, int h, const Option& opt)
  332. {
  333. Layer* interp = create_layer(LayerType::Interp);
  334. ParamDict pd;
  335. pd.set(0, 3);
  336. pd.set(3, h);
  337. pd.set(4, w);
  338. interp->load_param(pd);
  339. interp->create_pipeline(opt);
  340. interp->forward(src, dst, opt);
  341. interp->destroy_pipeline(opt);
  342. delete interp;
  343. }
  344. void convert_packing(const Mat& src, Mat& dst, int _elempack, const Option& opt)
  345. {
  346. Layer* packing = create_layer(LayerType::Packing);
  347. ParamDict pd;
  348. pd.set(0, _elempack);
  349. packing->load_param(pd);
  350. packing->create_pipeline(opt);
  351. packing->forward(src, dst, opt);
  352. packing->destroy_pipeline(opt);
  353. delete packing;
  354. }
  355. void flatten(const Mat& src, Mat& dst, const Option& opt)
  356. {
  357. Layer* flatten = create_layer(LayerType::Flatten);
  358. ParamDict pd;
  359. flatten->load_param(pd);
  360. flatten->create_pipeline(opt);
  361. flatten->forward(src, dst, opt);
  362. flatten->destroy_pipeline(opt);
  363. delete flatten;
  364. }
  365. void cast_float32_to_float16(const Mat& src, Mat& dst, const Option& opt)
  366. {
  367. Layer* cast = create_layer(LayerType::Cast);
  368. ParamDict pd;
  369. pd.set(0, 1);
  370. pd.set(1, 2);
  371. cast->load_param(pd);
  372. cast->create_pipeline(opt);
  373. cast->forward(src, dst, opt);
  374. cast->destroy_pipeline(opt);
  375. delete cast;
  376. }
  377. void cast_float16_to_float32(const Mat& src, Mat& dst, const Option& opt)
  378. {
  379. Layer* cast = create_layer(LayerType::Cast);
  380. ParamDict pd;
  381. pd.set(0, 2);
  382. pd.set(1, 1);
  383. cast->load_param(pd);
  384. cast->create_pipeline(opt);
  385. cast->forward(src, dst, opt);
  386. cast->destroy_pipeline(opt);
  387. delete cast;
  388. }
  389. void cast_int8_to_float32(const Mat& src, Mat& dst, const Option& opt)
  390. {
  391. Layer* cast = create_layer(LayerType::Cast);
  392. ParamDict pd;
  393. pd.set(0, 3);
  394. pd.set(1, 1);
  395. cast->load_param(pd);
  396. cast->create_pipeline(opt);
  397. cast->forward(src, dst, opt);
  398. cast->destroy_pipeline(opt);
  399. delete cast;
  400. }
  401. void cast_float32_to_bfloat16(const Mat& src, Mat& dst, const Option& opt)
  402. {
  403. Layer* cast = create_layer(LayerType::Cast);
  404. ParamDict pd;
  405. pd.set(0, 1);
  406. pd.set(1, 4);
  407. cast->load_param(pd);
  408. cast->create_pipeline(opt);
  409. cast->forward(src, dst, opt);
  410. cast->destroy_pipeline(opt);
  411. delete cast;
  412. }
  413. void cast_bfloat16_to_float32(const Mat& src, Mat& dst, const Option& opt)
  414. {
  415. Layer* cast = create_layer(LayerType::Cast);
  416. ParamDict pd;
  417. pd.set(0, 4);
  418. pd.set(1, 1);
  419. cast->load_param(pd);
  420. cast->create_pipeline(opt);
  421. cast->forward(src, dst, opt);
  422. cast->destroy_pipeline(opt);
  423. delete cast;
  424. }
  425. void quantize_float32_to_int8(const Mat& src, Mat& dst, float scale, const Option& opt)
  426. {
  427. Layer* quantize = create_layer(LayerType::Quantize);
  428. ParamDict pd;
  429. pd.set(0, scale);
  430. quantize->load_param(pd);
  431. quantize->create_pipeline(opt);
  432. quantize->forward(src, dst, opt);
  433. quantize->destroy_pipeline(opt);
  434. delete quantize;
  435. }
  436. void dequantize_int32_to_float32(Mat& m, float scale, const float* bias, int bias_data_size, const Option& opt)
  437. {
  438. Layer* dequantize = create_layer(LayerType::Dequantize);
  439. ParamDict pd;
  440. pd.set(0, scale);
  441. pd.set(1, bias ? 1 : 0);
  442. pd.set(2, bias_data_size);
  443. dequantize->load_param(pd);
  444. Mat weights[1];
  445. weights[0] = Mat(bias_data_size, (void*)bias);
  446. dequantize->load_model(ModelBinFromMatArray(weights));
  447. dequantize->create_pipeline(opt);
  448. dequantize->forward_inplace(m, opt);
  449. dequantize->destroy_pipeline(opt);
  450. delete dequantize;
  451. }
  452. void requantize_int8_to_int8(const Mat& src, Mat& dst, float scale_in, float scale_out, const float* bias, int bias_data_size, int fusion_relu, const Option& opt)
  453. {
  454. Layer* requantize = create_layer(LayerType::Requantize);
  455. ParamDict pd;
  456. pd.set(0, scale_in);
  457. pd.set(1, scale_out);
  458. pd.set(2, bias ? 1 : 0);
  459. pd.set(3, bias_data_size);
  460. pd.set(4, fusion_relu);
  461. requantize->load_param(pd);
  462. Mat weights[1];
  463. weights[0] = Mat(bias_data_size, (void*)bias);
  464. requantize->load_model(ModelBinFromMatArray(weights));
  465. requantize->create_pipeline(opt);
  466. requantize->forward(src, dst, opt);
  467. requantize->destroy_pipeline(opt);
  468. delete requantize;
  469. }
  470. } // namespace ncnn