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