1 /*
2 * Copyright (c) 2022-2023 Huawei Device Co., Ltd.
3 * Licensed under the Apache License, Version 2.0 (the "License");
4 * you may not use this file except in compliance with the License.
5 * You may obtain a copy of the License at
6 *
7 * http://www.apache.org/licenses/LICENSE-2.0
8 *
9 * Unless required by applicable law or agreed to in writing, software
10 * distributed under the License is distributed on an "AS IS" BASIS,
11 * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
12 * See the License for the specific language governing permissions and
13 * limitations under the License.
14 */
15
16 #include <cmath>
17 #include <fstream>
18 #ifdef ENABLE_OPENCL
19 #include <securec.h>
20 #endif // ENABLE_OPENCL
21
22 #ifdef USE_ROSEN_DRAWING
23 #include "include/core/SkData.h"
24 #include "image/bitmap.h"
25 #include "utils/data.h"
26 #endif
27
28 #include "base/log/ace_trace.h"
29 #include "base/log/log.h"
30 #include "base/thread/background_task_executor.h"
31 #include "base/utils/system_properties.h"
32 #include "core/image/image_cache.h"
33 #include "core/image/image_compressor.h"
34
35 namespace OHOS::Ace {
load(void)36 __attribute__((constructor)) void load(void)
37 {
38 #ifdef ENABLE_OPENCL
39 #ifdef __MUSL__
40 OHOS::InitOpenCL();
41 #endif
42 #endif
43 }
44
45 std::shared_ptr<ImageCompressor> ImageCompressor::instance_ = nullptr;
46 std::mutex ImageCompressor::instanceMutex_;
GetInstance()47 std::shared_ptr<ImageCompressor> ImageCompressor::GetInstance()
48 {
49 if (instance_ == nullptr) {
50 std::lock_guard<std::mutex> lock(instanceMutex_);
51 if (instance_ == nullptr) {
52 instance_.reset(new ImageCompressor());
53 instance_->Init();
54 }
55 }
56 return instance_;
57 }
58
Init()59 void ImageCompressor::Init()
60 {
61 #ifdef ENABLE_OPENCL
62 switch_ = SystemProperties::IsAstcEnabled();
63 if (switch_) {
64 clOk_ = OHOS::InitOpenCL();
65 maxErr_ = SystemProperties::GetAstcMaxError();
66 psnr_ = SystemProperties::GetAstcPsnr();
67 InitPartition();
68 InitRecords();
69 }
70 #endif // ENABLE_OPENCL
71 }
72
CanCompress()73 bool ImageCompressor::CanCompress()
74 {
75 #ifdef UPLOAD_GPU_DISABLED
76 return false;
77 #else
78 if (switch_ && clOk_) {
79 return true;
80 }
81 return false;
82 #endif
83 }
84
85 #ifdef ENABLE_OPENCL
LoadShaderBin(cl_context context,cl_device_id device_id)86 cl_program ImageCompressor::LoadShaderBin(cl_context context, cl_device_id device_id)
87 {
88 ACE_FUNCTION_TRACE();
89 std::unique_ptr<FILE, decltype(&fclose)> file(fopen(shader_path_.c_str(), "rb"), fclose);
90 if (!file) {
91 LOGE("load cl shader failed");
92 return nullptr;
93 }
94 #ifndef USE_ROSEN_DRAWING
95 auto data = SkData::MakeFromFILE(file.get());
96 #else
97 auto data = SkData::MakeFromFILE(file.get());
98 #endif
99 if (!data) {
100 return nullptr;
101 }
102 cl_int err;
103 size_t len = data->size();
104 auto ptr = (const unsigned char*) data->data();
105 cl_program p = clCreateProgramWithBinary(context, 1, &device_id, &len, &ptr, NULL, &err);
106 if (err) {
107 return nullptr;
108 }
109 LOGD("load cl shader");
110 return p;
111 }
112
CreateKernel()113 bool ImageCompressor::CreateKernel()
114 {
115 if (!context_ || !kernel_) {
116 cl_int err;
117 cl_platform_id platform_id;
118 cl_device_id device_id;
119 clGetPlatformIDs(1, &platform_id, NULL);
120 clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_GPU, 1, &device_id, NULL);
121 context_ = clCreateContext(0, 1, &device_id, NULL, NULL, &err);
122 queue_ = clCreateCommandQueueWithProperties(context_, device_id, 0, &err);
123
124 cl_program program = LoadShaderBin(context_, device_id);
125 clBuildProgram(program, 1, &device_id, compileOption_.c_str(), NULL, NULL);
126 ACE_SCOPED_TRACE("clCreateKernel");
127 kernel_ = clCreateKernel(program, "astc", &err);
128 clReleaseProgram(program);
129 }
130 if (!context_ || !kernel_ || !queue_) {
131 ReleaseResource();
132 LOGE("build opencl program failed");
133 clOk_ = false;
134 return false;
135 }
136 refCount_++;
137 return true;
138 }
139
CheckImageQuality(std::string key,uint32_t sumErr,uint32_t maxErr,int32_t width,int32_t height)140 bool ImageCompressor::CheckImageQuality(std::string key, uint32_t sumErr, uint32_t maxErr, int32_t width, int32_t height)
141 {
142 bool isOk = true;
143 float mse = (float)sumErr / (width * height);
144 float psnr = 10 * log10(255 * 255 / mse);
145 if (maxErr == 0 || psnr == 0 || maxErr > maxErr_ || (int32_t)psnr < psnr_) {
146 isOk = false;
147 std::lock_guard<std::mutex> mLock(recordsMutex_);
148 failedRecords_.insert(key);
149 }
150 LOGI("compress quality %{public}s [%{public}u, %{public}.2f] size(%{public}d×%{public}d) %{public}s",
151 key.c_str(), maxErr, psnr, width, height, isOk ? "ok" : "no");
152 return isOk;
153 }
154
ReleaseResource()155 void ImageCompressor::ReleaseResource()
156 {
157 ACE_FUNCTION_TRACE();
158 clReleaseKernel(kernel_);
159 kernel_ = NULL;
160 clReleaseCommandQueue(queue_);
161 queue_ = NULL;
162 clReleaseContext(context_);
163 context_ = NULL;
164 }
165 #endif // ENABLE_OPENCL
166
167 #ifndef USE_ROSEN_DRAWING
GpuCompress(std::string key,SkPixmap & pixmap,int32_t width,int32_t height)168 sk_sp<SkData> ImageCompressor::GpuCompress(std::string key, SkPixmap& pixmap, int32_t width, int32_t height)
169 #else
170 std::shared_ptr<RSData> ImageCompressor::GpuCompress(std::string key, RSBitmap& bitmap, int32_t width, int32_t height)
171 #endif
172 {
173 #ifdef ENABLE_OPENCL
174 std::lock_guard<std::mutex> lock(instanceMutex_);
175 if (width <= 0 || height <= 0 || !clOk_ || IsFailedImage(key) || width > maxSize_ || height > maxSize_) {
176 return nullptr;
177 }
178 if (!CreateKernel()) {
179 return nullptr;
180 }
181 ACE_SCOPED_TRACE("GpuCompress %d×%d", width, height);
182
183 cl_int err;
184
185 // Number of work items in each local work group
186 int32_t blockX = ceil((width + DIM - 1) / DIM);
187 int32_t blockY = ceil((height + DIM - 1) / DIM);
188 int32_t numBlocks = blockX * blockY;
189 size_t local[] = { DIM, DIM };
190 size_t global[2];
191 global[0] = (width % local[0] == 0 ? width : (width + local[0] - width % local[0]));
192 global[1] = (height % local[1] == 0 ? height : (height + local[1] - height % local[1]));
193
194 size_t astc_size = numBlocks * DIM * DIM;
195
196 cl_image_format image_format = { CL_RGBA, CL_UNORM_INT8 };
197 cl_image_desc desc = { CL_MEM_OBJECT_IMAGE2D, width, height };
198 cl_mem inputImage = clCreateImage(context_, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
199 #ifndef USE_ROSEN_DRAWING
200 &image_format, &desc, const_cast<void*>(pixmap.addr()), &err);
201 #else
202 &image_format, &desc, bitmap.GetPixels(), &err);
203 #endif
204 cl_mem astcResult = clCreateBuffer(context_, CL_MEM_ALLOC_HOST_PTR, astc_size, NULL, &err);
205 cl_mem partInfos = clCreateBuffer(context_, CL_MEM_COPY_HOST_PTR,
206 sizeof(PartInfo) * parts_.size(), &parts_[0], &err);
207
208 uint32_t* blockErrs = new uint32_t[numBlocks]{0};
209 cl_mem clErrs = clCreateBuffer(context_, CL_MEM_USE_HOST_PTR, sizeof(uint32_t) * numBlocks, blockErrs, &err);
210 err |= clSetKernelArg(kernel_, 0, sizeof(cl_mem), &inputImage);
211 err |= clSetKernelArg(kernel_, 1, sizeof(cl_mem), &astcResult);
212 err |= clSetKernelArg(kernel_, 2, sizeof(cl_mem), &partInfos);
213 err |= clSetKernelArg(kernel_, 3, sizeof(cl_mem), &clErrs);
214
215 err = clEnqueueNDRangeKernel(queue_, kernel_, 2, NULL, global, local, 0, NULL, NULL);
216
217 clFinish(queue_);
218
219 uint32_t max_val = 0, sum_val = 0;
220 err = clEnqueueReadBuffer(queue_, clErrs, CL_TRUE, 0, sizeof(uint32_t) * numBlocks, blockErrs, 0, NULL, NULL);
221 for (int32_t i = 0; i < numBlocks; i++) {
222 sum_val += blockErrs[i];
223 max_val = fmax(max_val, blockErrs[i]);
224 }
225
226 clReleaseMemObject(inputImage);
227 clReleaseMemObject(partInfos);
228 clReleaseMemObject(clErrs);
229 delete[] blockErrs;
230
231 if (!CheckImageQuality(key, sum_val, max_val, width, height)) {
232 clReleaseMemObject(astcResult);
233 return nullptr;
234 }
235
236 #ifndef USE_ROSEN_DRAWING
237 auto astc_data = SkData::MakeUninitialized(astc_size);
238 clEnqueueReadBuffer(queue_, astcResult, CL_TRUE, 0, astc_size, astc_data->writable_data(), 0, NULL, NULL);
239 #else
240 auto astc_data = std::make_shared<RSData>();
241 astc_data->BuildUninitialized(astc_size);
242 clEnqueueReadBuffer(queue_, astcResult, CL_TRUE, 0, astc_size, astc_data->WritableData(), 0, NULL, NULL);
243 #endif
244 clReleaseMemObject(astcResult);
245 return astc_data;
246 #else
247 return nullptr;
248 #endif // ENABLE_OPENCL
249 }
250
251
ScheduleReleaseTask()252 std::function<void()> ImageCompressor::ScheduleReleaseTask()
253 {
254 std::function<void()> task = [this]() {
255 #ifdef ENABLE_OPENCL
256 if (refCount_ > 0 && clOk_) {
257 refCount_--;
258 if (refCount_ <= 0) {
259 this->ReleaseResource();
260
261 // save failed records
262 std::ofstream saveFile(recordsPath_);
263 if (!saveFile.is_open()) {
264 return;
265 }
266 std::lock_guard<std::mutex> mLock(recordsMutex_);
267 for (auto s : failedRecords_) {
268 saveFile << s << "\n";
269 }
270 saveFile.close();
271 }
272 }
273 #endif // ENABLE_OPENCL
274 };
275 return task;
276 }
277
278 #ifndef USE_ROSEN_DRAWING
WriteToFile(std::string srcKey,sk_sp<SkData> compressedData,Size imgSize)279 void ImageCompressor::WriteToFile(std::string srcKey, sk_sp<SkData> compressedData, Size imgSize)
280 #else
281 void ImageCompressor::WriteToFile(std::string srcKey, std::shared_ptr<RSData> compressedData, Size imgSize)
282 #endif
283 {
284 if (!compressedData || srcKey.empty()) {
285 return;
286 }
287 #ifdef ENABLE_OPENCL
288 BackgroundTaskExecutor::GetInstance().PostTask(
289 [srcKey, compressedData, imgSize]() {
290 AstcHeader header;
291 int32_t xsize = imgSize.Width();
292 int32_t ysize = imgSize.Height();
293 header.magic[0] = MAGIC_FILE_CONSTANT & 0xFF;
294 header.magic[1] = (MAGIC_FILE_CONSTANT >> 8) & 0xFF;
295 header.magic[2] = (MAGIC_FILE_CONSTANT >> 16) & 0xFF;
296 header.magic[3] = (MAGIC_FILE_CONSTANT >> 24) & 0xFF;
297 header.blockdimX = DIM;
298 header.blockdimY = DIM;
299 header.blockdimZ = 1;
300 header.xsize[0] = xsize & 0xFF;
301 header.xsize[1] = (xsize >> 8) & 0xFF;
302 header.xsize[2] = (xsize >> 16) & 0xFF;
303 header.ysize[0] = ysize & 0xFF;
304 header.ysize[1] = (ysize >> 8) & 0xFF;
305 header.ysize[2] = (ysize >> 16) & 0xFF;
306 header.zsize[0] = 1;
307 header.zsize[1] = 0;
308 header.zsize[2] = 0;
309 LOGD("astc write file %{public}s size(%{public}d×%{public}d) (%{public}.2f×%{public}.2f)",
310 srcKey.c_str(), xsize, ysize, imgSize.Width(), imgSize.Height());
311
312 #ifndef USE_ROSEN_DRAWING
313 int32_t fileSize = compressedData->size() + sizeof(header);
314 sk_sp<SkData> toWrite = SkData::MakeUninitialized(fileSize);
315 uint8_t* toWritePtr = (uint8_t*) toWrite->writable_data();
316 #else
317 int32_t fileSize = compressedData->GetSize() + sizeof(header);
318 auto toWrite = std::shared_ptr<RSData>();
319 toWrite->BuildUninitialized(fileSize);
320 uint8_t* toWritePtr = (uint8_t*) toWrite->WritableData();
321 #endif
322 if (memcpy_s(toWritePtr, fileSize, &header, sizeof(header)) != EOK) {
323 LOGE("astc write file failed");
324 return;
325 }
326 #ifndef USE_ROSEN_DRAWING
327 if (memcpy_s(toWritePtr + sizeof(header), compressedData->size(),
328 compressedData->data(), compressedData->size()) != EOK) {
329 #else
330 if (memcpy_s(toWritePtr + sizeof(header), compressedData->GetSize(),
331 compressedData->GetData(), compressedData->GetSize()) != EOK) {
332 #endif
333 LOGE("astc write file failed");
334 return;
335 }
336
337 ImageCache::WriteCacheFile(srcKey, toWritePtr, fileSize, ".astc");
338 }, BgTaskPriority::LOW);
339 #endif
340 }
341
342 #ifndef USE_ROSEN_DRAWING
343 sk_sp<SkData> ImageCompressor::StripFileHeader(sk_sp<SkData> fileData)
344 {
345 if (fileData) {
346 auto imageData = SkData::MakeSubset(fileData.get(), sizeof(AstcHeader), fileData->size() - sizeof(AstcHeader));
347 if (!imageData->isEmpty()) {
348 return imageData;
349 }
350 }
351 return nullptr;
352 }
353 #else
354 std::shared_ptr<RSData> ImageCompressor::StripFileHeader(std::shared_ptr<RSData> fileData)
355 {
356 if (fileData) {
357 if (fileData->GetSize() <= sizeof(AstcHeader)) {
358 return nullptr;
359 }
360 auto imageData = std::make_shared<RSData>();
361 if (imageData->BuildWithCopy(static_cast<const uint8_t*>(fileData->GetData()) + sizeof(AstcHeader),
362 fileData->GetSize() - sizeof(AstcHeader))) {
363 return imageData;
364 }
365 }
366 return nullptr;
367 }
368 #endif
369
370 /**
371 * @brief Hash function used for procedural partition assignment.
372 *
373 * @param seed The hash seed.
374 *
375 * @return The hashed value.
376 */
377 static uint32_t Hash52(uint32_t seed)
378 {
379 seed ^= seed >> 15;
380
381 // (2^4 + 1) * (2^7 + 1) * (2^17 - 1)
382 seed *= 0xEEDE0891;
383 seed ^= seed >> 5;
384 seed += seed << 16;
385 seed ^= seed >> 7;
386 seed ^= seed >> 3;
387 seed ^= seed << 6;
388 seed ^= seed >> 17;
389 return seed;
390 }
391
392 /**
393 * @brief Select texel assignment for a single coordinate.
394 *
395 * @param seed The seed - the partition index from the block.
396 * @param x The texel X coordinate in the block.
397 * @param y The texel Y coordinate in the block.
398 * @param z The texel Z coordinate in the block.
399 * @param partitionCount The total partition count of this encoding.
400 * @param smallBlock @c true if the blockhas fewer than 32 texels.
401 *
402 * @return The assigned partition index for this texel.
403 */
404 static uint8_t SelectPartition(int32_t seed, int32_t x, int32_t y, int32_t z, int32_t partitionCount, bool smallBlock)
405 {
406 // For small blocks bias the coordinates to get better distribution
407 if (smallBlock) {
408 x <<= 1;
409 y <<= 1;
410 z <<= 1;
411 }
412
413 seed += (partitionCount - 1) * 1024;
414
415 uint32_t num = Hash52(seed);
416
417 uint8_t seed1 = num & 0xF;
418 uint8_t seed2 = (num >> 4) & 0xF;
419 uint8_t seed3 = (num >> 8) & 0xF;
420 uint8_t seed4 = (num >> 12) & 0xF;
421 uint8_t seed5 = (num >> 16) & 0xF;
422 uint8_t seed6 = (num >> 20) & 0xF;
423 uint8_t seed7 = (num >> 24) & 0xF;
424 uint8_t seed8 = (num >> 28) & 0xF;
425 uint8_t seed9 = (num >> 18) & 0xF;
426 uint8_t seed10 = (num >> 22) & 0xF;
427 uint8_t seed11 = (num >> 26) & 0xF;
428 uint8_t seed12 = ((num >> 30) | (num << 2)) & 0xF;
429
430 // Squaring all the seeds in order to bias their distribution towards lower values.
431 seed1 *= seed1;
432 seed2 *= seed2;
433 seed3 *= seed3;
434 seed4 *= seed4;
435 seed5 *= seed5;
436 seed6 *= seed6;
437 seed7 *= seed7;
438 seed8 *= seed8;
439 seed9 *= seed9;
440 seed10 *= seed10;
441 seed11 *= seed11;
442 seed12 *= seed12;
443
444 int32_t sh1, sh2;
445 if (seed & 1) {
446 sh1 = (seed & 2 ? 4 : 5);
447 sh2 = (partitionCount == 3 ? 6 : 5);
448 } else {
449 sh1 = (partitionCount == 3 ? 6 : 5);
450 sh2 = (seed & 2 ? 4 : 5);
451 }
452
453 int32_t sh3 = (seed & 0x10) ? sh1 : sh2;
454
455 seed1 >>= sh1;
456 seed2 >>= sh2;
457 seed3 >>= sh1;
458 seed4 >>= sh2;
459 seed5 >>= sh1;
460 seed6 >>= sh2;
461 seed7 >>= sh1;
462 seed8 >>= sh2;
463
464 seed9 >>= sh3;
465 seed10 >>= sh3;
466 seed11 >>= sh3;
467 seed12 >>= sh3;
468
469 int32_t a = seed1 * x + seed2 * y + seed11 * z + (num >> 14);
470 int32_t b = seed3 * x + seed4 * y + seed12 * z + (num >> 10);
471 int32_t c = seed5 * x + seed6 * y + seed9 * z + (num >> 6);
472 int32_t d = seed7 * x + seed8 * y + seed10 * z + (num >> 2);
473
474 // Apply the saw
475 a &= 0x3F;
476 b &= 0x3F;
477 c &= 0x3F;
478 d &= 0x3F;
479
480 // Remove some of the components if we are to output < 4 partitions_.
481 if (partitionCount <= 3) {
482 d = 0;
483 }
484
485 if (partitionCount <= 2) {
486 c = 0;
487 }
488
489 if (partitionCount <= 1) {
490 b = 0;
491 }
492
493 uint8_t partition;
494 if (a >= b && a >= c && a >= d) {
495 partition = 0;
496 } else if (b >= c && b >= d) {
497 partition = 1;
498 } else if (c >= d) {
499 partition = 2;
500 } else {
501 partition = 3;
502 }
503
504 return partition;
505 }
506
507 bool ImageCompressor::InitPartitionInfo(PartInfo *partInfos, int32_t part_index, int32_t part_count)
508 {
509 int32_t texIdx = 0;
510 int32_t counts[4] = {0};
511 for (int32_t y = 0; y < DIM; y++) {
512 for (int32_t x = 0; x < DIM; x++) {
513 int32_t part = SelectPartition(part_index, x, y, 0, part_count, true);
514 partInfos->bitmaps[part] |= 1u << texIdx;
515 counts[part]++;
516 texIdx++;
517 }
518 }
519 int32_t realPartCount = 0;
520 if (counts[0] == 0) {
521 realPartCount = 0;
522 } else if (counts[1] == 0) {
523 realPartCount = 1;
524 } else if (counts[2] == 0) {
525 realPartCount = 2;
526 } else if (counts[3] == 0) {
527 realPartCount = 3;
528 } else {
529 realPartCount = 4;
530 }
531 if (realPartCount == part_count) {
532 return true;
533 }
534 return false;
535 }
536
537 void ImageCompressor::InitPartition()
538 {
539 parts_.clear();
540 int32_t arrSize = sizeof(partitions_) / sizeof(partitions_[0]);
541 for (int32_t i = 0; i < arrSize; i++) {
542 PartInfo p = {};
543 if (InitPartitionInfo(&p, partitions_[i], 2)) {
544 p.partid = partitions_[i];
545 parts_.push_back(p);
546 LOGD("part id:%d %d %d", p.partid, p.bitmaps[0], p.bitmaps[1]);
547 }
548 }
549 compileOption_ = "-D PARTITION_SERACH_MAX=" + std::to_string(parts_.size());
550 }
551
552 #ifdef ENABLE_OPENCL
553 bool ImageCompressor::IsFailedImage(std::string key)
554 {
555 std::lock_guard<std::mutex> mLock(recordsMutex_);
556 return failedRecords_.find(key) != failedRecords_.end();
557 }
558 #endif
559
560 void ImageCompressor::InitRecords()
561 {
562 recordsPath_ = ImageCache::GetImageCacheFilePath("record") + ".txt";
563 std::ifstream openFile(recordsPath_);
564 if (!openFile.is_open()) {
565 return;
566 }
567 std::string line;
568 std::lock_guard<std::mutex> mLock(recordsMutex_);
569 while (!openFile.eof()) {
570 std::getline(openFile, line);
571 failedRecords_.insert(line);
572 }
573 openFile.close();
574 }
575 #ifdef FUZZTEST
576 void ImageCompressor::PartDoing()
577 {
578 InitPartition();
579 InitRecords();
580 }
581 #endif
582 } // namespace OHOS::Ace
583