Chromium Code Reviews| OLD | NEW |
|---|---|
| 1 | 1 |
| 2 /* | 2 /* |
| 3 * Copyright 2013 Google Inc. | 3 * Copyright 2013 Google Inc. |
| 4 * | 4 * |
| 5 * Use of this source code is governed by a BSD-style license that can be | 5 * Use of this source code is governed by a BSD-style license that can be |
| 6 * found in the LICENSE file. | 6 * found in the LICENSE file. |
| 7 */ | 7 */ |
| 8 | 8 |
| 9 #include <cstring> | 9 #include <cstring> |
| 10 | 10 |
| (...skipping 51 matching lines...) Expand 10 before | Expand all | Expand 10 after Loading... | |
| 62 bool SkCLImageDiffer::loadKernelSource(const char source[], const char name[], c l_kernel* kernel) { | 62 bool SkCLImageDiffer::loadKernelSource(const char source[], const char name[], c l_kernel* kernel) { |
| 63 // Build the kernel source | 63 // Build the kernel source |
| 64 size_t sourceLen = strlen(source); | 64 size_t sourceLen = strlen(source); |
| 65 cl_program program = clCreateProgramWithSource(fContext, 1, &source, &source Len, NULL); | 65 cl_program program = clCreateProgramWithSource(fContext, 1, &source, &source Len, NULL); |
| 66 cl_int programErr = clBuildProgram(program, 1, &fDevice, "", NULL, NULL); | 66 cl_int programErr = clBuildProgram(program, 1, &fDevice, "", NULL, NULL); |
| 67 if (CL_SUCCESS != programErr) { | 67 if (CL_SUCCESS != programErr) { |
| 68 SkDebugf("Program creation failed: %s\n", cl_error_to_string(programErr) ); | 68 SkDebugf("Program creation failed: %s\n", cl_error_to_string(programErr) ); |
| 69 | 69 |
| 70 // Attempt to get information about why the build failed | 70 // Attempt to get information about why the build failed |
| 71 char buildLog[4096]; | 71 char buildLog[4096]; |
| 72 clGetProgramBuildInfo(program, fDevice, CL_PROGRAM_BUILD_LOG, sizeof(bui ldLog), buildLog, NULL); | 72 clGetProgramBuildInfo(program, fDevice, CL_PROGRAM_BUILD_LOG, sizeof(bui ldLog), |
| 73 buildLog, NULL); | |
| 73 SkDebugf("Build log: %s\n", buildLog); | 74 SkDebugf("Build log: %s\n", buildLog); |
| 74 | 75 |
| 75 return false; | 76 return false; |
| 76 } | 77 } |
| 77 | 78 |
| 78 cl_int kernelErr; | 79 cl_int kernelErr; |
| 79 *kernel = clCreateKernel(program, name, &kernelErr); | 80 *kernel = clCreateKernel(program, name, &kernelErr); |
| 80 if (CL_SUCCESS != kernelErr) { | 81 if (CL_SUCCESS != kernelErr) { |
| 81 SkDebugf("Kernel creation failed: %s\n", cl_error_to_string(kernelErr)); | 82 SkDebugf("Kernel creation failed: %s\n", cl_error_to_string(kernelErr)); |
| 82 return false; | 83 return false; |
| (...skipping 35 matching lines...) Expand 10 before | Expand all | Expand 10 after Loading... | |
| 118 SkDebugf("Input image creation failed: %s\n", cl_error_to_string(imageEr r)); | 119 SkDebugf("Input image creation failed: %s\n", cl_error_to_string(imageEr r)); |
| 119 return false; | 120 return false; |
| 120 } | 121 } |
| 121 | 122 |
| 122 return true; | 123 return true; |
| 123 } | 124 } |
| 124 | 125 |
| 125 | 126 |
| 126 //////////////////////////////////////////////////////////////// | 127 //////////////////////////////////////////////////////////////// |
| 127 | 128 |
| 129 struct SkDifferentPixelsImageDiffer::QueuedDiff { | |
| 130 bool finished; | |
| 131 double result; | |
| 132 int numDiffPixels; | |
| 133 SkIPoint* poi; | |
| 134 cl_mem baseline; | |
| 135 cl_mem test; | |
| 136 cl_mem resultsBuffer; | |
| 137 cl_mem poiBuffer; | |
| 138 }; | |
| 139 | |
| 128 const char* SkDifferentPixelsImageDiffer::getName() { | 140 const char* SkDifferentPixelsImageDiffer::getName() { |
| 129 return "different_pixels"; | 141 return "different_pixels"; |
| 130 } | 142 } |
| 131 | 143 |
| 132 int SkDifferentPixelsImageDiffer::queueDiff(SkBitmap * baseline, SkBitmap * test ) { | 144 int SkDifferentPixelsImageDiffer::queueDiff(SkBitmap * baseline, SkBitmap * test ) { |
| 133 int diffID = fQueuedDiffs.count(); | 145 int diffID = fQueuedDiffs.count(); |
| 134 double startTime = get_seconds(); | 146 double startTime = get_seconds(); |
| 135 QueuedDiff* diff = fQueuedDiffs.push(); | 147 QueuedDiff* diff = fQueuedDiffs.push(); |
| 136 | 148 |
| 149 // If we never end up running the kernel, include some safe defaults in the result. | |
| 150 diff->finished = false; | |
| 151 diff->result = -1.0; | |
| 152 diff->numDiffPixels = 0; | |
| 153 diff->poi = NULL; | |
| 154 | |
| 137 // Ensure the images are comparable | 155 // Ensure the images are comparable |
| 138 if (baseline->width() != test->width() || baseline->height() != test->height () || | 156 if (baseline->width() != test->width() || baseline->height() != test->height () || |
| 139 baseline->width() <= 0 || baseline->height() <= 0) { | 157 baseline->width() <= 0 || baseline->height() <= 0) { |
| 140 diff->finished = true; | 158 diff->finished = true; |
| 141 diff->result = 0.0; | |
| 142 return diffID; | 159 return diffID; |
| 143 } | 160 } |
| 144 | 161 |
| 145 // Upload images to the CL device | 162 // Upload images to the CL device |
| 146 if (!this->makeImage2D(baseline, &diff->baseline) || !this->makeImage2D(test , &diff->test)) { | 163 if (!this->makeImage2D(baseline, &diff->baseline) || !this->makeImage2D(test , &diff->test)) { |
| 147 diff->finished = true; | 164 diff->finished = true; |
| 148 diff->result = 0.0; | |
| 149 fIsGood = false; | 165 fIsGood = false; |
| 150 return -1; | 166 return -1; |
| 151 } | 167 } |
| 152 | 168 |
| 153 // A small hack that makes calculating percentage difference easier later on . | 169 // A small hack that makes calculating percentage difference easier later on . |
| 154 diff->result = 1.0 / ((double)baseline->width() * baseline->height()); | 170 diff->result = 1.0 / ((double)baseline->width() * baseline->height()); |
| 155 | 171 |
| 156 // Make a buffer to store results into | 172 // Make a buffer to store results into |
| 157 int numDiffPixels = 0; | 173 int zero = 0; |
|
bsalomon
2013/06/28 14:47:17
?
Zach Reizner
2013/06/28 14:52:58
I have to initialize OpenCL buffers with pointers.
bsalomon
2013/06/28 14:56:37
ok, how about
static const int kZero;
| |
| 158 diff->resultsBuffer = clCreateBuffer(fContext, CL_MEM_READ_WRITE | CL_MEM_CO PY_HOST_PTR, | 174 diff->resultsBuffer = clCreateBuffer(fContext, CL_MEM_READ_WRITE | CL_MEM_CO PY_HOST_PTR, |
| 159 sizeof(int), &numDiffPixels, NULL); | 175 sizeof(int), &zero, NULL); |
| 176 | |
| 177 diff->poiBuffer = clCreateBuffer(fContext, CL_MEM_WRITE_ONLY, | |
| 178 sizeof(int) * 2 * baseline->width() * basel ine->height(), | |
| 179 NULL, NULL); | |
| 160 | 180 |
| 161 // Set all kernel arguments | 181 // Set all kernel arguments |
| 162 cl_int setArgErr = clSetKernelArg(fKernel, 0, sizeof(cl_mem), &diff->baselin e); | 182 cl_int setArgErr = clSetKernelArg(fKernel, 0, sizeof(cl_mem), &diff->baselin e); |
| 163 setArgErr |= clSetKernelArg(fKernel, 1, sizeof(cl_mem), &diff->test); | 183 setArgErr |= clSetKernelArg(fKernel, 1, sizeof(cl_mem), &diff->test); |
| 164 setArgErr |= clSetKernelArg(fKernel, 2, sizeof(cl_mem), &diff->results Buffer); | 184 setArgErr |= clSetKernelArg(fKernel, 2, sizeof(cl_mem), &diff->results Buffer); |
| 185 setArgErr |= clSetKernelArg(fKernel, 3, sizeof(cl_mem), &diff->poiBuff er); | |
| 165 if (CL_SUCCESS != setArgErr) { | 186 if (CL_SUCCESS != setArgErr) { |
| 166 SkDebugf("Set arg failed: %s\n", cl_error_to_string(setArgErr)); | 187 SkDebugf("Set arg failed: %s\n", cl_error_to_string(setArgErr)); |
| 167 fIsGood = false; | 188 fIsGood = false; |
| 168 return -1; | 189 return -1; |
| 169 } | 190 } |
| 170 | 191 |
| 171 // Queue this diff on the CL device | 192 // Queue this diff on the CL device |
| 172 cl_event event; | 193 cl_event event; |
| 173 const size_t workSize[] = { baseline->width(), baseline->height() }; | 194 const size_t workSize[] = { baseline->width(), baseline->height() }; |
| 174 cl_int enqueueErr; | 195 cl_int enqueueErr; |
| 175 enqueueErr = clEnqueueNDRangeKernel(fCommandQueue, fKernel, 2, NULL, workSiz e, NULL, 0, NULL, &event); | 196 enqueueErr = clEnqueueNDRangeKernel(fCommandQueue, fKernel, 2, NULL, workSiz e, |
| 197 NULL, 0, NULL, &event); | |
| 176 if (CL_SUCCESS != enqueueErr) { | 198 if (CL_SUCCESS != enqueueErr) { |
| 177 SkDebugf("Enqueue failed: %s\n", cl_error_to_string(enqueueErr)); | 199 SkDebugf("Enqueue failed: %s\n", cl_error_to_string(enqueueErr)); |
| 178 fIsGood = false; | 200 fIsGood = false; |
| 179 return -1; | 201 return -1; |
| 180 } | 202 } |
| 181 | 203 |
| 182 // This makes things totally synchronous. Actual queue is not ready yet | 204 // This makes things totally synchronous. Actual queue is not ready yet |
| 183 clWaitForEvents(1, &event); | 205 clWaitForEvents(1, &event); |
| 184 diff->finished = true; | 206 diff->finished = true; |
| 185 | 207 |
| 186 // Immediate read back the results | 208 // Immediate read back the results |
| 187 clEnqueueReadBuffer(fCommandQueue, diff->resultsBuffer, CL_TRUE, 0, sizeof(i nt), &numDiffPixels, 0, NULL, NULL); | 209 clEnqueueReadBuffer(fCommandQueue, diff->resultsBuffer, CL_TRUE, 0, |
| 188 diff->result *= (double)numDiffPixels; | 210 sizeof(int), &diff->numDiffPixels, 0, NULL, NULL); |
| 211 diff->result *= (double)diff->numDiffPixels; | |
| 189 diff->result = (1.0 - diff->result); | 212 diff->result = (1.0 - diff->result); |
| 190 | 213 |
| 214 diff->poi = SkNEW_ARRAY(SkIPoint, diff->numDiffPixels); | |
| 215 clEnqueueReadBuffer(fCommandQueue, diff->poiBuffer, CL_TRUE, 0, | |
| 216 sizeof(SkIPoint) * diff->numDiffPixels, diff->poi, 0, NU LL, NULL); | |
| 217 | |
| 191 // Release all the buffers created | 218 // Release all the buffers created |
| 219 clReleaseMemObject(diff->poiBuffer); | |
| 192 clReleaseMemObject(diff->resultsBuffer); | 220 clReleaseMemObject(diff->resultsBuffer); |
| 193 clReleaseMemObject(diff->baseline); | 221 clReleaseMemObject(diff->baseline); |
| 194 clReleaseMemObject(diff->test); | 222 clReleaseMemObject(diff->test); |
| 195 | 223 |
| 196 SkDebugf("Time: %f\n", (get_seconds() - startTime)); | 224 SkDebugf("Time: %f\n", (get_seconds() - startTime)); |
| 197 | 225 |
| 198 return diffID; | 226 return diffID; |
| 199 } | 227 } |
| 200 | 228 |
| 229 void SkDifferentPixelsImageDiffer::deleteDiff(int id) { | |
| 230 QueuedDiff* diff = &fQueuedDiffs[id]; | |
| 231 if (NULL != diff->poi) { | |
| 232 SkDELETE_ARRAY(diff->poi); | |
| 233 diff->poi = NULL; | |
| 234 } | |
| 235 } | |
| 236 | |
| 201 bool SkDifferentPixelsImageDiffer::isFinished(int id) { | 237 bool SkDifferentPixelsImageDiffer::isFinished(int id) { |
| 202 return fQueuedDiffs[id].finished; | 238 return fQueuedDiffs[id].finished; |
| 203 } | 239 } |
| 204 | 240 |
| 205 double SkDifferentPixelsImageDiffer::getResult(int id) { | 241 double SkDifferentPixelsImageDiffer::getResult(int id) { |
| 206 return fQueuedDiffs[id].result; | 242 return fQueuedDiffs[id].result; |
| 207 } | 243 } |
| 208 | 244 |
| 245 int SkDifferentPixelsImageDiffer::getPointsOfInterestCount(int id) { | |
| 246 return fQueuedDiffs[id].numDiffPixels; | |
| 247 } | |
| 248 | |
| 249 SkIPoint* SkDifferentPixelsImageDiffer::getPointsOfInterest(int id) { | |
| 250 return fQueuedDiffs[id].poi; | |
| 251 } | |
| 209 | 252 |
| 210 bool SkDifferentPixelsImageDiffer::onInit() { | 253 bool SkDifferentPixelsImageDiffer::onInit() { |
| 211 if (!loadKernelFile("experimental/skpdiff/diff_pixels.cl", "diff", &fKernel) ) { | 254 if (!loadKernelFile("experimental/skpdiff/diff_pixels.cl", "diff", &fKernel) ) { |
| 212 return false; | 255 return false; |
| 213 } | 256 } |
| 214 | 257 |
| 215 return true; | 258 return true; |
| 216 } | 259 } |
| OLD | NEW |