Parallel Image
Searching
Using PostgreSQL
and PgOpenCL
Tim Child
CEO
3DMashUp
Speaker Bio
• Tim Child
• 35 years experience of software development
• Formerly
• VP Oracle Corporation
• VP BEA Systems Inc.
• VP Informix
• Leader at Illustra, Autodesk, Navteq, Intuit, …
• 30+ years experience in 3D, CAD, GIS and DBMS
Outline
• Speaker’s Bio
• Outline
• Goals
• Image Applications
• GPU Comparison
• OpenCL GPU/CPU Programming Language
• PgOpenCL
• System Architecture
• Image Data Types
• Image Processing
• Workflow
• Example Image = Operator
• Future Directions
• Q &A
GPU Accelerated Database
Overall Goals
• Develop New Applications
o Develop new GPU Accelerated Database Applications that are computationally
intensive.
• Ease of Use
o Make use GPU accelerated code easier to use
o Make GPU accelerated code more mainstream to Information Technology
• Data Scalability
o Scale GPU application data size
• Enhance existing database internal operations
Image Applications
• E-Commerce
o Custom catalog images
• Medical
o X-Ray, CT Scan, MRI
• Earth Sciences, GIS
o Remote Sensing, Aerial Photography, LIDAR
• Industrial
o QA, Metrology
• Games, Entertainment
o VR, AR, Social Media
• Arts
o Photography
• Security
o Biometrics (Face, Iris, Finger prints)
• Many others
GPU Comparison
Vendor NVidia ATI Radeon Intel
Architecture Fermi Evergreen Nehalem
Cores/ ALU 448 1408 4
Simple Simple Complex
Transistors 3.1 B 2.15 B 731 M
Clock 1.5 G Hz 851 M Hz 3 G Hz
Peak Float 1500 G 2720 G 96 G
Performance FLOP / s FLOP / s FLOP / s
Peak Double 750 G 544 G 48 G
Performance FLOP / s FLOP / s FLOP / s
Memory ~ 190 G / s ~ 153 G / s ~ 30 G / s
Bandwidth
Power 250 W 253 W 80 W
Consumption
SIMD / Vector Many Many SSE4+
Instructions
OpenCL
• OpenCL - Open Compute Language
o Subset of C 99
o Open Specification
o Proposed by Apple
o Many Companies Collaborated on the Specification
o Portable, Device Agnostic (Nvidia, AMD, Intel, ARM)
o Specification maintained by Khronos Group
• PgOpenCL
o OpenCL as a PostgreSQL Procedural Language
PgOpenCL
• New PostgreSQL Procedural Language
o Language Handler
• Determines N Dimensional Range Work Group Size ( Number of Threads)
• Maps Arguments to Buffers
• Retains Buffers for Reuse
• Calls Kernel Function
• Returns results
o Language Validator
• Creates Function – Kernel Binding
o Parameter Modes, Names, Types, Qualifiers and Attribute
• Compiles
o Syntax Checking
o Generate Program Binary
• New data types
o cl_double4, cl_double8, ….
o Image2D,
• System Admin Pseudo-Tables
o Platform, Device, Run-Time, …
PgOpenCL
System Architecture
PostgreSQL
Sever
PgOpenCL
PgOpenCL
Web HTTP Web SQL SQL
SQL
Browser Server Statement Procedure
Procedure
PCIe Bus
TCP/IP
App
PostgreSQL GPGPU
Server
Disk I/O Tables
TCP/IP
PostgreSQL
Client
Image Data Types
2D or 3D
Channel Order
R G B A or A R G B or R G B or I
Channel Type
1 2 3 4 or 1 2 or 1
byte, short, half, integer, float
Image Processing
• Geometric Transformation
o Scale, Rotate, Warp, Interpolation, …
• Color Transformation
o Gray Scale, Color Balance, Histograms
• Mensuration
o Area, Centroid, Perimeter, ..
• Noise Reduction,
o Noise reduction, de-blurring, sharpening,
• Spatial Filtering
o Gaussian Filter, Laplace Filter, Sobel Filter, ..
• Transformations
o DCT, DFT, FFT
PgOpenCL Image2D Type
• New UDT
CREATE TYPE opencl.image2d
(INPUT="opencl.image2d_in", OUTPUT="opencl.image2d_out", DEFAULT='',
INTERNALLENGTH=-1, ALIGNMENT=int4, STORAGE=EXTENDED);
ALTER TYPE opencl.image2d OWNER TO postgres;
COMMENT ON TYPE opencl.image2d IS 'PgOpenCL 2D Image type';
• Input Format
insert into images (image) values
('CL_ARGB CL_UNSIGNED_INT8 1 1 4 /////w==');
Channel Channel Data Width, Height, Base64 Encoded
Order Type Row Pitch Pixel Stream
• Casts from - to bytea type
Workflow
Basic Example
Image Operator =
Example C Code
/** Now check each Pixel **/
for ( y = 0; y < height; y++)
{
for( x = 0;x x < width; x++)
{
if ( pixelA[ x, y ] != pixelB[ x, y ]
{
PG_RETURN_BOOL(false);
}
}
}
Example OpenCL Code
kernel void imageequalskernel(__read_only image2d_t inputImageA, __read_only image2d_t
inputImageB, __global int * result)
{
int2 coord = (int2)(get_global_id(0), get_global_id(1));
Uint4 pixelA;
Uint4 pixelB;
/** If the coordinates are within range **/
if (coord.x < (get_global_size(0) - 1) && coord.y < get_global_size(1) - 1)
{
/** Read and compare each pixel channel, if they're not equal increment the result
counter **/
pixelA = read_imageui( inputImageA, imageSampler, (int2)(coord.x, coord.y) );
pixelB = read_imageui( inputImageB, imageSampler, (int2)(coord.x, coord.y) );
/** compare each RGBA component **/
if ( pixelA.x != pixelB.x || pixelA.y != pixelB.y || pixelA.z != pixelB.z ||
pixelA.w != pixelB.w )
atomic_inc(&result[0]);
}
}
SQL Query
create type imagetype as enum ( 'jpg' );
create temporary table images
(
id serial,
name text,
type imagetype,
image image2d
);
insert into images (name, type, image) values ('test',
'jpg', 'CL_ARGB CL_UNSIGNED_INT8 1 1 4 /////w==');
select imageequalskernel(image, image) = ‘,0-’::integer*1+;
Wishful Thinking
Possible Nest Step Future Directions
• Develop Image Processing Applications
o Rounds out function
o Provides an example
• Server Side Compression/Decompressions
o Jpg
o Tiff
• Client Side Support
o C Library
o Java Library
• New Image Types
o Image3D
• New Formats
o Tiff
o Geo-Tiff
o DICOM
• Enhance Language Manager
o Async
o Image Processing Pipelines
Summary
• PostgreSQL – Centralized Data and Image Storage Platform
• Minimizing Data Transfer - Data Remains on Server
• Fast - Runs at GPU Speed
• Integrated Data and Image Queries
• Standards – OpenCL, SQL
Q&A
• PgOpenCL
• Twitter @3DMashUp
• Blog www.scribd.com/3dmashup
• OpenCL
• www.khronos.org/opencl/
• www.amd.com/us/products/technologies/stream-technology/opencl/
• http://software.intel.com/en-us/articles/intel-opencl-sdk
• http://www.nvidia.com/object/cuda_opencl_new.html
