bool SkCLImageDiffer::loadKernelSource(const char source[], const char name[], cl_kernel* kernel) { // Build the kernel source size_t sourceLen = strlen(source); cl_program program = clCreateProgramWithSource(fContext, 1, &source, &sourceLen, NULL); cl_int programErr = clBuildProgram(program, 1, &fDevice, "", NULL, NULL); if (CL_SUCCESS != programErr) { SkDebugf("Program creation failed: %s\n", cl_error_to_string(programErr)); // Attempt to get information about why the build failed char buildLog[4096]; clGetProgramBuildInfo(program, fDevice, CL_PROGRAM_BUILD_LOG, sizeof(buildLog), buildLog, NULL); SkDebugf("Build log: %s\n", buildLog); return false; } cl_int kernelErr; *kernel = clCreateKernel(program, name, &kernelErr); if (CL_SUCCESS != kernelErr) { SkDebugf("Kernel creation failed: %s\n", cl_error_to_string(kernelErr)); return false; } return true; }
bool SkCLImageDiffer::init(cl_device_id device, cl_context context) { fContext = context; fDevice = device; cl_int queueErr; fCommandQueue = clCreateCommandQueue(fContext, fDevice, 0, &queueErr); if (CL_SUCCESS != queueErr) { SkDebugf("Command queue creation failed: %s\n", cl_error_to_string(queueErr)); fIsGood = false; return false; } fIsGood = this->onInit(); return fIsGood; }
bool SkCLImageDiffer::makeImage2D(SkBitmap* bitmap, cl_mem* image) const { cl_int imageErr; cl_image_format bitmapFormat; switch (bitmap->colorType()) { case kAlpha_8_SkColorType: bitmapFormat.image_channel_order = CL_A; bitmapFormat.image_channel_data_type = CL_UNSIGNED_INT8; break; case kRGB_565_SkColorType: bitmapFormat.image_channel_order = CL_RGB; bitmapFormat.image_channel_data_type = CL_UNORM_SHORT_565; break; case kN32_SkColorType: bitmapFormat.image_channel_order = CL_RGBA; bitmapFormat.image_channel_data_type = CL_UNSIGNED_INT8; break; default: SkDebugf("Image format is unsupported\n"); return false; } // Upload the bitmap data to OpenCL bitmap->lockPixels(); *image = clCreateImage2D(fContext, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, &bitmapFormat, bitmap->width(), bitmap->height(), bitmap->rowBytes(), bitmap->getPixels(), &imageErr); bitmap->unlockPixels(); if (CL_SUCCESS != imageErr) { SkDebugf("Input image creation failed: %s\n", cl_error_to_string(imageErr)); return false; } return true; }
bool SkDifferentPixelsMetric::diff(SkBitmap* baseline, SkBitmap* test, bool computeMask, Result* result) const { double startTime = get_seconds(); if (!fIsGood) { return false; } // If we never end up running the kernel, include some safe defaults in the result. result->poiCount = 0; // Ensure the images are comparable if (baseline->width() != test->width() || baseline->height() != test->height() || baseline->width() <= 0 || baseline->height() <= 0 || baseline->config() != test->config()) { return false; } cl_mem baselineImage; cl_mem testImage; cl_mem resultsBuffer; // Upload images to the CL device if (!this->makeImage2D(baseline, &baselineImage) || !this->makeImage2D(test, &testImage)) { SkDebugf("creation of openCL images failed"); return false; } // A small hack that makes calculating percentage difference easier later on. result->result = 1.0 / ((double)baseline->width() * baseline->height()); // Make a buffer to store results into. It must be initialized with pointers to memory. static const int kZero = 0; // We know OpenCL won't write to it because we use CL_MEM_COPY_HOST_PTR resultsBuffer = clCreateBuffer(fContext, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeof(int), (int*)&kZero, NULL); // Set all kernel arguments cl_int setArgErr = clSetKernelArg(fKernel, 0, sizeof(cl_mem), &baselineImage); setArgErr |= clSetKernelArg(fKernel, 1, sizeof(cl_mem), &testImage); setArgErr |= clSetKernelArg(fKernel, 2, sizeof(cl_mem), &resultsBuffer); if (CL_SUCCESS != setArgErr) { SkDebugf("Set arg failed: %s\n", cl_error_to_string(setArgErr)); return false; } // Queue this diff on the CL device 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); if (CL_SUCCESS != enqueueErr) { SkDebugf("Enqueue failed: %s\n", cl_error_to_string(enqueueErr)); return false; } // This makes things totally synchronous. Actual queue is not ready yet clWaitForEvents(1, &event); // Immediate read back the results clEnqueueReadBuffer(fCommandQueue, resultsBuffer, CL_TRUE, 0, sizeof(int), &result->poiCount, 0, NULL, NULL); result->result *= (double)result->poiCount; result->result = (1.0 - result->result); // Release all the buffers created clReleaseMemObject(resultsBuffer); clReleaseMemObject(baselineImage); clReleaseMemObject(testImage); result->timeElapsed = get_seconds() - startTime; return true; }