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_file_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 return p;
110 }
111
CreateKernel()112 bool ImageCompressor::CreateKernel()
113 {
114 if (!context_ || !kernel_) {
115 cl_int err;
116 cl_platform_id platform_id;
117 cl_device_id device_id;
118 clGetPlatformIDs(1, &platform_id, NULL);
119 clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_GPU, 1, &device_id, NULL);
120 context_ = clCreateContext(0, 1, &device_id, NULL, NULL, &err);
121 queue_ = clCreateCommandQueueWithProperties(context_, device_id, 0, &err);
122
123 cl_program program = LoadShaderBin(context_, device_id);
124 clBuildProgram(program, 1, &device_id, compileOption_.c_str(), NULL, NULL);
125 ACE_SCOPED_TRACE("clCreateKernel");
126 kernel_ = clCreateKernel(program, "astc", &err);
127 clReleaseProgram(program);
128 }
129 if (!context_ || !kernel_ || !queue_) {
130 ReleaseResource();
131 LOGE("build opencl program failed");
132 clOk_ = false;
133 return false;
134 }
135 refCount_++;
136 return true;
137 }
138
CheckImageQuality(std::string key,uint32_t sumErr,uint32_t maxErr,int32_t width,int32_t height)139 bool ImageCompressor::CheckImageQuality(std::string key, uint32_t sumErr, uint32_t maxErr, int32_t width, int32_t height)
140 {
141 bool isOk = true;
142 float mse = (float)sumErr / (width * height);
143 float psnr = 10 * log10(255 * 255 / mse);
144 if (maxErr == 0 || psnr == 0 || maxErr > maxErr_ || (int32_t)psnr < psnr_) {
145 isOk = false;
146 std::lock_guard<std::mutex> mLock(recordsMutex_);
147 failedRecords_.insert(key);
148 }
149 LOGI("compress quality %{public}s [%{public}u, %{public}.2f] size(%{public}d×%{public}d) %{public}s",
150 key.c_str(), maxErr, psnr, width, height, isOk ? "ok" : "no");
151 return isOk;
152 }
153
ReleaseResource()154 void ImageCompressor::ReleaseResource()
155 {
156 ACE_FUNCTION_TRACE();
157 clReleaseKernel(kernel_);
158 kernel_ = NULL;
159 clReleaseCommandQueue(queue_);
160 queue_ = NULL;
161 clReleaseContext(context_);
162 context_ = NULL;
163 }
164 #endif // ENABLE_OPENCL
165
166 #ifndef USE_ROSEN_DRAWING
GpuCompress(std::string key,SkPixmap & pixmap,int32_t width,int32_t height)167 sk_sp<SkData> ImageCompressor::GpuCompress(std::string key, SkPixmap& pixmap, int32_t width, int32_t height)
168 #else
169 std::shared_ptr<RSData> ImageCompressor::GpuCompress(std::string key, RSBitmap& bitmap, int32_t width, int32_t height)
170 #endif
171 {
172 #ifdef ENABLE_OPENCL
173 std::lock_guard<std::mutex> lock(instanceMutex_);
174 if (width <= 0 || height <= 0 || !clOk_ || IsFailedImage(key) || width > maxSize_ || height > maxSize_) {
175 return nullptr;
176 }
177 if (!CreateKernel()) {
178 return nullptr;
179 }
180 ACE_SCOPED_TRACE("GpuCompress %d×%d", width, height);
181
182 cl_int err;
183
184 // Number of work items in each local work group
185 int32_t blockX = ceil((width + DIM - 1) / DIM);
186 int32_t blockY = ceil((height + DIM - 1) / DIM);
187 int32_t numBlocks = blockX * blockY;
188 size_t local[] = { DIM, DIM };
189 size_t global[2];
190 global[0] = (width % local[0] == 0 ? width : (width + local[0] - width % local[0]));
191 global[1] = (height % local[1] == 0 ? height : (height + local[1] - height % local[1]));
192
193 size_t astc_size = numBlocks * DIM * DIM;
194
195 cl_image_format image_format = { CL_RGBA, CL_UNORM_INT8 };
196 cl_image_desc desc = { CL_MEM_OBJECT_IMAGE2D, width, height };
197 cl_mem inputImage = clCreateImage(context_, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
198 #ifndef USE_ROSEN_DRAWING
199 &image_format, &desc, const_cast<void*>(pixmap.addr()), &err);
200 #else
201 &image_format, &desc, bitmap.GetPixels(), &err);
202 #endif
203 cl_mem astcResult = clCreateBuffer(context_, CL_MEM_ALLOC_HOST_PTR, astc_size, NULL, &err);
204 cl_mem partInfos = clCreateBuffer(context_, CL_MEM_COPY_HOST_PTR,
205 sizeof(PartInfo) * parts_.size(), &parts_[0], &err);
206
207 uint32_t* blockErrs = new uint32_t[numBlocks]{0};
208 cl_mem clErrs = clCreateBuffer(context_, CL_MEM_USE_HOST_PTR, sizeof(uint32_t) * numBlocks, blockErrs, &err);
209 err |= clSetKernelArg(kernel_, 0, sizeof(cl_mem), &inputImage);
210 err |= clSetKernelArg(kernel_, 1, sizeof(cl_mem), &astcResult);
211 err |= clSetKernelArg(kernel_, 2, sizeof(cl_mem), &partInfos);
212 err |= clSetKernelArg(kernel_, 3, sizeof(cl_mem), &clErrs);
213
214 err = clEnqueueNDRangeKernel(queue_, kernel_, 2, NULL, global, local, 0, NULL, NULL);
215
216 clFinish(queue_);
217
218 uint32_t max_val = 0, sum_val = 0;
219 err = clEnqueueReadBuffer(queue_, clErrs, CL_TRUE, 0, sizeof(uint32_t) * numBlocks, blockErrs, 0, NULL, NULL);
220 for (int32_t i = 0; i < numBlocks; i++) {
221 sum_val += blockErrs[i];
222 max_val = fmax(max_val, blockErrs[i]);
223 }
224
225 clReleaseMemObject(inputImage);
226 clReleaseMemObject(partInfos);
227 clReleaseMemObject(clErrs);
228 delete[] blockErrs;
229
230 if (!CheckImageQuality(key, sum_val, max_val, width, height)) {
231 clReleaseMemObject(astcResult);
232 return nullptr;
233 }
234
235 #ifndef USE_ROSEN_DRAWING
236 auto astc_data = SkData::MakeUninitialized(astc_size);
237 clEnqueueReadBuffer(queue_, astcResult, CL_TRUE, 0, astc_size, astc_data->writable_data(), 0, NULL, NULL);
238 #else
239 auto astc_data = std::make_shared<RSData>();
240 astc_data->BuildUninitialized(astc_size);
241 clEnqueueReadBuffer(queue_, astcResult, CL_TRUE, 0, astc_size, astc_data->WritableData(), 0, NULL, NULL);
242 #endif
243 clReleaseMemObject(astcResult);
244 return astc_data;
245 #else
246 return nullptr;
247 #endif // ENABLE_OPENCL
248 }
249
250
ScheduleReleaseTask()251 std::function<void()> ImageCompressor::ScheduleReleaseTask()
252 {
253 #ifdef ENABLE_OPENCL
254 std::function<void()> task = [this]() {
255 if (refCount_ > 0 && clOk_) {
256 refCount_--;
257 if (refCount_ <= 0) {
258 this->ReleaseResource();
259
260 // save failed records
261 std::ofstream saveFile(recordsPath_);
262 if (!saveFile.is_open()) {
263 return;
264 }
265 std::lock_guard<std::mutex> mLock(recordsMutex_);
266 for (auto s : failedRecords_) {
267 saveFile << s << "\n";
268 }
269 saveFile.close();
270 }
271 }
272 #else
273 std::function<void()> task = []() {
274 #endif // ENABLE_OPENCL
275 };
276
277 return task;
278 }
279
280 #ifndef USE_ROSEN_DRAWING
281 void ImageCompressor::WriteToFile(std::string srcKey, sk_sp<SkData> compressedData, Size imgSize)
282 #else
283 void ImageCompressor::WriteToFile(std::string srcKey, std::shared_ptr<RSData> compressedData, Size imgSize)
284 #endif
285 {
286 if (!compressedData || srcKey.empty()) {
287 return;
288 }
289 #ifdef ENABLE_OPENCL
290 BackgroundTaskExecutor::GetInstance().PostTask(
291 [srcKey, compressedData, imgSize]() {
292 AstcHeader header;
293 int32_t xsize = imgSize.Width();
294 int32_t ysize = imgSize.Height();
295 header.magic[0] = MAGIC_FILE_CONSTANT & 0xFF;
296 header.magic[1] = (MAGIC_FILE_CONSTANT >> 8) & 0xFF;
297 header.magic[2] = (MAGIC_FILE_CONSTANT >> 16) & 0xFF;
298 header.magic[3] = (MAGIC_FILE_CONSTANT >> 24) & 0xFF;
299 header.blockdimX = DIM;
300 header.blockdimY = DIM;
301 header.blockdimZ = 1;
302 header.xsize[0] = xsize & 0xFF;
303 header.xsize[1] = (xsize >> 8) & 0xFF;
304 header.xsize[2] = (xsize >> 16) & 0xFF;
305 header.ysize[0] = ysize & 0xFF;
306 header.ysize[1] = (ysize >> 8) & 0xFF;
307 header.ysize[2] = (ysize >> 16) & 0xFF;
308 header.zsize[0] = 1;
309 header.zsize[1] = 0;
310 header.zsize[2] = 0;
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 ImageFileCache::GetInstance().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 }
547 }
548 compileOption_ = "-D PARTITION_SERACH_MAX=" + std::to_string(parts_.size());
549 }
550
551 #ifdef ENABLE_OPENCL
552 bool ImageCompressor::IsFailedImage(std::string key)
553 {
554 std::lock_guard<std::mutex> mLock(recordsMutex_);
555 return failedRecords_.find(key) != failedRecords_.end();
556 }
557 #endif
558
559 void ImageCompressor::InitRecords()
560 {
561 recordsPath_ = ImageFileCache::GetInstance().GetImageCacheFilePath("record") + ".txt";
562 std::ifstream openFile(recordsPath_);
563 if (!openFile.is_open()) {
564 return;
565 }
566 std::string line;
567 std::lock_guard<std::mutex> mLock(recordsMutex_);
568 while (!openFile.eof()) {
569 std::getline(openFile, line);
570 failedRecords_.insert(line);
571 }
572 openFile.close();
573 }
574 #ifdef FUZZTEST
575 void ImageCompressor::PartDoing()
576 {
577 InitPartition();
578 InitRecords();
579 }
580 #endif
581 } // namespace OHOS::Ace
582