Chromium Code Reviews
chromiumcodereview-hr@appspot.gserviceaccount.com (chromiumcodereview-hr) | Please choose your nickname with Settings | Help | Chromium Project | Gerrit Changes | Sign out
(67)

Side by Side Diff: experimental/skpdiff/SkDifferentPixelsMetric_opencl.cpp

Issue 19787006: ports for mac, ios, android, linux, windows (Closed) Base URL: https://skia.googlecode.com/svn/trunk
Patch Set: Created 7 years, 5 months ago
Use n/p to move between diff chunks; N/P to move between comments. Draft comments are only viewable by you.
Jump to:
View unified diff | Download patch | Annotate | Revision Log
OLDNEW
1 /* 1 /*
2 * Copyright 2013 Google Inc. 2 * Copyright 2013 Google Inc.
3 * 3 *
4 * Use of this source code is governed by a BSD-style license that can be 4 * Use of this source code is governed by a BSD-style license that can be
5 * found in the LICENSE file. 5 * found in the LICENSE file.
6 */ 6 */
7 7
8 #include "SkBitmap.h" 8 #include "SkBitmap.h"
9 9
10 #include "SkDifferentPixelsMetric.h" 10 #include "SkDifferentPixelsMetric.h"
11 #include "skpdiff_util.h" 11 #include "skpdiff_util.h"
12 12
13 static const char kDifferentPixelsKernelSource[] = 13 static const char kDifferentPixelsKernelSource[] =
14 "#pragma OPENCL_EXTENSION cl_khr_global_int32_base_atomics \n" 14 "#pragma OPENCL_EXTENSION cl_khr_global_int32_base_atomics \n"
15 " \n" 15 " \n"
16 "const sampler_t gInSampler = CLK_NORMALIZED_COORDS_FALSE | \n" 16 "const sampler_t gInSampler = CLK_NORMALIZED_COORDS_FALSE | \n"
17 " CLK_ADDRESS_CLAMP_TO_EDGE | \n" 17 " CLK_ADDRESS_CLAMP_TO_EDGE | \n"
18 " CLK_FILTER_NEAREST; \n" 18 " CLK_FILTER_NEAREST; \n"
19 " \n" 19 " \n"
20 "__kernel void diff(read_only image2d_t baseline, read_only image2d_t test, \n" 20 "__kernel void diff(read_only image2d_t baseline, read_only image2d_t test, \n"
21 " __global int* result, __global int2* poi) { \n" 21 " __global int* result, __global int2* poi) { \n"
22 " int2 coord = (int2)(get_global_id(0), get_global_id(1)); \n" 22 " int2 coord = (int2)(get_global_id(0), get_global_id(1)); \n"
23 " uint4 baselinePixel = read_imageui(baseline, gInSampler, coord); \n" 23 " uint4 baselinePixel = read_imageui(baseline, gInSampler, coord); \n"
24 " uint4 testPixel = read_imageui(test, gInSampler, coord); \n" 24 " uint4 testPixel = read_imageui(test, gInSampler, coord); \n"
25 " int4 pixelCompare = baselinePixel == testPixel; \n"
26 " if (baselinePixel.x != testPixel.x || \n" 25 " if (baselinePixel.x != testPixel.x || \n"
27 " baselinePixel.y != testPixel.y || \n" 26 " baselinePixel.y != testPixel.y || \n"
28 " baselinePixel.z != testPixel.z || \n" 27 " baselinePixel.z != testPixel.z || \n"
29 " baselinePixel.w != testPixel.w) { \n" 28 " baselinePixel.w != testPixel.w) { \n"
30 " \n" 29 " \n"
31 " int poiIndex = atomic_inc(result); \n" 30 " int poiIndex = atomic_inc(result); \n"
32 " poi[poiIndex] = coord; \n" 31 " poi[poiIndex] = coord; \n"
33 " } \n" 32 " } \n"
34 "} \n"; 33 "} \n";
35 34
(...skipping 77 matching lines...) Expand 10 before | Expand all | Expand 10 after
113 // This makes things totally synchronous. Actual queue is not ready yet 112 // This makes things totally synchronous. Actual queue is not ready yet
114 clWaitForEvents(1, &event); 113 clWaitForEvents(1, &event);
115 diff->finished = true; 114 diff->finished = true;
116 115
117 // Immediate read back the results 116 // Immediate read back the results
118 clEnqueueReadBuffer(fCommandQueue, diff->resultsBuffer, CL_TRUE, 0, 117 clEnqueueReadBuffer(fCommandQueue, diff->resultsBuffer, CL_TRUE, 0,
119 sizeof(int), &diff->numDiffPixels, 0, NULL, NULL); 118 sizeof(int), &diff->numDiffPixels, 0, NULL, NULL);
120 diff->result *= (double)diff->numDiffPixels; 119 diff->result *= (double)diff->numDiffPixels;
121 diff->result = (1.0 - diff->result); 120 diff->result = (1.0 - diff->result);
122 121
123 diff->poi = SkNEW_ARRAY(SkIPoint, diff->numDiffPixels); 122 // Reading a buffer of size zero can cause issues on some (Mac) OpenCL platf orms.
124 clEnqueueReadBuffer(fCommandQueue, diff->poiBuffer, CL_TRUE, 0, 123 if (diff->numDiffPixels > 0) {
124 diff->poi = SkNEW_ARRAY(SkIPoint, diff->numDiffPixels);
125 clEnqueueReadBuffer(fCommandQueue, diff->poiBuffer, CL_TRUE, 0,
125 sizeof(SkIPoint) * diff->numDiffPixels, diff->poi, 0, NU LL, NULL); 126 sizeof(SkIPoint) * diff->numDiffPixels, diff->poi, 0, NU LL, NULL);
127 }
126 128
127 // Release all the buffers created 129 // Release all the buffers created
128 clReleaseMemObject(diff->poiBuffer); 130 clReleaseMemObject(diff->poiBuffer);
129 clReleaseMemObject(diff->resultsBuffer); 131 clReleaseMemObject(diff->resultsBuffer);
130 clReleaseMemObject(diff->baseline); 132 clReleaseMemObject(diff->baseline);
131 clReleaseMemObject(diff->test); 133 clReleaseMemObject(diff->test);
132 134
133 SkDebugf("Time: %f\n", (get_seconds() - startTime)); 135 SkDebugf("Time: %f\n", (get_seconds() - startTime));
134 136
135 return diffID; 137 return diffID;
(...skipping 23 matching lines...) Expand all
159 return fQueuedDiffs[id].poi; 161 return fQueuedDiffs[id].poi;
160 } 162 }
161 163
162 bool SkDifferentPixelsMetric::onInit() { 164 bool SkDifferentPixelsMetric::onInit() {
163 if (!this->loadKernelSource(kDifferentPixelsKernelSource, "diff", &fKernel)) { 165 if (!this->loadKernelSource(kDifferentPixelsKernelSource, "diff", &fKernel)) {
164 return false; 166 return false;
165 } 167 }
166 168
167 return true; 169 return true;
168 } 170 }
OLDNEW

Powered by Google App Engine
This is Rietveld 408576698