OLD | NEW |
(Empty) | |
| 1 |
| 2 /* |
| 3 * Copyright 2013 Google Inc. |
| 4 * |
| 5 * Use of this source code is governed by a BSD-style license that can be |
| 6 * found in the LICENSE file. |
| 7 */ |
| 8 |
| 9 #include <cstring> |
| 10 |
| 11 #include "SkBitmap.h" |
| 12 #include "SkStream.h" |
| 13 |
| 14 #include "SkCLImageDiffer.h" |
| 15 #include "skpdiff_util.h" |
| 16 |
| 17 SkCLImageDiffer::SkCLImageDiffer() { |
| 18 fIsGood = false; |
| 19 } |
| 20 |
| 21 |
| 22 bool SkCLImageDiffer::init(cl_device_id device, cl_context context) { |
| 23 fContext = context; |
| 24 fDevice = device; |
| 25 |
| 26 cl_int queueErr; |
| 27 fCommandQueue = clCreateCommandQueue(fContext, fDevice, 0, &queueErr); |
| 28 if (CL_SUCCESS != queueErr) { |
| 29 SkDebugf("Command queue creation failed: %s\n", cl_error_to_string(queue
Err)); |
| 30 fIsGood = false; |
| 31 return false; |
| 32 } |
| 33 |
| 34 fIsGood = this->onInit(); |
| 35 return fIsGood; |
| 36 } |
| 37 |
| 38 bool SkCLImageDiffer::loadKernelFile(const char file[], const char name[], cl_ke
rnel* kernel) { |
| 39 // Open the kernel source file |
| 40 SkFILEStream sourceStream(file); |
| 41 if (!sourceStream.isValid()) { |
| 42 SkDebugf("Failed to open kernel source file"); |
| 43 return false; |
| 44 } |
| 45 |
| 46 return loadKernelStream(&sourceStream, name, kernel); |
| 47 } |
| 48 |
| 49 bool SkCLImageDiffer::loadKernelStream(SkStream* stream, const char name[], cl_k
ernel* kernel) { |
| 50 // Read the kernel source into memory |
| 51 SkString sourceString; |
| 52 sourceString.resize(stream->getLength()); |
| 53 size_t bytesRead = stream->read(sourceString.writable_str(), sourceString.si
ze()); |
| 54 if (bytesRead != sourceString.size()) { |
| 55 SkDebugf("Failed to read kernel source file"); |
| 56 return false; |
| 57 } |
| 58 |
| 59 return loadKernelSource(sourceString.c_str(), name, kernel); |
| 60 } |
| 61 |
| 62 bool SkCLImageDiffer::loadKernelSource(const char source[], const char name[], c
l_kernel* kernel) { |
| 63 // Build the kernel source |
| 64 size_t sourceLen = strlen(source); |
| 65 cl_program program = clCreateProgramWithSource(fContext, 1, &source, &source
Len, NULL); |
| 66 cl_int programErr = clBuildProgram(program, 1, &fDevice, "", NULL, NULL); |
| 67 if (CL_SUCCESS != programErr) { |
| 68 SkDebugf("Program creation failed: %s\n", cl_error_to_string(programErr)
); |
| 69 |
| 70 // Attempt to get information about why the build failed |
| 71 char buildLog[4096]; |
| 72 clGetProgramBuildInfo(program, fDevice, CL_PROGRAM_BUILD_LOG, sizeof(bui
ldLog), buildLog, NULL); |
| 73 SkDebugf("Build log: %s\n", buildLog); |
| 74 |
| 75 return false; |
| 76 } |
| 77 |
| 78 cl_int kernelErr; |
| 79 *kernel = clCreateKernel(program, name, &kernelErr); |
| 80 if (CL_SUCCESS != kernelErr) { |
| 81 SkDebugf("Kernel creation failed: %s\n", cl_error_to_string(kernelErr)); |
| 82 return false; |
| 83 } |
| 84 |
| 85 return true; |
| 86 } |
| 87 |
| 88 bool SkCLImageDiffer::makeImage2D(SkBitmap* bitmap, cl_mem* image) { |
| 89 cl_int imageErr; |
| 90 cl_image_format bitmapFormat; |
| 91 switch (bitmap->config()) { |
| 92 case SkBitmap::kA8_Config: |
| 93 bitmapFormat.image_channel_order = CL_A; |
| 94 bitmapFormat.image_channel_data_type = CL_UNSIGNED_INT8; |
| 95 break; |
| 96 case SkBitmap::kRGB_565_Config: |
| 97 bitmapFormat.image_channel_order = CL_RGB; |
| 98 bitmapFormat.image_channel_data_type = CL_UNORM_SHORT_565; |
| 99 break; |
| 100 case SkBitmap::kARGB_8888_Config: |
| 101 bitmapFormat.image_channel_order = CL_RGBA; |
| 102 bitmapFormat.image_channel_data_type = CL_UNSIGNED_INT8; |
| 103 break; |
| 104 default: |
| 105 SkDebugf("Image format is unsupported\n"); |
| 106 return false; |
| 107 } |
| 108 |
| 109 // Upload the bitmap data to OpenCL |
| 110 bitmap->lockPixels(); |
| 111 *image = clCreateImage2D(fContext, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, |
| 112 &bitmapFormat, bitmap->width(), bitmap->height(), |
| 113 bitmap->rowBytes(), bitmap->getPixels(), |
| 114 &imageErr); |
| 115 bitmap->unlockPixels(); |
| 116 |
| 117 if (CL_SUCCESS != imageErr) { |
| 118 SkDebugf("Input image creation failed: %s\n", cl_error_to_string(imageEr
r)); |
| 119 return false; |
| 120 } |
| 121 |
| 122 return true; |
| 123 } |
| 124 |
| 125 |
| 126 //////////////////////////////////////////////////////////////// |
| 127 |
| 128 const char* SkDifferentPixelsImageDiffer::getName() { |
| 129 return "Find Different Pixels"; |
| 130 } |
| 131 |
| 132 int SkDifferentPixelsImageDiffer::queueDiff(SkBitmap * baseline, SkBitmap * test
) { |
| 133 int diffID = fQueuedDiffs.count(); |
| 134 double startTime = get_seconds(); |
| 135 QueuedDiff* diff = fQueuedDiffs.push(); |
| 136 |
| 137 // Ensure the images are comparable |
| 138 if (baseline->width() != test->width() || baseline->height() != test->height
() || |
| 139 baseline->width() <= 0 || baseline->height() <= 0) { |
| 140 diff->finished = true; |
| 141 diff->result = 0.0; |
| 142 return diffID; |
| 143 } |
| 144 |
| 145 // Upload images to the CL device |
| 146 if (!this->makeImage2D(baseline, &diff->baseline) || !this->makeImage2D(test
, &diff->test)) { |
| 147 diff->finished = true; |
| 148 diff->result = 0.0; |
| 149 fIsGood = false; |
| 150 return -1; |
| 151 } |
| 152 |
| 153 // A small hack that makes calculating percentage difference easier later on
. |
| 154 diff->result = 1.0 / ((double)baseline->width() * baseline->height()); |
| 155 |
| 156 // Make a buffer to store results into |
| 157 int numDiffPixels = 0; |
| 158 diff->resultsBuffer = clCreateBuffer(fContext, CL_MEM_READ_WRITE | CL_MEM_CO
PY_HOST_PTR, |
| 159 sizeof(int), &numDiffPixels, NULL); |
| 160 |
| 161 // Set all kernel arguments |
| 162 cl_int setArgErr = clSetKernelArg(fKernel, 0, sizeof(cl_mem), &diff->baselin
e); |
| 163 setArgErr |= clSetKernelArg(fKernel, 1, sizeof(cl_mem), &diff->test); |
| 164 setArgErr |= clSetKernelArg(fKernel, 2, sizeof(cl_mem), &diff->results
Buffer); |
| 165 if (CL_SUCCESS != setArgErr) { |
| 166 SkDebugf("Set arg failed: %s\n", cl_error_to_string(setArgErr)); |
| 167 fIsGood = false; |
| 168 return -1; |
| 169 } |
| 170 |
| 171 // Queue this diff on the CL device |
| 172 cl_event event; |
| 173 const size_t workSize[] = { baseline->width(), baseline->height() }; |
| 174 cl_int enqueueErr; |
| 175 enqueueErr = clEnqueueNDRangeKernel(fCommandQueue, fKernel, 2, NULL, workSiz
e, NULL, 0, NULL, &event); |
| 176 if (CL_SUCCESS != enqueueErr) { |
| 177 SkDebugf("Enqueue failed: %s\n", cl_error_to_string(enqueueErr)); |
| 178 fIsGood = false; |
| 179 return -1; |
| 180 } |
| 181 |
| 182 // This makes things totally synchronous. Actual queue is not ready yet |
| 183 clWaitForEvents(1, &event); |
| 184 diff->finished = true; |
| 185 |
| 186 // Immediate read back the results |
| 187 clEnqueueReadBuffer(fCommandQueue, diff->resultsBuffer, CL_TRUE, 0, sizeof(i
nt), &numDiffPixels, 0, NULL, NULL); |
| 188 diff->result *= (double)numDiffPixels; |
| 189 diff->result = (1.0 - diff->result); |
| 190 SkDebugf("Time: %f\n", (get_seconds() - startTime)); |
| 191 |
| 192 return diffID; |
| 193 } |
| 194 |
| 195 bool SkDifferentPixelsImageDiffer::isFinished(int id) { |
| 196 return fQueuedDiffs[id].finished; |
| 197 } |
| 198 |
| 199 double SkDifferentPixelsImageDiffer::getResult(int id) { |
| 200 return fQueuedDiffs[id].result; |
| 201 } |
| 202 |
| 203 |
| 204 bool SkDifferentPixelsImageDiffer::onInit() { |
| 205 if (!loadKernelFile("experimental/skpdiff/diff_pixels.cl", "diff", &fKernel)
) { |
| 206 return false; |
| 207 } |
| 208 |
| 209 return true; |
| 210 } |
OLD | NEW |