Chromium Code Reviews| Index: experimental/skpdiff/SkCLImageDiffer.cpp |
| diff --git a/experimental/skpdiff/SkCLImageDiffer.cpp b/experimental/skpdiff/SkCLImageDiffer.cpp |
| index 50eb2b6fc8f4a0befaff8c6a6dcae7b108cc1230..30d7d75b4fc41412130a89d633c486d9da5f5d54 100644 |
| --- a/experimental/skpdiff/SkCLImageDiffer.cpp |
| +++ b/experimental/skpdiff/SkCLImageDiffer.cpp |
| @@ -69,7 +69,8 @@ bool SkCLImageDiffer::loadKernelSource(const char source[], const char name[], c |
| // Attempt to get information about why the build failed |
| char buildLog[4096]; |
| - clGetProgramBuildInfo(program, fDevice, CL_PROGRAM_BUILD_LOG, sizeof(buildLog), buildLog, NULL); |
| + clGetProgramBuildInfo(program, fDevice, CL_PROGRAM_BUILD_LOG, sizeof(buildLog), |
| + buildLog, NULL); |
| SkDebugf("Build log: %s\n", buildLog); |
| return false; |
| @@ -125,6 +126,17 @@ bool SkCLImageDiffer::makeImage2D(SkBitmap* bitmap, cl_mem* image) { |
| //////////////////////////////////////////////////////////////// |
| +struct SkDifferentPixelsImageDiffer::QueuedDiff { |
| + bool finished; |
| + double result; |
| + int numDiffPixels; |
| + SkIPoint* poi; |
| + cl_mem baseline; |
| + cl_mem test; |
| + cl_mem resultsBuffer; |
| + cl_mem poiBuffer; |
| +}; |
| + |
| const char* SkDifferentPixelsImageDiffer::getName() { |
| return "different_pixels"; |
| } |
| @@ -134,18 +146,22 @@ int SkDifferentPixelsImageDiffer::queueDiff(SkBitmap * baseline, SkBitmap * test |
| double startTime = get_seconds(); |
| QueuedDiff* diff = fQueuedDiffs.push(); |
| + // If we never end up running the kernel, include some safe defaults in the result. |
| + diff->finished = false; |
| + diff->result = -1.0; |
| + diff->numDiffPixels = 0; |
| + diff->poi = NULL; |
| + |
| // Ensure the images are comparable |
| if (baseline->width() != test->width() || baseline->height() != test->height() || |
| baseline->width() <= 0 || baseline->height() <= 0) { |
| diff->finished = true; |
| - diff->result = 0.0; |
| return diffID; |
| } |
| // Upload images to the CL device |
| if (!this->makeImage2D(baseline, &diff->baseline) || !this->makeImage2D(test, &diff->test)) { |
| diff->finished = true; |
| - diff->result = 0.0; |
| fIsGood = false; |
| return -1; |
| } |
| @@ -154,14 +170,19 @@ int SkDifferentPixelsImageDiffer::queueDiff(SkBitmap * baseline, SkBitmap * test |
| diff->result = 1.0 / ((double)baseline->width() * baseline->height()); |
| // Make a buffer to store results into |
| - int numDiffPixels = 0; |
| + 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;
|
| diff->resultsBuffer = clCreateBuffer(fContext, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, |
| - sizeof(int), &numDiffPixels, NULL); |
| + sizeof(int), &zero, NULL); |
| + |
| + diff->poiBuffer = clCreateBuffer(fContext, CL_MEM_WRITE_ONLY, |
| + sizeof(int) * 2 * baseline->width() * baseline->height(), |
| + NULL, NULL); |
| // Set all kernel arguments |
| cl_int setArgErr = clSetKernelArg(fKernel, 0, sizeof(cl_mem), &diff->baseline); |
| setArgErr |= clSetKernelArg(fKernel, 1, sizeof(cl_mem), &diff->test); |
| setArgErr |= clSetKernelArg(fKernel, 2, sizeof(cl_mem), &diff->resultsBuffer); |
| + setArgErr |= clSetKernelArg(fKernel, 3, sizeof(cl_mem), &diff->poiBuffer); |
| if (CL_SUCCESS != setArgErr) { |
| SkDebugf("Set arg failed: %s\n", cl_error_to_string(setArgErr)); |
| fIsGood = false; |
| @@ -172,7 +193,8 @@ int SkDifferentPixelsImageDiffer::queueDiff(SkBitmap * baseline, SkBitmap * test |
| cl_event event; |
| const size_t workSize[] = { baseline->width(), baseline->height() }; |
| cl_int enqueueErr; |
| - enqueueErr = clEnqueueNDRangeKernel(fCommandQueue, fKernel, 2, NULL, workSize, NULL, 0, NULL, &event); |
| + enqueueErr = clEnqueueNDRangeKernel(fCommandQueue, fKernel, 2, NULL, workSize, |
| + NULL, 0, NULL, &event); |
| if (CL_SUCCESS != enqueueErr) { |
| SkDebugf("Enqueue failed: %s\n", cl_error_to_string(enqueueErr)); |
| fIsGood = false; |
| @@ -184,11 +206,17 @@ int SkDifferentPixelsImageDiffer::queueDiff(SkBitmap * baseline, SkBitmap * test |
| diff->finished = true; |
| // Immediate read back the results |
| - clEnqueueReadBuffer(fCommandQueue, diff->resultsBuffer, CL_TRUE, 0, sizeof(int), &numDiffPixels, 0, NULL, NULL); |
| - diff->result *= (double)numDiffPixels; |
| + clEnqueueReadBuffer(fCommandQueue, diff->resultsBuffer, CL_TRUE, 0, |
| + sizeof(int), &diff->numDiffPixels, 0, NULL, NULL); |
| + diff->result *= (double)diff->numDiffPixels; |
| diff->result = (1.0 - diff->result); |
| + diff->poi = SkNEW_ARRAY(SkIPoint, diff->numDiffPixels); |
| + clEnqueueReadBuffer(fCommandQueue, diff->poiBuffer, CL_TRUE, 0, |
| + sizeof(SkIPoint) * diff->numDiffPixels, diff->poi, 0, NULL, NULL); |
| + |
| // Release all the buffers created |
| + clReleaseMemObject(diff->poiBuffer); |
| clReleaseMemObject(diff->resultsBuffer); |
| clReleaseMemObject(diff->baseline); |
| clReleaseMemObject(diff->test); |
| @@ -198,6 +226,14 @@ int SkDifferentPixelsImageDiffer::queueDiff(SkBitmap * baseline, SkBitmap * test |
| return diffID; |
| } |
| +void SkDifferentPixelsImageDiffer::deleteDiff(int id) { |
| + QueuedDiff* diff = &fQueuedDiffs[id]; |
| + if (NULL != diff->poi) { |
| + SkDELETE_ARRAY(diff->poi); |
| + diff->poi = NULL; |
| + } |
| +} |
| + |
| bool SkDifferentPixelsImageDiffer::isFinished(int id) { |
| return fQueuedDiffs[id].finished; |
| } |
| @@ -206,6 +242,13 @@ double SkDifferentPixelsImageDiffer::getResult(int id) { |
| return fQueuedDiffs[id].result; |
| } |
| +int SkDifferentPixelsImageDiffer::getPointsOfInterestCount(int id) { |
| + return fQueuedDiffs[id].numDiffPixels; |
| +} |
| + |
| +SkIPoint* SkDifferentPixelsImageDiffer::getPointsOfInterest(int id) { |
| + return fQueuedDiffs[id].poi; |
| +} |
| bool SkDifferentPixelsImageDiffer::onInit() { |
| if (!loadKernelFile("experimental/skpdiff/diff_pixels.cl", "diff", &fKernel)) { |