Merged changes in the trunk up to revision 54421.

Conflicts resolved:
release/datafiles/startup.blend
release/scripts/startup/bl_ui/properties_render.py
source/blender/SConscript
source/blender/blenloader/intern/readfile.c

69 files changed, 3952 insertions, 368 deletions

diff --git a/intern/CMakeLists.txt b/intern/CMakeLists.txt
index be797c45ba1..bb8cd7aaf70 100644
--- a/intern/CMakeLists.txt
+++ b/intern/CMakeLists.txt
@@ -69,6 +69,14 @@ if(WITH_INTERNATIONAL)
 	add_subdirectory(locale)
 endif()
 
+if(WITH_BULLET)
+	add_subdirectory(rigidbody)
+endif()
+
+if(WITH_COMPOSITOR)
+	add_subdirectory(opencl)
+endif()
+
 # only windows needs utf16 converter
 if(WIN32)
 	add_subdirectory(utfconv)
diff --git a/intern/SConscript b/intern/SConscript
index a35c99bbbaa..828c1adc20d 100644
--- a/intern/SConscript
+++ b/intern/SConscript
@@ -59,6 +59,12 @@ if env['WITH_BF_BOOLEAN']:
 if env['WITH_BF_INTERNATIONAL']:
     SConscript(['locale/SConscript'])
 
+if env['WITH_BF_BULLET']:
+    SConscript (['rigidbody/SConscript'])
+
+if env['WITH_BF_COMPOSITOR']:
+    SConscript (['opencl/SConscript'])
+
 if env['OURPLATFORM'] in ('win32-vc', 'win32-mingw', 'win64-mingw', 'linuxcross', 'win64-vc'):
     SConscript(['utfconv/SConscript'])
 
diff --git a/intern/cycles/CMakeLists.txt b/intern/cycles/CMakeLists.txt
index 048a2a50a7f..226218ae512 100644
--- a/intern/cycles/CMakeLists.txt
+++ b/intern/cycles/CMakeLists.txt
@@ -13,10 +13,12 @@ if(WITH_RAYOPTIMIZATION AND SUPPORT_SSE_BUILD)
 endif()
 
 if(WIN32 AND MSVC)
-	set(CYCLES_OPTIMIZED_KERNEL_FLAGS "/arch:SSE2 -D_CRT_SECURE_NO_WARNINGS /fp:fast /EHsc")
+	set(CYCLES_SSE2_KERNEL_FLAGS "/arch:SSE2 -D_CRT_SECURE_NO_WARNINGS /fp:fast /EHsc")
+	set(CYCLES_SSE3_KERNEL_FLAGS "/arch:SSE2 -D_CRT_SECURE_NO_WARNINGS /fp:fast /EHsc")
 	set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} /fp:fast -D_CRT_SECURE_NO_WARNINGS /EHsc")
 elseif(CMAKE_COMPILER_IS_GNUCC)
-	set(CYCLES_OPTIMIZED_KERNEL_FLAGS "-ffast-math -msse -msse2 -msse3 -mfpmath=sse")
+	set(CYCLES_SSE2_KERNEL_FLAGS "-ffast-math -msse -msse2 -mfpmath=sse")
+	set(CYCLES_SSE3_KERNEL_FLAGS "-ffast-math -msse -msse2 -msse3 -mfpmath=sse")
 	set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -ffast-math")
 endif()
 
@@ -44,6 +46,7 @@ endif()
 
 if(WITH_CYCLES_OSL)
 	add_definitions(-DWITH_OSL)
+	add_definitions(-DOSL_STATIC_LIBRARY)
 	include_directories(${OSL_INCLUDES})
 endif()
 
diff --git a/intern/cycles/SConscript b/intern/cycles/SConscript
index 44a17ac0cd6..8a8ef9cce39 100644
--- a/intern/cycles/SConscript
+++ b/intern/cycles/SConscript
@@ -36,7 +36,8 @@ sources = cycles.Glob('bvh/*.cpp') + cycles.Glob('device/*.cpp') + cycles.Glob('
 
 sources.remove(path.join('util', 'util_view.cpp'))
 sources.remove(path.join('render', 'film_response.cpp'))
-sources.remove(path.join('kernel', 'kernel_optimized.cpp'))
+sources.remove(path.join('kernel', 'kernel_sse2.cpp'))
+sources.remove(path.join('kernel', 'kernel_sse3.cpp'))
 
 incs = [] 
 defs = []
@@ -51,6 +52,7 @@ defs.append('WITH_CUDA')
 
 if env['WITH_BF_CYCLES_OSL']:
     defs.append('WITH_OSL')
+    defs.append('OSL_STATIC_LIBRARY')
     incs.append(cycles['BF_OSL_INC'])
 
 incs.extend('. bvh render device kernel kernel/osl kernel/svm util subd'.split())
@@ -72,21 +74,29 @@ if env['OURPLATFORM'] in ('win32-vc', 'win32-mingw', 'linuxcross', 'win64-vc', '
 
 # optimized kernel
 if env['WITH_BF_RAYOPTIMIZATION']:
-    optim_cxxflags = Split(env['CXXFLAGS'])
+    sse2_cxxflags = Split(env['CXXFLAGS'])
+    sse3_cxxflags = Split(env['CXXFLAGS'])
 
     if env['OURPLATFORM'] == 'win32-vc':
-        optim_cxxflags.append('/arch:SSE2 -D_CRT_SECURE_NO_WARNINGS /fp:fast /EHsc'.split())
+        sse2_cxxflags.append('/arch:SSE2 -D_CRT_SECURE_NO_WARNINGS /fp:fast /EHsc'.split())
+        sse3_cxxflags.append('/arch:SSE2 -D_CRT_SECURE_NO_WARNINGS /fp:fast /EHsc'.split())
     elif env['OURPLATFORM'] == 'win64-vc':
-        optim_cxxflags.append('-D_CRT_SECURE_NO_WARNINGS /fp:fast /EHsc'.split())
+        sse2_cxxflags.append('-D_CRT_SECURE_NO_WARNINGS /fp:fast /EHsc'.split())
+        sse3_cxxflags.append('-D_CRT_SECURE_NO_WARNINGS /fp:fast /EHsc'.split())
     else:
-        optim_cxxflags.append('-ffast-math -msse -msse2 -msse3 -mfpmath=sse'.split())
+        sse2_cxxflags.append('-ffast-math -msse -msse2 -mfpmath=sse'.split())
+        sse3_cxxflags.append('-ffast-math -msse -msse2 -msse3 -mfpmath=sse'.split())
     
     defs.append('WITH_OPTIMIZED_KERNEL')
     optim_defs = defs[:]
-    optim_sources = [path.join('kernel', 'kernel_optimized.cpp')]
 
-    cycles_optim = cycles.Clone()
-    cycles_optim.BlenderLib('bf_intern_cycles_optimized', optim_sources, incs, optim_defs, libtype=['intern'], priority=[10], cxx_compileflags=optim_cxxflags)
+    cycles_sse3 = cycles.Clone()
+    sse3_sources = [path.join('kernel', 'kernel_sse3.cpp')]
+    cycles_sse3.BlenderLib('bf_intern_cycles_sse3', sse3_sources, incs, optim_defs, libtype=['intern'], priority=[10], cxx_compileflags=sse3_cxxflags)
+
+    cycles_sse2 = cycles.Clone()
+    sse2_sources = [path.join('kernel', 'kernel_sse2.cpp')]
+    cycles_sse2.BlenderLib('bf_intern_cycles_sse2', sse2_sources, incs, optim_defs, libtype=['intern'], priority=[10], cxx_compileflags=sse2_cxxflags)
 
 cycles.BlenderLib('bf_intern_cycles', sources, incs, defs, libtype=['intern'], priority=[0], cxx_compileflags=cxxflags)
 
diff --git a/intern/cycles/blender/addon/properties.py b/intern/cycles/blender/addon/properties.py
index 3661274ae43..65f245fba69 100644
--- a/intern/cycles/blender/addon/properties.py
+++ b/intern/cycles/blender/addon/properties.py
@@ -66,9 +66,11 @@ enum_panorama_types = (
 
 enum_curve_presets = (
     ('CUSTOM', "Custom", "Set general parameters"),
-    ('TANGENT_SHADING', "Tangent Normal", "Use planar geometry and tangent normals"),
-    ('TRUE_NORMAL', "True Normal", "Use true normals (good for thin strands)"),
-    ('ACCURATE_PRESET', "Accurate", "Use best settings (suitable for glass materials)"),
+    ('FAST_PLANES', "Fast Planes", "Use camera facing triangles (fast but memory intensive)"),
+    ('TANGENT_SHADING', "Tangent Normal", "Use planar line segments and tangent normals"),
+    ('TRUE_NORMAL', "True Normal", "Use true normals with line segments(good for thin strands)"),
+    ('ACCURATE_PRESET', "Accurate", "Use best line segment settings (suitable for glass materials)"),
+    ('SMOOTH_CURVES', "Smooth Curves", "Use smooth cardinal curves (slowest)"),
     )
 
 enum_curve_primitives = (
@@ -462,8 +464,8 @@ class CyclesMaterialSettings(bpy.types.PropertyGroup):
                 type=cls,
                 )
         cls.sample_as_light = BoolProperty(
-                name="Sample as Lamp",
-                description="Use direct light sampling for this material, "
+                name="Multiple Importance Sample",
+                description="Use multiple importance sampling for this material, "
                             "disabling may reduce overall noise for large "
                             "objects that emit little light compared to other light sources",
                 default=True,
@@ -499,6 +501,12 @@ class CyclesLampSettings(bpy.types.PropertyGroup):
                 min=1, max=10000,
                 default=1,
                 )
+        cls.use_multiple_importance_sampling = BoolProperty(
+                name="Multiple Importance Sample",
+                description="Use multiple importance sampling for the lamp, "
+                            "reduces noise for area lamps and sharp glossy materials",
+                default=False,
+                )
 
     @classmethod
     def unregister(cls):
@@ -514,8 +522,8 @@ class CyclesWorldSettings(bpy.types.PropertyGroup):
                 type=cls,
                 )
         cls.sample_as_light = BoolProperty(
-                name="Sample as Lamp",
-                description="Use direct light sampling for the environment, "
+                name="Multiple Importance Sample",
+                description="Use multiple importance sampling for the environment, "
                             "enabling for non-solid colors is recommended",
                 default=False,
                 )
diff --git a/intern/cycles/blender/addon/ui.py b/intern/cycles/blender/addon/ui.py
index 9d0b6d96f47..b9fbb85581e 100644
--- a/intern/cycles/blender/addon/ui.py
+++ b/intern/cycles/blender/addon/ui.py
@@ -168,6 +168,19 @@ class CyclesRender_PT_film(CyclesButtonsPanel, Panel):
         if cscene.filter_type != 'BOX':
             sub.prop(cscene, "filter_width", text="Width")
 
+        layout.separator()
+
+        rd = scene.render
+        col = layout.column()
+
+        split = col.split(percentage=0.40)
+        split.prop(rd, "use_antialiasing", "OpenGL AA")
+        row = split.row()
+        row.active = rd.use_antialiasing
+        row.prop(rd, "antialiasing_samples", expand=True)
+
+        col.prop(rd, "alpha_mode", text="OpenGL Alpha")
+
 
 class CyclesRender_PT_performance(CyclesButtonsPanel, Panel):
     bl_label = "Performance"
@@ -552,6 +565,19 @@ def panel_node_draw(layout, id_data, output_type, input_name):
     return True
 
 
+class CyclesLamp_PT_preview(CyclesButtonsPanel, Panel):
+    bl_label = "Preview"
+    bl_context = "data"
+    bl_options = {'DEFAULT_CLOSED'}
+
+    @classmethod
+    def poll(cls, context):
+        return context.lamp and CyclesButtonsPanel.poll(context)
+
+    def draw(self, context):
+        self.layout.template_preview(context.lamp)
+
+
 class CyclesLamp_PT_lamp(CyclesButtonsPanel, Panel):
     bl_label = "Lamp"
     bl_context = "data"
@@ -591,6 +617,8 @@ class CyclesLamp_PT_lamp(CyclesButtonsPanel, Panel):
         col = split.column()
         col.prop(clamp, "cast_shadow")
 
+        layout.prop(clamp, "use_multiple_importance_sampling")
+
         if lamp.type == 'HEMI':
             layout.label(text="Not supported, interpreted as sun lamp.")
 
@@ -636,6 +664,19 @@ class CyclesLamp_PT_spot(CyclesButtonsPanel, Panel):
         col.prop(lamp, "show_cone")
 
 
+class CyclesWorld_PT_preview(CyclesButtonsPanel, Panel):
+    bl_label = "Preview"
+    bl_context = "world"
+    bl_options = {'DEFAULT_CLOSED'}
+
+    @classmethod
+    def poll(cls, context):
+        return context.world and CyclesButtonsPanel.poll(context)
+
+    def draw(self, context):
+        self.layout.template_preview(context.world)
+
+
 class CyclesWorld_PT_surface(CyclesButtonsPanel, Panel):
     bl_label = "Surface"
     bl_context = "world"
@@ -722,6 +763,19 @@ class CyclesWorld_PT_settings(CyclesButtonsPanel, Panel):
             sub.prop(cworld, "samples")
 
 
+class CyclesMaterial_PT_preview(CyclesButtonsPanel, Panel):
+    bl_label = "Preview"
+    bl_context = "material"
+    bl_options = {'DEFAULT_CLOSED'}
+
+    @classmethod
+    def poll(cls, context):
+        return context.material and CyclesButtonsPanel.poll(context)
+
+    def draw(self, context):
+        self.layout.template_preview(context.material)
+
+
 class CyclesMaterial_PT_surface(CyclesButtonsPanel, Panel):
     bl_label = "Surface"
     bl_context = "material"
@@ -798,9 +852,10 @@ class CyclesMaterial_PT_settings(CyclesButtonsPanel, Panel):
         col.prop(mat, "diffuse_color", text="Viewport Color")
 
         col = split.column()
-        col.prop(cmat, "sample_as_light")
         col.prop(mat, "pass_index")
 
+        layout.prop(cmat, "sample_as_light")
+
 
 class CyclesTexture_PT_context(CyclesButtonsPanel, Panel):
     bl_label = ""
@@ -838,24 +893,8 @@ class CyclesTexture_PT_context(CyclesButtonsPanel, Panel):
                 split = layout.split(percentage=0.2)
                 split.label(text="Type:")
                 split.prop(tex, "type", text="")
-
-
-class CyclesTexture_PT_nodes(CyclesButtonsPanel, Panel):
-    bl_label = "Nodes"
-    bl_context = "texture"
-
-    @classmethod
-    def poll(cls, context):
-        tex = context.texture
-        return (tex and tex.use_nodes) and CyclesButtonsPanel.poll(context)
-
-    def draw(self, context):
-        layout = self.layout
-
-        tex = context.texture
-        panel_node_draw(layout, tex, 'OUTPUT_TEXTURE', 'Color')
-
-
+                
+                
 class CyclesTexture_PT_node(CyclesButtonsPanel, Panel):
     bl_label = "Node"
     bl_context = "texture"
@@ -879,14 +918,12 @@ class CyclesTexture_PT_mapping(CyclesButtonsPanel, Panel):
 
     @classmethod
     def poll(cls, context):
-        tex = context.texture
         node = context.texture_node
-        return (node or (tex and tex.use_nodes)) and CyclesButtonsPanel.poll(context)
+        return node and CyclesButtonsPanel.poll(context)
 
     def draw(self, context):
         layout = self.layout
-
-        # tex = context.texture
+        
         node = context.texture_node
 
         mapping = node.texture_mapping
@@ -912,15 +949,13 @@ class CyclesTexture_PT_colors(CyclesButtonsPanel, Panel):
 
     @classmethod
     def poll(cls, context):
-        # tex = context.texture
         # node = context.texture_node
         return False
-        #return (node or (tex and tex.use_nodes)) and CyclesButtonsPanel.poll(context)
+        #return node and CyclesButtonsPanel.poll(context)
 
     def draw(self, context):
         layout = self.layout
-
-        # tex = context.texture
+        
         node = context.texture_node
 
         mapping = node.color_mapping
@@ -1141,12 +1176,14 @@ def get_panels():
         types.RENDER_PT_stamp,
         types.SCENE_PT_scene,
         types.SCENE_PT_color_management,
+        types.SCENE_PT_custom_props,
         types.SCENE_PT_audio,
         types.SCENE_PT_unit,
         types.SCENE_PT_keying_sets,
         types.SCENE_PT_keying_set_paths,
         types.SCENE_PT_physics,
         types.WORLD_PT_context_world,
+        types.WORLD_PT_custom_props,
         types.DATA_PT_context_mesh,
         types.DATA_PT_context_camera,
         types.DATA_PT_context_lamp,
@@ -1169,6 +1206,11 @@ def get_panels():
         types.DATA_PT_custom_props_camera,
         types.DATA_PT_custom_props_lamp,
         types.DATA_PT_custom_props_speaker,
+        types.DATA_PT_custom_props_arm,
+        types.DATA_PT_custom_props_curve,
+        types.DATA_PT_custom_props_lattice,
+        types.DATA_PT_custom_props_metaball,
+        types.TEXTURE_PT_custom_props,
         types.TEXTURE_PT_clouds,
         types.TEXTURE_PT_wood,
         types.TEXTURE_PT_marble,
@@ -1188,6 +1230,7 @@ def get_panels():
         types.TEXTURE_PT_influence,
         types.TEXTURE_PT_colors,
         types.PARTICLE_PT_context_particles,
+        types.PARTICLE_PT_custom_props,
         types.PARTICLE_PT_emission,
         types.PARTICLE_PT_hair_dynamics,
         types.PARTICLE_PT_cache,
@@ -1204,10 +1247,9 @@ def get_panels():
         types.PARTICLE_PT_field_weights,
         types.PARTICLE_PT_force_fields,
         types.PARTICLE_PT_vertexgroups,
-        types.PARTICLE_PT_custom_props,
-        types.MATERIAL_PT_preview,
-        types.DATA_PT_preview,
-        types.WORLD_PT_preview,
+        types.MATERIAL_PT_custom_props,
+        types.BONE_PT_custom_props,
+        types.OBJECT_PT_custom_props,
         )
 
 
diff --git a/intern/cycles/blender/blender_curves.cpp b/intern/cycles/blender/blender_curves.cpp
index cf2c018f1a2..5b8d495084d 100644
--- a/intern/cycles/blender/blender_curves.cpp
+++ b/intern/cycles/blender/blender_curves.cpp
@@ -37,14 +37,15 @@ void curveinterp_v3_v3v3v3v3(float3 *p, float3 *v1, float3 *v2, float3 *v3, floa
 void interp_weights(float t, float data[4], int type);
 float shaperadius(float shape, float root, float tip, float time);
 void InterpolateKeySegments(int seg, int segno, int key, int curve, float3 *keyloc, float *time, ParticleCurveData *CData, int interpolation);
-bool ObtainCacheParticleUV(Mesh *mesh, BL::Mesh *b_mesh, BL::Object *b_ob, ParticleCurveData *CData, bool use_parents, int uv_num);
-bool ObtainCacheParticleVcol(Mesh *mesh, BL::Mesh *b_mesh, BL::Object *b_ob, ParticleCurveData *CData, bool use_parents, int vcol_num);
-bool ObtainCacheParticleData(Mesh *mesh, BL::Mesh *b_mesh, BL::Object *b_ob, ParticleCurveData *CData, bool use_parents);
+bool ObtainCacheParticleUV(Mesh *mesh, BL::Mesh *b_mesh, BL::Object *b_ob, ParticleCurveData *CData, bool use_parents, bool background, int uv_num);
+bool ObtainCacheParticleVcol(Mesh *mesh, BL::Mesh *b_mesh, BL::Object *b_ob, ParticleCurveData *CData, bool use_parents, bool background, int vcol_num);
+bool ObtainCacheParticleData(Mesh *mesh, BL::Mesh *b_mesh, BL::Object *b_ob, ParticleCurveData *CData, bool use_parents, bool background);
+void ExportCurveSegments(Mesh *mesh, ParticleCurveData *CData, int interpolation, int segments);
 void ExportCurveTrianglePlanes(Mesh *mesh, ParticleCurveData *CData, int interpolation, bool use_smooth, int segments, float3 RotCam);
 void ExportCurveTriangleRibbons(Mesh *mesh, ParticleCurveData *CData, int interpolation, bool use_smooth, int segments);
 void ExportCurveTriangleGeometry(Mesh *mesh, ParticleCurveData *CData, int interpolation, bool use_smooth, int resolution, int segments);
-void ExportCurveSegments(Mesh *mesh, ParticleCurveData *CData, int interpolation, int segments);
-void ExportCurveTriangleUVs(Mesh *mesh, ParticleCurveData *CData, int interpolation, bool use_smooth, int segments, int vert_offset, int resol);
+void ExportCurveTriangleUV(Mesh *mesh, ParticleCurveData *CData, int interpolation, int segments, int vert_offset, int resol, float3 *uvdata);
+void ExportCurveTriangleVcol(Mesh *mesh, ParticleCurveData *CData, int interpolation, int segments, int vert_offset, int resol, float3 *fdata);
 
 ParticleCurveData::ParticleCurveData()
 {
@@ -151,7 +152,7 @@ void InterpolateKeySegments(int seg, int segno, int key, int curve, float3 *keyl
 		curveinterp_v3_v3v3v3v3(keyloc, &ckey_loc1, &ckey_loc2, &ckey_loc3, &ckey_loc4, t);
 }
 
-bool ObtainCacheParticleData(Mesh *mesh, BL::Mesh *b_mesh, BL::Object *b_ob, ParticleCurveData *CData, bool use_parents)
+bool ObtainCacheParticleData(Mesh *mesh, BL::Mesh *b_mesh, BL::Object *b_ob, ParticleCurveData *CData, bool use_parents, bool background)
 {
 
 	int curvenum = 0;
@@ -176,12 +177,10 @@ bool ObtainCacheParticleData(Mesh *mesh, BL::Mesh *b_mesh, BL::Object *b_ob, Par
 
 				int mi = clamp(b_psys.settings().material()-1, 0, mesh->used_shaders.size()-1);
 				int shader = mesh->used_shaders[mi];
-				int draw_step = b_psys.settings().draw_step();
+				int draw_step = background ? b_psys.settings().render_step() : b_psys.settings().draw_step();
 				int ren_step = (int)pow((float)2.0f,(float)draw_step);
-				/*b_psys.settings().render_step(draw_step);*/
-
 				int totparts = b_psys.particles.length();
-				int totchild = b_psys.child_particles.length() * b_psys.settings().draw_percentage() / 100;
+				int totchild = background ? b_psys.child_particles.length() : (int)((float)b_psys.child_particles.length() * (float)b_psys.settings().draw_percentage() / 100.0f);
 				int totcurves = totchild;
 				
 				if(use_parents || b_psys.settings().child_type() == 0)
@@ -240,7 +239,7 @@ bool ObtainCacheParticleData(Mesh *mesh, BL::Mesh *b_mesh, BL::Object *b_ob, Par
 
 }
 
-bool ObtainCacheParticleUV(Mesh *mesh, BL::Mesh *b_mesh, BL::Object *b_ob, ParticleCurveData *CData, bool use_parents, int uv_num)
+bool ObtainCacheParticleUV(Mesh *mesh, BL::Mesh *b_mesh, BL::Object *b_ob, ParticleCurveData *CData, bool use_parents, bool background, int uv_num)
 {
 #if 0
 	int keyno = 0;
@@ -269,13 +268,10 @@ bool ObtainCacheParticleUV(Mesh *mesh, BL::Mesh *b_mesh, BL::Object *b_ob, Parti
 #if 0
 				int mi = clamp(b_psys.settings().material()-1, 0, mesh->used_shaders.size()-1);
 				int shader = mesh->used_shaders[mi];
-				int draw_step = b_psys.settings().draw_step();
-				int ren_step = (int)pow((float)2.0f,(float)draw_step);
-				b_psys.settings().render_step(draw_step);
 #endif
 
 				int totparts = b_psys.particles.length();
-				int totchild = b_psys.child_particles.length() * b_psys.settings().draw_percentage() / 100;
+				int totchild = background ? b_psys.child_particles.length() : (int)((float)b_psys.child_particles.length() * (float)b_psys.settings().draw_percentage() / 100.0f);
 				int totcurves = totchild;
 				
 				if (use_parents || b_psys.settings().child_type() == 0)
@@ -314,7 +310,7 @@ bool ObtainCacheParticleUV(Mesh *mesh, BL::Mesh *b_mesh, BL::Object *b_ob, Parti
 
 }
 
-bool ObtainCacheParticleVcol(Mesh *mesh, BL::Mesh *b_mesh, BL::Object *b_ob, ParticleCurveData *CData, bool use_parents, int vcol_num)
+bool ObtainCacheParticleVcol(Mesh *mesh, BL::Mesh *b_mesh, BL::Object *b_ob, ParticleCurveData *CData, bool use_parents, bool background, int vcol_num)
 {
 #if 0
 	int keyno = 0;
@@ -342,12 +338,9 @@ bool ObtainCacheParticleVcol(Mesh *mesh, BL::Mesh *b_mesh, BL::Object *b_ob, Par
 #if 0
 				int mi = clamp(b_psys.settings().material()-1, 0, mesh->used_shaders.size()-1);
 				int shader = mesh->used_shaders[mi];
-				int draw_step = b_psys.settings().draw_step();
-				int ren_step = (int)pow((float)2.0f,(float)draw_step);
-				b_psys.settings().render_step(draw_step);
 #endif
 				int totparts = b_psys.particles.length();
-				int totchild = b_psys.child_particles.length() * b_psys.settings().draw_percentage() / 100;
+				int totchild = background ? b_psys.child_particles.length() : (int)((float)b_psys.child_particles.length() * (float)b_psys.settings().draw_percentage() / 100.0f);
 				int totcurves = totchild;
 				
 				if (use_parents || b_psys.settings().child_type() == 0)
@@ -386,6 +379,18 @@ bool ObtainCacheParticleVcol(Mesh *mesh, BL::Mesh *b_mesh, BL::Object *b_ob, Par
 
 }
 
+static void set_resolution(Mesh *mesh, BL::Mesh *b_mesh, BL::Object *b_ob, BL::Scene *scene, bool render)
+{
+	BL::Object::modifiers_iterator b_mod;
+	for(b_ob->modifiers.begin(b_mod); b_mod != b_ob->modifiers.end(); ++b_mod) {
+		if ((b_mod->type() == b_mod->type_PARTICLE_SYSTEM) && (b_mod->show_viewport()) && (b_mod->show_render())) {
+			BL::ParticleSystemModifier psmd((const PointerRNA)b_mod->ptr);
+			BL::ParticleSystem b_psys((const PointerRNA)psmd.particle_system().ptr);
+			b_psys.set_resolution(*scene, *b_ob, (render)? 2: 1);
+		}
+	}
+}
+
 void ExportCurveTrianglePlanes(Mesh *mesh, ParticleCurveData *CData, int interpolation, bool use_smooth, int segments, float3 RotCam)
 {
 	int vertexno = mesh->verts.size();
@@ -423,13 +428,13 @@ void ExportCurveTrianglePlanes(Mesh *mesh, ParticleCurveData *CData, int interpo
 
 					float radius = shaperadius(CData->psys_shape[sys], CData->psys_rootradius[sys], CData->psys_tipradius[sys], time);
 
-					if(CData->psys_closetip[sys] && (subv == segments) && (curvekey == CData->curve_firstkey[curve] + CData->curve_keynum[curve] - 1))
-						radius = shaperadius(CData->psys_shape[sys], CData->psys_rootradius[sys], 0.0f, 0.95f);
-
-					if((curvekey == CData->curve_firstkey[curve] + CData->curve_keynum[curve] - 1) && (subv == segments))
+					if((curvekey == CData->curve_firstkey[curve] + CData->curve_keynum[curve] - 2) && (subv == segments))
 						radius = shaperadius(CData->psys_shape[sys], CData->psys_rootradius[sys], CData->psys_tipradius[sys], 0.95f);
 
-					xbasis = normalize(cross(v1,RotCam - ickey_loc));
+					if(CData->psys_closetip[sys] && (subv == segments) && (curvekey == CData->curve_firstkey[curve] + CData->curve_keynum[curve] - 2))
+						radius = shaperadius(CData->psys_shape[sys], CData->psys_rootradius[sys], 0.0f, 0.95f);
+
+					xbasis = normalize(cross(RotCam - ickey_loc,v1));
 					float3 ickey_loc_shfl = ickey_loc - radius * xbasis;
 					float3 ickey_loc_shfr = ickey_loc + radius * xbasis;
 					mesh->verts.push_back(ickey_loc_shfl);
@@ -537,12 +542,12 @@ void ExportCurveTriangleRibbons(Mesh *mesh, ParticleCurveData *CData, int interp
 
 					float radius = shaperadius(CData->psys_shape[sys], CData->psys_rootradius[sys], CData->psys_tipradius[sys], time);
 
-					if(CData->psys_closetip[sys] && (subv == segments) && (curvekey == CData->curve_firstkey[curve] + CData->curve_keynum[curve] - 1))
-						radius = shaperadius(CData->psys_shape[sys], CData->psys_rootradius[sys], 0.0f, 0.95f);
-
-					if((curvekey == CData->curve_firstkey[curve] + CData->curve_keynum[curve] - 1) && (subv == segments))
+					if((curvekey == CData->curve_firstkey[curve] + CData->curve_keynum[curve] - 2) && (subv == segments))
 						radius = shaperadius(CData->psys_shape[sys], CData->psys_rootradius[sys], CData->psys_tipradius[sys], 0.95f);
 
+					if(CData->psys_closetip[sys] && (subv == segments) && (curvekey == CData->curve_firstkey[curve] + CData->curve_keynum[curve] - 2))
+						radius = shaperadius(CData->psys_shape[sys], CData->psys_rootradius[sys], 0.0f, 0.95f);
+
 					float3 ickey_loc_shfl = ickey_loc - radius * xbasis;
 					float3 ickey_loc_shfr = ickey_loc + radius * xbasis;
 					mesh->verts.push_back(ickey_loc_shfl);
@@ -654,12 +659,12 @@ void ExportCurveTriangleGeometry(Mesh *mesh, ParticleCurveData *CData, int inter
 
 					float radius = shaperadius(CData->psys_shape[sys], CData->psys_rootradius[sys], CData->psys_tipradius[sys], time);
 
-					if(CData->psys_closetip[sys] && (subv == segments) && (curvekey == CData->curve_firstkey[curve] + CData->curve_keynum[curve] - 1))
-						radius = shaperadius(CData->psys_shape[sys], CData->psys_rootradius[sys], 0.0f, 0.95f);
-
-					if((curvekey == CData->curve_firstkey[curve] + CData->curve_keynum[curve] - 1) && (subv == segments))
+					if((curvekey == CData->curve_firstkey[curve] + CData->curve_keynum[curve] - 2) && (subv == segments))
 						radius = shaperadius(CData->psys_shape[sys], CData->psys_rootradius[sys], CData->psys_tipradius[sys], 0.95f);
 
+					if(CData->psys_closetip[sys] && (subv == segments) && (curvekey == CData->curve_firstkey[curve] + CData->curve_keynum[curve] - 2))
+						radius = shaperadius(CData->psys_shape[sys], CData->psys_rootradius[sys], 0.0f, 0.95f);
+
 					float angle = 2 * M_PI_F / (float)resolution;
 					for(int section = 0 ; section < resolution; section++) {
 						float3 ickey_loc_shf = ickey_loc + radius * (cosf(angle * section) * xbasis + sinf(angle * section) * ybasis);
@@ -759,17 +764,14 @@ static void ExportCurveSegments(Scene *scene, Mesh *mesh, ParticleCurveData *CDa
 	}
 }
 
-void ExportCurveTriangleUVs(Mesh *mesh, ParticleCurveData *CData, int interpolation, bool use_smooth, int segments, int vert_offset, int resol)
+void ExportCurveTriangleUV(Mesh *mesh, ParticleCurveData *CData, int interpolation, int segments, int vert_offset, int resol, float3 *uvdata)
 {
+	if(uvdata == NULL)
+		return;
+
 	float time = 0.0f;
 	float prevtime = 0.0f;
 
-	Attribute *attr = mesh->attributes.find(ATTR_STD_UV);
-	if (attr == NULL)
-		return;
-		
-	float3 *uvdata = attr->data_float3();
-
 	int vertexindex = vert_offset;
 
 	for( int sys = 0; sys < CData->psys_firstcurve.size() ; sys++) {
@@ -818,6 +820,58 @@ void ExportCurveTriangleUVs(Mesh *mesh, ParticleCurveData *CData, int interpolat
 	}
 
 }
+
+void ExportCurveTriangleVcol(Mesh *mesh, ParticleCurveData *CData, int interpolation, int segments, int vert_offset, int resol, float3 *fdata)
+{
+	if(fdata == NULL)
+		return;
+
+	float time = 0.0f;
+//	float prevtime = 0.0f; // UNUSED
+
+	int vertexindex = vert_offset;
+
+	for( int sys = 0; sys < CData->psys_firstcurve.size() ; sys++) {
+		for( int curve = CData->psys_firstcurve[sys]; curve < CData->psys_firstcurve[sys] + CData->psys_curvenum[sys] ; curve++) {
+
+			for( int curvekey = CData->curve_firstkey[curve]; curvekey < CData->curve_firstkey[curve] + CData->curve_keynum[curve] - 1; curvekey++) {
+
+				int subv = 1;
+
+				if (curvekey == CData->curve_firstkey[curve])
+					subv = 0;
+				
+				for (; subv <= segments; subv++) {
+
+					float3 ickey_loc = make_float3(0.0f,0.0f,0.0f);
+
+					InterpolateKeySegments(subv, segments, curvekey, curve, &ickey_loc, &time, CData , interpolation);
+
+					if(subv!=0) {
+						for(int section = 0 ; section < resol; section++) {
+							fdata[vertexindex] = color_srgb_to_scene_linear(CData->curve_vcol[curve]);
+							vertexindex++;
+							fdata[vertexindex] = color_srgb_to_scene_linear(CData->curve_vcol[curve]);
+							vertexindex++;
+							fdata[vertexindex] = color_srgb_to_scene_linear(CData->curve_vcol[curve]);
+							vertexindex++;
+							fdata[vertexindex] = color_srgb_to_scene_linear(CData->curve_vcol[curve]);
+							vertexindex++;
+							fdata[vertexindex] = color_srgb_to_scene_linear(CData->curve_vcol[curve]);
+							vertexindex++;
+							fdata[vertexindex] = color_srgb_to_scene_linear(CData->curve_vcol[curve]);
+							vertexindex++;
+						}
+					}
+
+					// prevtime = time;  // UNUSED
+				}
+			}
+		}
+	}
+
+}
+
 /* Hair Curve Sync */
 
 void BlenderSync::sync_curve_settings()
@@ -863,6 +917,13 @@ void BlenderSync::sync_curve_settings()
 		curve_system_manager->use_joined = false;
 
 		switch(preset) {
+			case CURVE_FAST_PLANES:
+				/*camera facing planes*/
+				curve_system_manager->primitive = CURVE_TRIANGLES;
+				curve_system_manager->triangle_method = CURVE_CAMERA_TRIANGLES;
+				curve_system_manager->use_smooth = true;
+				curve_system_manager->resolution = 1;
+				break;
 			case CURVE_TANGENT_SHADING:
 				/*tangent shading*/
 				curve_system_manager->line_method = CURVE_UNCORRECTED;
@@ -890,6 +951,12 @@ void BlenderSync::sync_curve_settings()
 				curve_system_manager->use_tangent_normal_geometry = false;
 				curve_system_manager->use_tangent_normal_correction = false;
 				break;
+			case CURVE_SMOOTH_CURVES:
+				/*Cardinal curves preset*/
+				curve_system_manager->primitive = CURVE_SEGMENTS;
+				curve_system_manager->use_backfacing = true;
+				curve_system_manager->subdivisions = 4;
+				break;
 		}
 		
 	}
@@ -940,14 +1007,20 @@ void BlenderSync::sync_curves(Mesh *mesh, BL::Mesh b_mesh, BL::Object b_ob, bool
 	bool use_smooth = scene->curve_system_manager->use_smooth;
 	bool use_parents = scene->curve_system_manager->use_parents;
 	bool export_tgs = scene->curve_system_manager->use_joined;
+	size_t vert_num = mesh->verts.size();
+	size_t tri_num = mesh->triangles.size();
+	int used_res = 1;
 
 	/* extract particle hair data - should be combined with connecting to mesh later*/
 
 	ParticleCurveData CData;
 
-	ObtainCacheParticleData(mesh, &b_mesh, &b_ob, &CData, use_parents);
+	if(!preview)
+		set_resolution(mesh, &b_mesh, &b_ob, &b_scene, true);
+
+	ObtainCacheParticleData(mesh, &b_mesh, &b_ob, &CData, use_parents, !preview);
 
-	/* attach strands to mesh */
+	/* obtain camera parameters */
 	BL::Object b_CamOb = b_scene.camera();
 	float3 RotCam = make_float3(0.0f, 0.0f, 0.0f);
 	if(b_CamOb) {
@@ -957,22 +1030,16 @@ void BlenderSync::sync_curves(Mesh *mesh, BL::Mesh b_mesh, BL::Object b_ob, bool
 		RotCam = transform_point(&itfm, make_float3(ctfm.x.w, ctfm.y.w, ctfm.z.w));
 	}
 
+	/* add hair geometry to mesh */
 	if(primitive == CURVE_TRIANGLES){
-		int vert_num = mesh->triangles.size() * 3;
-		ObtainCacheParticleUV(mesh, &b_mesh, &b_ob, &CData, use_parents, 0);
-		if(triangle_method == CURVE_CAMERA_TRIANGLES) {
+		if(triangle_method == CURVE_CAMERA_TRIANGLES)
 			ExportCurveTrianglePlanes(mesh, &CData, interpolation, use_smooth, segments, RotCam);
-			ExportCurveTriangleUVs(mesh, &CData, interpolation, use_smooth, segments, vert_num, 1);
-		}
-		else if(triangle_method == CURVE_RIBBON_TRIANGLES) {
+		else if(triangle_method == CURVE_RIBBON_TRIANGLES)
 			ExportCurveTriangleRibbons(mesh, &CData, interpolation, use_smooth, segments);
-			ExportCurveTriangleUVs(mesh, &CData, interpolation, use_smooth, segments, vert_num, 1);
-		}
 		else {
 			ExportCurveTriangleGeometry(mesh, &CData, interpolation, use_smooth, resolution, segments);
-			ExportCurveTriangleUVs(mesh, &CData, interpolation, use_smooth, segments, vert_num, resolution);
+			used_res = resolution;
 		}
-
 	}
 	else {
 		ExportCurveSegments(scene, mesh, &CData, interpolation, segments);
@@ -990,14 +1057,24 @@ void BlenderSync::sync_curves(Mesh *mesh, BL::Mesh b_mesh, BL::Object b_ob, bool
 				data_tangent[ck] = tg;
 			}
 		}
+	}
+
 
-		/* generated coordinates from first key. we should ideally get this from
-		 * blender to handle deforming objects */
-		{
-			if(mesh->need_attribute(scene, ATTR_STD_GENERATED)) {
-				float3 loc, size;
-				mesh_texture_space(b_mesh, loc, size);
+	/* generated coordinates from first key. we should ideally get this from
+	 * blender to handle deforming objects */
+	{
+		if(mesh->need_attribute(scene, ATTR_STD_GENERATED)) {
+			float3 loc, size;
+			mesh_texture_space(b_mesh, loc, size);
 
+			if(primitive == CURVE_TRIANGLES) {
+				Attribute *attr_generated = mesh->attributes.add(ATTR_STD_GENERATED);
+				float3 *generated = attr_generated->data_float3();
+
+				for(size_t i = vert_num; i < mesh->verts.size(); i++)
+					generated[i] = mesh->verts[i]*size - loc;
+			}
+			else {
 				Attribute *attr_generated = mesh->curve_attributes.add(ATTR_STD_GENERATED);
 				float3 *generated = attr_generated->data_float3();
 				size_t i = 0;
@@ -1008,63 +1085,88 @@ void BlenderSync::sync_curves(Mesh *mesh, BL::Mesh b_mesh, BL::Object b_ob, bool
 				}
 			}
 		}
+	}
 
-		/* create vertex color attributes */
-		{
-			BL::Mesh::tessface_vertex_colors_iterator l;
-			int vcol_num = 0;
+	/* create vertex color attributes */
+	{
+		BL::Mesh::tessface_vertex_colors_iterator l;
+		int vcol_num = 0;
 
-			for(b_mesh.tessface_vertex_colors.begin(l); l != b_mesh.tessface_vertex_colors.end(); ++l, vcol_num++) {
-				if(!mesh->need_attribute(scene, ustring(l->name().c_str())))
-					continue;
+		for(b_mesh.tessface_vertex_colors.begin(l); l != b_mesh.tessface_vertex_colors.end(); ++l, vcol_num++) {
+			if(!mesh->need_attribute(scene, ustring(l->name().c_str())))
+				continue;
+
+			ObtainCacheParticleVcol(mesh, &b_mesh, &b_ob, &CData, use_parents, !preview, vcol_num);
+
+			if(primitive == CURVE_TRIANGLES) {
+
+				Attribute *attr_vcol = mesh->attributes.add(
+					ustring(l->name().c_str()), TypeDesc::TypeColor, ATTR_ELEMENT_CORNER);
+
+				float3 *fdata = attr_vcol->data_float3();
 
+				ExportCurveTriangleVcol(mesh, &CData, interpolation, segments, tri_num * 3, used_res, fdata);
+			}
+			else {
 				Attribute *attr_vcol = mesh->curve_attributes.add(
 					ustring(l->name().c_str()), TypeDesc::TypeColor, ATTR_ELEMENT_CURVE);
 
-				ObtainCacheParticleVcol(mesh, &b_mesh, &b_ob, &CData, use_parents, vcol_num);
-
-				float3 *vcol = attr_vcol->data_float3();
+				float3 *fdata = attr_vcol->data_float3();
 
-				if(vcol) {
+				if(fdata) {
 					for(size_t curve = 0; curve < CData.curve_vcol.size() ;curve++)
-						vcol[curve] = color_srgb_to_scene_linear(CData.curve_vcol[curve]);
+						fdata[curve] = color_srgb_to_scene_linear(CData.curve_vcol[curve]);
 				}
 			}
 		}
+	}
 
-		/* create uv map attributes */
-		{
-			BL::Mesh::tessface_uv_textures_iterator l;
-			int uv_num = 0;
+	/* create UV attributes */
+	{
+		BL::Mesh::tessface_uv_textures_iterator l;
+		int uv_num = 0;
 
-			for(b_mesh.tessface_uv_textures.begin(l); l != b_mesh.tessface_uv_textures.end(); ++l, uv_num++) {
-				bool active_render = l->active_render();
-				AttributeStandard std = (active_render)? ATTR_STD_UV: ATTR_STD_NONE;
-				ustring name = ustring(l->name().c_str());
+		for(b_mesh.tessface_uv_textures.begin(l); l != b_mesh.tessface_uv_textures.end(); ++l, uv_num++) {
+			bool active_render = l->active_render();
+			AttributeStandard std = (active_render)? ATTR_STD_UV: ATTR_STD_NONE;
+			ustring name = ustring(l->name().c_str());
 
-				/* UV map */
-				if(mesh->need_attribute(scene, name) || mesh->need_attribute(scene, std)) {
-					Attribute *attr;
+			/* UV map */
+			if(mesh->need_attribute(scene, name) || mesh->need_attribute(scene, std)) {
+				Attribute *attr_uv;
 
+				ObtainCacheParticleUV(mesh, &b_mesh, &b_ob, &CData, use_parents, !preview, uv_num);
+
+				if(primitive == CURVE_TRIANGLES) {
 					if(active_render)
-						attr = mesh->curve_attributes.add(std, name);
+						attr_uv = mesh->attributes.add(std, name);
 					else
-						attr = mesh->curve_attributes.add(name, TypeDesc::TypePoint,  ATTR_ELEMENT_CURVE);
+						attr_uv = mesh->attributes.add(name, TypeDesc::TypePoint, ATTR_ELEMENT_CORNER);
 
-					ObtainCacheParticleUV(mesh, &b_mesh, &b_ob, &CData, use_parents, uv_num);
+					float3 *uv = attr_uv->data_float3();
 
-					float3 *uv = attr->data_float3();
+					ExportCurveTriangleUV(mesh, &CData, interpolation, segments, tri_num * 3, used_res, uv);
+				}
+				else {
+					if(active_render)
+						attr_uv = mesh->curve_attributes.add(std, name);
+					else
+						attr_uv = mesh->curve_attributes.add(name, TypeDesc::TypePoint,  ATTR_ELEMENT_CURVE);
+
+					float3 *uv = attr_uv->data_float3();
 
 					if(uv) {
-						for(size_t curve = 0; curve < CData.curve_uv.size() ;curve++)
+						for(size_t curve = 0; curve < CData.curve_uv.size(); curve++)
 							uv[curve] = CData.curve_uv[curve];
 					}
-				}			
+				}
 			}
 		}
-
 	}
 
+	if(!preview)
+		set_resolution(mesh, &b_mesh, &b_ob, &b_scene, false);
+
 	mesh->compute_bounds();
 }
 
diff --git a/intern/cycles/blender/blender_mesh.cpp b/intern/cycles/blender/blender_mesh.cpp
index 1dd7800dfa4..c61a7fafb57 100644
--- a/intern/cycles/blender/blender_mesh.cpp
+++ b/intern/cycles/blender/blender_mesh.cpp
@@ -313,7 +313,7 @@ static void create_mesh(Scene *scene, Mesh *mesh, BL::Mesh b_mesh, const vector<
 			std = (active_render)? ATTR_STD_UV_TANGENT: ATTR_STD_NONE;
 			name = ustring((string(l->name().c_str()) + ".tangent").c_str());
 
-			if(mesh->need_attribute(scene, name) || mesh->need_attribute(scene, std)) {
+			if(mesh->need_attribute(scene, name) || (active_render && mesh->need_attribute(scene, std))) {
 				std = (active_render)? ATTR_STD_UV_TANGENT_SIGN: ATTR_STD_NONE;
 				name = ustring((string(l->name().c_str()) + ".tangent_sign").c_str());
 				bool need_sign = (mesh->need_attribute(scene, name) || mesh->need_attribute(scene, std));
@@ -435,7 +435,7 @@ Mesh *BlenderSync::sync_mesh(BL::Object b_ob, bool object_updated, bool hide_tri
 	mesh_synced.insert(mesh);
 
 	/* create derived mesh */
-	BL::Mesh b_mesh = object_to_mesh(b_ob, b_scene, true, !preview);
+	BL::Mesh b_mesh = object_to_mesh(b_data, b_ob, b_scene, true, !preview);
 	PointerRNA cmesh = RNA_pointer_get(&b_ob_data.ptr, "cycles");
 
 	vector<Mesh::Triangle> oldtriangle = mesh->triangles;
@@ -507,7 +507,7 @@ void BlenderSync::sync_mesh_motion(BL::Object b_ob, Mesh *mesh, int motion)
 		return;
 
 	/* get derived mesh */
-	BL::Mesh b_mesh = object_to_mesh(b_ob, b_scene, true, !preview);
+	BL::Mesh b_mesh = object_to_mesh(b_data, b_ob, b_scene, true, !preview);
 
 	if(b_mesh) {
 		BL::Mesh::vertices_iterator v;
diff --git a/intern/cycles/blender/blender_object.cpp b/intern/cycles/blender/blender_object.cpp
index e9bcea70ab6..a9d37dbed99 100644
--- a/intern/cycles/blender/blender_object.cpp
+++ b/intern/cycles/blender/blender_object.cpp
@@ -156,6 +156,7 @@ void BlenderSync::sync_light(BL::Object b_parent, int persistent_id[OBJECT_PERSI
 	/* shadow */
 	PointerRNA clamp = RNA_pointer_get(&b_lamp.ptr, "cycles");
 	light->cast_shadow = get_boolean(clamp, "cast_shadow");
+	light->use_mis = get_boolean(clamp, "use_multiple_importance_sampling");
 	light->samples = get_int(clamp, "samples");
 
 	/* tag */
diff --git a/intern/cycles/blender/blender_python.cpp b/intern/cycles/blender/blender_python.cpp
index 676fba76ddf..a10f3b63033 100644
--- a/intern/cycles/blender/blender_python.cpp
+++ b/intern/cycles/blender/blender_python.cpp
@@ -54,7 +54,7 @@ static PyObject *create_func(PyObject *self, PyObject *args)
 	PyObject *pyengine, *pyuserpref, *pydata, *pyscene, *pyregion, *pyv3d, *pyrv3d;
 	int preview_osl;
 
-	if(!PyArg_ParseTuple(args, "OOOOOOOp", &pyengine, &pyuserpref, &pydata, &pyscene, &pyregion, &pyv3d, &pyrv3d, &preview_osl))
+	if(!PyArg_ParseTuple(args, "OOOOOOOi", &pyengine, &pyuserpref, &pydata, &pyscene, &pyregion, &pyv3d, &pyrv3d, &preview_osl))
 		return NULL;
 
 	/* RNA */
diff --git a/intern/cycles/blender/blender_session.cpp b/intern/cycles/blender/blender_session.cpp
index 650d3d387ee..951f7022a10 100644
--- a/intern/cycles/blender/blender_session.cpp
+++ b/intern/cycles/blender/blender_session.cpp
@@ -111,9 +111,9 @@ void BlenderSession::create_session()
 	b_engine.use_highlight_tiles(session_params.progressive_refine == false);
 
 	/* setup callbacks for builtin image support */
-	scene->image_manager->builtin_image_info_cb = function_bind(&BlenderSession::builtin_image_info, this, _1, _2, _3, _4, _5);
-	scene->image_manager->builtin_image_pixels_cb = function_bind(&BlenderSession::builtin_image_pixels, this, _1, _2);
-	scene->image_manager->builtin_image_float_pixels_cb = function_bind(&BlenderSession::builtin_image_float_pixels, this, _1, _2);
+	scene->image_manager->builtin_image_info_cb = function_bind(&BlenderSession::builtin_image_info, this, _1, _2, _3, _4, _5, _6);
+	scene->image_manager->builtin_image_pixels_cb = function_bind(&BlenderSession::builtin_image_pixels, this, _1, _2, _3);
+	scene->image_manager->builtin_image_float_pixels_cb = function_bind(&BlenderSession::builtin_image_float_pixels, this, _1, _2, _3);
 }
 
 void BlenderSession::reset_session(BL::BlendData b_data_, BL::Scene b_scene_)
@@ -493,11 +493,10 @@ bool BlenderSession::draw(int w, int h)
 		}
 		else {
 			/* update camera from 3d view */
-			bool need_update = scene->camera->need_update;
 
 			sync->sync_view(b_v3d, b_rv3d, w, h);
 
-			if(scene->camera->need_update && !need_update)
+			if(scene->camera->need_update)
 				reset = true;
 
 			session->scene->mutex.unlock();
@@ -622,23 +621,19 @@ void BlenderSession::test_cancel()
 /* builtin image file name is actually an image datablock name with
  * absolute sequence frame number concatenated via '@' character
  *
- * this function splits image id name and frame number from a
- * builtin image name
+ * this function splits frame from builtin name
  */
-void BlenderSession::builtin_name_split(const string &builtin_name, string &name, int &frame)
+int BlenderSession::builtin_image_frame(const string &builtin_name)
 {
 	int last = builtin_name.find_last_of('@');
-	name = builtin_name.substr(0, last);
-	frame = atoi(builtin_name.substr(last + 1, builtin_name.size() - last - 1).c_str());
+	return atoi(builtin_name.substr(last + 1, builtin_name.size() - last - 1).c_str());
 }
 
-void BlenderSession::builtin_image_info(const string &builtin_name, bool &is_float, int &width, int &height, int &channels)
+void BlenderSession::builtin_image_info(const string &builtin_name, void *builtin_data, bool &is_float, int &width, int &height, int &channels)
 {
-	string name;
-	int frame;
-	builtin_name_split(builtin_name, name, frame);
-
-	BL::Image b_image = b_data.images[name];
+	PointerRNA ptr;
+	RNA_id_pointer_create((ID*)builtin_data, &ptr);
+	BL::Image b_image(ptr);
 
 	if(b_image) {
 		is_float = b_image.is_float();
@@ -654,13 +649,13 @@ void BlenderSession::builtin_image_info(const string &builtin_name, bool &is_flo
 	}
 }
 
-bool BlenderSession::builtin_image_pixels(const string &builtin_name, unsigned char *pixels)
+bool BlenderSession::builtin_image_pixels(const string &builtin_name, void *builtin_data, unsigned char *pixels)
 {
-	string name;
-	int frame;
-	builtin_name_split(builtin_name, name, frame);
+	int frame = builtin_image_frame(builtin_name);
 
-	BL::Image b_image = b_data.images[name];
+	PointerRNA ptr;
+	RNA_id_pointer_create((ID*)builtin_data, &ptr);
+	BL::Image b_image(ptr);
 
 	if(b_image) {
 		int width = b_image.size()[0];
@@ -696,13 +691,13 @@ bool BlenderSession::builtin_image_pixels(const string &builtin_name, unsigned c
 	return false;
 }
 
-bool BlenderSession::builtin_image_float_pixels(const string &builtin_name, float *pixels)
+bool BlenderSession::builtin_image_float_pixels(const string &builtin_name, void *builtin_data, float *pixels)
 {
-	string name;
-	int frame;
-	builtin_name_split(builtin_name, name, frame);
+	int frame = builtin_image_frame(builtin_name);
 
-	BL::Image b_image = b_data.images[name];
+	PointerRNA ptr;
+	RNA_id_pointer_create((ID*)builtin_data, &ptr);
+	BL::Image b_image(ptr);
 
 	if(b_image) {
 		int width = b_image.size()[0];
diff --git a/intern/cycles/blender/blender_session.h b/intern/cycles/blender/blender_session.h
index 686ff3d1be9..0210df9c044 100644
--- a/intern/cycles/blender/blender_session.h
+++ b/intern/cycles/blender/blender_session.h
@@ -94,10 +94,10 @@ protected:
 	void do_write_update_render_result(BL::RenderResult b_rr, BL::RenderLayer b_rlay, RenderTile& rtile, bool do_update_only);
 	void do_write_update_render_tile(RenderTile& rtile, bool do_update_only);
 
-	void builtin_name_split(const string &builtin_name, string &name, int &frame);
-	void builtin_image_info(const string &builtin_name, bool &is_float, int &width, int &height, int &channels);
-	bool builtin_image_pixels(const string &builtin_name, unsigned char *pixels);
-	bool builtin_image_float_pixels(const string &builtin_name, float *pixels);
+	int builtin_image_frame(const string &builtin_name);
+	void builtin_image_info(const string &builtin_name, void *builtin_data, bool &is_float, int &width, int &height, int &channels);
+	bool builtin_image_pixels(const string &builtin_name, void *builtin_data, unsigned char *pixels);
+	bool builtin_image_float_pixels(const string &builtin_name, void *builtin_data, float *pixels);
 };
 
 CCL_NAMESPACE_END
diff --git a/intern/cycles/blender/blender_shader.cpp b/intern/cycles/blender/blender_shader.cpp
index b1eaedba47a..0f7dc15db19 100644
--- a/intern/cycles/blender/blender_shader.cpp
+++ b/intern/cycles/blender/blender_shader.cpp
@@ -529,11 +529,11 @@ static ShaderNode *add_node(Scene *scene, BL::BlendData b_data, BL::Scene b_scen
 					int scene_frame = b_scene.frame_current();
 					int image_frame = image_user_frame_number(b_image_node.image_user(), scene_frame);
 					image->filename = b_image.name() + "@" + string_printf("%d", image_frame);
-					image->is_builtin = true;
+					image->builtin_data = b_image.ptr.data;
 				}
 				else {
 					image->filename = image_user_file_path(b_image_node.image_user(), b_image, b_scene.frame_current());
-					image->is_builtin = false;
+					image->builtin_data = NULL;
 				}
 
 				image->animated = b_image_node.image_user().use_auto_refresh();
@@ -558,11 +558,12 @@ static ShaderNode *add_node(Scene *scene, BL::BlendData b_data, BL::Scene b_scen
 					int scene_frame = b_scene.frame_current();
 					int image_frame = image_user_frame_number(b_env_node.image_user(), scene_frame);
 					env->filename = b_image.name() + "@" + string_printf("%d", image_frame);
-					env->is_builtin = true;
+					env->builtin_data = b_image.ptr.data;
 				}
 				else {
 					env->filename = image_user_file_path(b_env_node.image_user(), b_image, b_scene.frame_current());
 					env->animated = b_env_node.image_user().use_auto_refresh();
+					env->builtin_data = NULL;
 				}
 			}
 			env->color_space = EnvironmentTextureNode::color_space_enum[(int)b_env_node.color_space()];
diff --git a/intern/cycles/blender/blender_util.h b/intern/cycles/blender/blender_util.h
index f134416f2d0..0947b5a2e9d 100644
--- a/intern/cycles/blender/blender_util.h
+++ b/intern/cycles/blender/blender_util.h
@@ -39,9 +39,9 @@ float *BKE_image_get_float_pixels_for_frame(void *image, int frame);
 
 CCL_NAMESPACE_BEGIN
 
-static inline BL::Mesh object_to_mesh(BL::Object self, BL::Scene scene, bool apply_modifiers, bool render)
+static inline BL::Mesh object_to_mesh(BL::BlendData data, BL::Object object, BL::Scene scene, bool apply_modifiers, bool render)
 {
-	return self.to_mesh(scene, apply_modifiers, (render)? 2: 1);
+	return data.meshes.new_from_object(scene, object, apply_modifiers, (render)? 2: 1);
 }
 
 static inline void colorramp_to_array(BL::ColorRamp ramp, float4 *data, int size)
diff --git a/intern/cycles/device/CMakeLists.txt b/intern/cycles/device/CMakeLists.txt
index 661d43ab036..fe2368b7ea8 100644
--- a/intern/cycles/device/CMakeLists.txt
+++ b/intern/cycles/device/CMakeLists.txt
@@ -23,7 +23,9 @@ set(SRC
 )
 
 if(WITH_NETWORK)
-	list(APPEND SRC device_network.cpp)
+	list(APPEND SRC
+		device_network.cpp
+	)
 endif()
 
 set(SRC_HEADERS
diff --git a/intern/cycles/device/device_cpu.cpp b/intern/cycles/device/device_cpu.cpp
index a1d7706a34e..1915245bb55 100644
--- a/intern/cycles/device/device_cpu.cpp
+++ b/intern/cycles/device/device_cpu.cpp
@@ -58,7 +58,8 @@ public:
 #endif
 
 		/* do now to avoid thread issues */
-		system_cpu_support_optimized();
+		system_cpu_support_sse2();
+		system_cpu_support_sse3();
 	}
 
 	~CPUDevice()
@@ -170,7 +171,7 @@ public:
 			int end_sample = tile.start_sample + tile.num_samples;
 
 #ifdef WITH_OPTIMIZED_KERNEL
-			if(system_cpu_support_optimized()) {
+			if(system_cpu_support_sse2()) {
 				for(int sample = start_sample; sample < end_sample; sample++) {
 					if (task.get_cancel() || task_pool.cancelled()) {
 						if(task.need_finish_queue == false)
@@ -179,7 +180,26 @@ public:
 
 					for(int y = tile.y; y < tile.y + tile.h; y++) {
 						for(int x = tile.x; x < tile.x + tile.w; x++) {
-							kernel_cpu_optimized_path_trace(&kg, render_buffer, rng_state,
+							kernel_cpu_sse2_path_trace(&kg, render_buffer, rng_state,
+								sample, x, y, tile.offset, tile.stride);
+						}
+					}
+
+					tile.sample = sample + 1;
+
+					task.update_progress(tile);
+				}
+			}
+			else if(system_cpu_support_sse3()) {
+				for(int sample = start_sample; sample < end_sample; sample++) {
+					if (task.get_cancel() || task_pool.cancelled()) {
+						if(task.need_finish_queue == false)
+							break;
+					}
+
+					for(int y = tile.y; y < tile.y + tile.h; y++) {
+						for(int x = tile.x; x < tile.x + tile.w; x++) {
+							kernel_cpu_sse3_path_trace(&kg, render_buffer, rng_state,
 								sample, x, y, tile.offset, tile.stride);
 						}
 					}
@@ -227,10 +247,16 @@ public:
 	void thread_tonemap(DeviceTask& task)
 	{
 #ifdef WITH_OPTIMIZED_KERNEL
-		if(system_cpu_support_optimized()) {
+		if(system_cpu_support_sse2()) {
+			for(int y = task.y; y < task.y + task.h; y++)
+				for(int x = task.x; x < task.x + task.w; x++)
+					kernel_cpu_sse2_tonemap(&kernel_globals, (uchar4*)task.rgba, (float*)task.buffer,
+						task.sample, task.resolution, x, y, task.offset, task.stride);
+		}
+		else if(system_cpu_support_sse3()) {
 			for(int y = task.y; y < task.y + task.h; y++)
 				for(int x = task.x; x < task.x + task.w; x++)
-					kernel_cpu_optimized_tonemap(&kernel_globals, (uchar4*)task.rgba, (float*)task.buffer,
+					kernel_cpu_sse3_tonemap(&kernel_globals, (uchar4*)task.rgba, (float*)task.buffer,
 						task.sample, task.resolution, x, y, task.offset, task.stride);
 		}
 		else
@@ -252,9 +278,17 @@ public:
 #endif
 
 #ifdef WITH_OPTIMIZED_KERNEL
-		if(system_cpu_support_optimized()) {
+		if(system_cpu_support_sse2()) {
+			for(int x = task.shader_x; x < task.shader_x + task.shader_w; x++) {
+				kernel_cpu_sse2_shader(&kg, (uint4*)task.shader_input, (float4*)task.shader_output, task.shader_eval_type, x);
+
+				if(task_pool.cancelled())
+					break;
+			}
+		}
+		else if(system_cpu_support_sse3()) {
 			for(int x = task.shader_x; x < task.shader_x + task.shader_w; x++) {
-				kernel_cpu_optimized_shader(&kg, (uint4*)task.shader_input, (float4*)task.shader_output, task.shader_eval_type, x);
+				kernel_cpu_sse3_shader(&kg, (uint4*)task.shader_input, (float4*)task.shader_output, task.shader_eval_type, x);
 
 				if(task_pool.cancelled())
 					break;
diff --git a/intern/cycles/kernel/CMakeLists.txt b/intern/cycles/kernel/CMakeLists.txt
index 6d5b9a063a0..e83756b7c8a 100644
--- a/intern/cycles/kernel/CMakeLists.txt
+++ b/intern/cycles/kernel/CMakeLists.txt
@@ -12,7 +12,8 @@ set(INC_SYS
 
 set(SRC
 	kernel.cpp
-	kernel_optimized.cpp
+	kernel_sse2.cpp
+	kernel_sse3.cpp
 	kernel.cl
 	kernel.cu
 )
@@ -149,7 +150,8 @@ include_directories(SYSTEM ${INC_SYS})
 add_library(cycles_kernel ${SRC} ${SRC_HEADERS} ${SRC_CLOSURE_HEADERS} ${SRC_SVM_HEADERS})
 
 if(WITH_CYCLES_OPTIMIZED_KERNEL)
-	set_source_files_properties(kernel_optimized.cpp PROPERTIES COMPILE_FLAGS "${CYCLES_OPTIMIZED_KERNEL_FLAGS}")
+	set_source_files_properties(kernel_sse2.cpp PROPERTIES COMPILE_FLAGS "${CYCLES_SSE2_KERNEL_FLAGS}")
+	set_source_files_properties(kernel_sse3.cpp PROPERTIES COMPILE_FLAGS "${CYCLES_SSE3_KERNEL_FLAGS}")
 endif()
 
 if(WITH_CYCLES_CUDA)
diff --git a/intern/cycles/kernel/kernel.h b/intern/cycles/kernel/kernel.h
index 26c0bcd6d1a..20ea5a61906 100644
--- a/intern/cycles/kernel/kernel.h
+++ b/intern/cycles/kernel/kernel.h
@@ -44,11 +44,18 @@ void kernel_cpu_shader(KernelGlobals *kg, uint4 *input, float4 *output,
 	int type, int i);
 
 #ifdef WITH_OPTIMIZED_KERNEL
-void kernel_cpu_optimized_path_trace(KernelGlobals *kg, float *buffer, unsigned int *rng_state,
+void kernel_cpu_sse2_path_trace(KernelGlobals *kg, float *buffer, unsigned int *rng_state,
 	int sample, int x, int y, int offset, int stride);
-void kernel_cpu_optimized_tonemap(KernelGlobals *kg, uchar4 *rgba, float *buffer,
+void kernel_cpu_sse2_tonemap(KernelGlobals *kg, uchar4 *rgba, float *buffer,
 	int sample, int resolution, int x, int y, int offset, int stride);
-void kernel_cpu_optimized_shader(KernelGlobals *kg, uint4 *input, float4 *output,
+void kernel_cpu_sse2_shader(KernelGlobals *kg, uint4 *input, float4 *output,
+	int type, int i);
+
+void kernel_cpu_sse3_path_trace(KernelGlobals *kg, float *buffer, unsigned int *rng_state,
+	int sample, int x, int y, int offset, int stride);
+void kernel_cpu_sse3_tonemap(KernelGlobals *kg, uchar4 *rgba, float *buffer,
+	int sample, int resolution, int x, int y, int offset, int stride);
+void kernel_cpu_sse3_shader(KernelGlobals *kg, uint4 *input, float4 *output,
 	int type, int i);
 #endif
 
diff --git a/intern/cycles/kernel/kernel_accumulate.h b/intern/cycles/kernel/kernel_accumulate.h
index 86177301357..1cf230634fc 100644
--- a/intern/cycles/kernel/kernel_accumulate.h
+++ b/intern/cycles/kernel/kernel_accumulate.h
@@ -220,7 +220,7 @@ __device_inline void path_radiance_accum_ao(PathRadiance *L, float3 throughput,
 #endif
 }
 
-__device_inline void path_radiance_accum_light(PathRadiance *L, float3 throughput, BsdfEval *bsdf_eval, float3 shadow, int bounce, bool is_lamp)
+__device_inline void path_radiance_accum_light(PathRadiance *L, float3 throughput, BsdfEval *bsdf_eval, float3 shadow, float shadow_fac, int bounce, bool is_lamp)
 {
 #ifdef __PASSES__
 	if(L->use_light_pass) {
@@ -233,9 +233,9 @@ __device_inline void path_radiance_accum_light(PathRadiance *L, float3 throughpu
 			if(is_lamp) {
 				float3 sum = throughput*(bsdf_eval->diffuse + bsdf_eval->glossy + bsdf_eval->transmission);
 
-				L->shadow.x += shadow.x;
-				L->shadow.y += shadow.y;
-				L->shadow.z += shadow.z;
+				L->shadow.x += shadow.x*shadow_fac;
+				L->shadow.y += shadow.y*shadow_fac;
+				L->shadow.z += shadow.z*shadow_fac;
 				L->shadow.w += average(sum);
 			}
 		}
diff --git a/intern/cycles/kernel/kernel_bvh.h b/intern/cycles/kernel/kernel_bvh.h
index 1a85b5bbefd..2b9ebf35d0c 100644
--- a/intern/cycles/kernel/kernel_bvh.h
+++ b/intern/cycles/kernel/kernel_bvh.h
@@ -126,21 +126,21 @@ __device_inline void bvh_node_intersect(KernelGlobals *kg,
 
 	/* intersect ray against child nodes */
 	float3 ood = P * idir;
-	float c0lox = n0xy.x * idir.x - ood.x;
-	float c0hix = n0xy.y * idir.x - ood.x;
-	float c0loy = n0xy.z * idir.y - ood.y;
-	float c0hiy = n0xy.w * idir.y - ood.y;
-	float c0loz = nz.x * idir.z - ood.z;
-	float c0hiz = nz.y * idir.z - ood.z;
+	NO_EXTENDED_PRECISION float c0lox = n0xy.x * idir.x - ood.x;
+	NO_EXTENDED_PRECISION float c0hix = n0xy.y * idir.x - ood.x;
+	NO_EXTENDED_PRECISION float c0loy = n0xy.z * idir.y - ood.y;
+	NO_EXTENDED_PRECISION float c0hiy = n0xy.w * idir.y - ood.y;
+	NO_EXTENDED_PRECISION float c0loz = nz.x * idir.z - ood.z;
+	NO_EXTENDED_PRECISION float c0hiz = nz.y * idir.z - ood.z;
 	NO_EXTENDED_PRECISION float c0min = max4(min(c0lox, c0hix), min(c0loy, c0hiy), min(c0loz, c0hiz), 0.0f);
 	NO_EXTENDED_PRECISION float c0max = min4(max(c0lox, c0hix), max(c0loy, c0hiy), max(c0loz, c0hiz), t);
 
-	float c1loz = nz.z * idir.z - ood.z;
-	float c1hiz = nz.w * idir.z - ood.z;
-	float c1lox = n1xy.x * idir.x - ood.x;
-	float c1hix = n1xy.y * idir.x - ood.x;
-	float c1loy = n1xy.z * idir.y - ood.y;
-	float c1hiy = n1xy.w * idir.y - ood.y;
+	NO_EXTENDED_PRECISION float c1loz = nz.z * idir.z - ood.z;
+	NO_EXTENDED_PRECISION float c1hiz = nz.w * idir.z - ood.z;
+	NO_EXTENDED_PRECISION float c1lox = n1xy.x * idir.x - ood.x;
+	NO_EXTENDED_PRECISION float c1hix = n1xy.y * idir.x - ood.x;
+	NO_EXTENDED_PRECISION float c1loy = n1xy.z * idir.y - ood.y;
+	NO_EXTENDED_PRECISION float c1hiy = n1xy.w * idir.y - ood.y;
 	NO_EXTENDED_PRECISION float c1min = max4(min(c1lox, c1hix), min(c1loy, c1hiy), min(c1loz, c1hiz), 0.0f);
 	NO_EXTENDED_PRECISION float c1max = min4(max(c1lox, c1hix), max(c1loy, c1hiy), max(c1loz, c1hiz), t);
 
diff --git a/intern/cycles/kernel/kernel_emission.h b/intern/cycles/kernel/kernel_emission.h
index 54bc0717b60..4048bbd9dfc 100644
--- a/intern/cycles/kernel/kernel_emission.h
+++ b/intern/cycles/kernel/kernel_emission.h
@@ -76,7 +76,7 @@ __device float3 direct_emissive_eval(KernelGlobals *kg, float rando,
 
 __device bool direct_emission(KernelGlobals *kg, ShaderData *sd, int lindex,
 	float randt, float rando, float randu, float randv, Ray *ray, BsdfEval *eval,
-	int *lamp)
+	bool *is_lamp)
 {
 	LightSample ls;
 
@@ -92,12 +92,6 @@ __device bool direct_emission(KernelGlobals *kg, ShaderData *sd, int lindex,
 		light_sample(kg, randt, randu, randv, sd->time, sd->P, &ls);
 	}
 
-	/* return lamp index for MIS */
-	if(ls.use_mis)
-		*lamp = ls.lamp;
-	else
-		*lamp= ~0;
-
 	if(ls.pdf == 0.0f)
 		return false;
 
@@ -114,7 +108,7 @@ __device bool direct_emission(KernelGlobals *kg, ShaderData *sd, int lindex,
 
 	shader_bsdf_eval(kg, sd, ls.D, eval, &bsdf_pdf);
 
-	if(ls.use_mis) {
+	if(ls.shader & SHADER_USE_MIS) {
 		/* multiple importance sampling */
 		float mis_weight = power_heuristic(ls.pdf, bsdf_pdf);
 		light_eval *= mis_weight;
@@ -146,6 +140,9 @@ __device bool direct_emission(KernelGlobals *kg, ShaderData *sd, int lindex,
 		ray->t = 0.0f;
 	}
 
+	/* return if it's a lamp for shadow pass */
+	*is_lamp = (ls.prim == ~0);
+
 	return true;
 }
 
diff --git a/intern/cycles/kernel/kernel_light.h b/intern/cycles/kernel/kernel_light.h
index 6ba3e439329..10a32226f17 100644
--- a/intern/cycles/kernel/kernel_light.h
+++ b/intern/cycles/kernel/kernel_light.h
@@ -31,7 +31,6 @@ typedef struct LightSample {
 	int prim;			/* primitive id for triangle/curve ligths */
 	int shader;			/* shader id */
 	int lamp;			/* lamp id */
-	int use_mis;		/* for lamps with size zero */
 	LightType type;		/* type of light */
 } LightSample;
 
@@ -218,11 +217,10 @@ __device void lamp_light_sample(KernelGlobals *kg, int lamp,
 
 	LightType type = (LightType)__float_as_int(data0.x);
 	ls->type = type;
-#ifdef __LAMP_MIS__
-	ls->use_mis = true;
-#else
-	ls->use_mis = false;
-#endif
+	ls->shader = __float_as_int(data1.x);
+	ls->object = ~0;
+	ls->prim = ~0;
+	ls->lamp = lamp;
 
 	if(type == LIGHT_DISTANT) {
 		/* distant light */
@@ -233,10 +231,6 @@ __device void lamp_light_sample(KernelGlobals *kg, int lamp,
 
 		if(radius > 0.0f)
 			D = distant_light_sample(D, radius, randu, randv);
-#ifdef __LAMP_MIS__
-		else
-			ls->use_mis = false;
-#endif
 
 		ls->P = D;
 		ls->Ng = D;
@@ -257,9 +251,6 @@ __device void lamp_light_sample(KernelGlobals *kg, int lamp,
 		ls->D = -D;
 		ls->t = FLT_MAX;
 		ls->eval_fac = 1.0f;
-#ifndef __LAMP_MIS__
-		ls->use_mis = true;
-#endif
 	}
 #endif
 	else {
@@ -271,10 +262,6 @@ __device void lamp_light_sample(KernelGlobals *kg, int lamp,
 			if(radius > 0.0f)
 				/* sphere light */
 				ls->P += sphere_light_sample(P, ls->P, radius, randu, randv);
-#ifdef __LAMP_MIS__
-			else
-				ls->use_mis = false;
-#endif
 
 			ls->D = normalize_len(ls->P - P, &ls->t);
 			ls->Ng = -ls->D;
@@ -304,13 +291,6 @@ __device void lamp_light_sample(KernelGlobals *kg, int lamp,
 
 			float invarea = data2.x;
 
-			if(invarea == 0.0f) {
-#ifdef __LAMP_MIS__
-				ls->use_mis = false;
-#endif
-				invarea = 1.0f;
-			}
-
 			ls->eval_fac = 0.25f*invarea;
 			ls->pdf = invarea;
 		}
@@ -318,11 +298,6 @@ __device void lamp_light_sample(KernelGlobals *kg, int lamp,
 		ls->eval_fac *= kernel_data.integrator.inv_pdf_lights;
 		ls->pdf *= lamp_light_pdf(kg, ls->Ng, -ls->D, ls->t);
 	}
-
-	ls->shader = __float_as_int(data1.x);
-	ls->object = ~0;
-	ls->prim = ~0;
-	ls->lamp = lamp;
 }
 
 __device bool lamp_light_eval(KernelGlobals *kg, int lamp, float3 P, float3 D, float t, LightSample *ls)
@@ -336,7 +311,9 @@ __device bool lamp_light_eval(KernelGlobals *kg, int lamp, float3 P, float3 D, f
 	ls->object = ~0;
 	ls->prim = ~0;
 	ls->lamp = lamp;
-	ls->use_mis = false; /* flag not used for eval */
+
+	if(!(ls->shader & SHADER_USE_MIS))
+		return false;
 
 	if(type == LIGHT_DISTANT) {
 		/* distant light */
@@ -475,7 +452,7 @@ __device void triangle_light_sample(KernelGlobals *kg, int prim, int object,
 	ls->object = object;
 	ls->prim = prim;
 	ls->lamp = ~0;
-	ls->use_mis = true;
+	ls->shader |= SHADER_USE_MIS;
 	ls->t = 0.0f;
 	ls->type = LIGHT_AREA;
 	ls->eval_fac = 1.0f;
@@ -529,11 +506,10 @@ __device void curve_segment_light_sample(KernelGlobals *kg, int prim, int object
 	ls->object = object;
 	ls->prim = prim;
 	ls->lamp = ~0;
-	ls->use_mis = true;
 	ls->t = 0.0f;
 	ls->type = LIGHT_STRAND;
 	ls->eval_fac = 1.0f;
-	ls->shader = __float_as_int(v00.z);
+	ls->shader = __float_as_int(v00.z) | SHADER_USE_MIS;
 
 	object_transform_light_sample(kg, ls, object, time);
 }
diff --git a/intern/cycles/kernel/kernel_path.h b/intern/cycles/kernel/kernel_path.h
index 1a5df66e6c2..9b13f4d2a42 100644
--- a/intern/cycles/kernel/kernel_path.h
+++ b/intern/cycles/kernel/kernel_path.h
@@ -254,7 +254,7 @@ __device float4 kernel_path_progressive(KernelGlobals *kg, RNG *rng, int sample,
 		bool hit = scene_intersect(kg, &ray, visibility, &isect);
 
 #ifdef __LAMP_MIS__
-		if(kernel_data.integrator.pdf_lights > 0.0f && !(state.flag & PATH_RAY_CAMERA)) {
+		if(kernel_data.integrator.use_lamp_mis && !(state.flag & PATH_RAY_CAMERA)) {
 			/* ray starting from previous non-transparent bounce */
 			Ray light_ray;
 
@@ -399,20 +399,19 @@ __device float4 kernel_path_progressive(KernelGlobals *kg, RNG *rng, int sample,
 
 				Ray light_ray;
 				BsdfEval L_light;
-				int lamp;
+				bool is_lamp;
 
 #ifdef __OBJECT_MOTION__
 				light_ray.time = sd.time;
 #endif
 
-				if(direct_emission(kg, &sd, -1, light_t, light_o, light_u, light_v, &light_ray, &L_light, &lamp)) {
+				if(direct_emission(kg, &sd, -1, light_t, light_o, light_u, light_v, &light_ray, &L_light, &is_lamp)) {
 					/* trace shadow ray */
 					float3 shadow;
 
 					if(!shadow_blocked(kg, &state, &light_ray, &shadow)) {
 						/* accumulate */
-						bool is_lamp = (lamp != ~0);
-						path_radiance_accum_light(&L, throughput, &L_light, shadow, state.bounce, is_lamp);
+						path_radiance_accum_light(&L, throughput, &L_light, shadow, 1.0f, state.bounce, is_lamp);
 					}
 				}
 			}
@@ -501,7 +500,7 @@ __device void kernel_path_indirect(KernelGlobals *kg, RNG *rng, int sample, Ray
 		bool hit = scene_intersect(kg, &ray, visibility, &isect);
 
 #ifdef __LAMP_MIS__
-		if(kernel_data.integrator.pdf_lights > 0.0f && !(state.flag & PATH_RAY_CAMERA)) {
+		if(kernel_data.integrator.use_lamp_mis && !(state.flag & PATH_RAY_CAMERA)) {
 			/* ray starting from previous non-transparent bounce */
 			Ray light_ray;
 
@@ -612,21 +611,20 @@ __device void kernel_path_indirect(KernelGlobals *kg, RNG *rng, int sample, Ray
 
 				Ray light_ray;
 				BsdfEval L_light;
-				int lamp;
+				bool is_lamp;
 
 #ifdef __OBJECT_MOTION__
 				light_ray.time = sd.time;
 #endif
 
 				/* sample random light */
-				if(direct_emission(kg, &sd, -1, light_t, light_o, light_u, light_v, &light_ray, &L_light, &lamp)) {
+				if(direct_emission(kg, &sd, -1, light_t, light_o, light_u, light_v, &light_ray, &L_light, &is_lamp)) {
 					/* trace shadow ray */
 					float3 shadow;
 
 					if(!shadow_blocked(kg, &state, &light_ray, &shadow)) {
 						/* accumulate */
-						bool is_lamp = (lamp != ~0);
-						path_radiance_accum_light(L, throughput, &L_light, shadow, state.bounce, is_lamp);
+						path_radiance_accum_light(L, throughput, &L_light, shadow, 1.0f, state.bounce, is_lamp);
 					}
 				}
 			}
@@ -819,7 +817,7 @@ __device float4 kernel_path_non_progressive(KernelGlobals *kg, RNG *rng, int sam
 		if(sd.flag & SD_BSDF_HAS_EVAL) {
 			Ray light_ray;
 			BsdfEval L_light;
-			int lamp;
+			bool is_lamp;
 
 #ifdef __OBJECT_MOTION__
 			light_ray.time = sd.time;
@@ -837,14 +835,13 @@ __device float4 kernel_path_non_progressive(KernelGlobals *kg, RNG *rng, int sam
 					float light_u = path_rng(kg, rng, sample*num_samples + j, rng_offset + PRNG_LIGHT_U);
 					float light_v = path_rng(kg, rng, sample*num_samples + j, rng_offset + PRNG_LIGHT_V);
 
-					if(direct_emission(kg, &sd, i, 0.0f, 0.0f, light_u, light_v, &light_ray, &L_light, &lamp)) {
+					if(direct_emission(kg, &sd, i, 0.0f, 0.0f, light_u, light_v, &light_ray, &L_light, &is_lamp)) {
 						/* trace shadow ray */
 						float3 shadow;
 
 						if(!shadow_blocked(kg, &state, &light_ray, &shadow)) {
 							/* accumulate */
-							bool is_lamp = (lamp != ~0);
-							path_radiance_accum_light(&L, throughput*num_samples_inv, &L_light, shadow, state.bounce, is_lamp);
+							path_radiance_accum_light(&L, throughput*num_samples_inv, &L_light, shadow, num_samples_inv, state.bounce, is_lamp);
 						}
 					}
 				}
@@ -867,14 +864,13 @@ __device float4 kernel_path_non_progressive(KernelGlobals *kg, RNG *rng, int sam
 					if(kernel_data.integrator.num_all_lights)
 						light_t = 0.5f*light_t;
 
-					if(direct_emission(kg, &sd, -1, light_t, 0.0f, light_u, light_v, &light_ray, &L_light, &lamp)) {
+					if(direct_emission(kg, &sd, -1, light_t, 0.0f, light_u, light_v, &light_ray, &L_light, &is_lamp)) {
 						/* trace shadow ray */
 						float3 shadow;
 
 						if(!shadow_blocked(kg, &state, &light_ray, &shadow)) {
 							/* accumulate */
-							bool is_lamp = (lamp != ~0);
-							path_radiance_accum_light(&L, throughput*num_samples_inv, &L_light, shadow, state.bounce, is_lamp);
+							path_radiance_accum_light(&L, throughput*num_samples_inv, &L_light, shadow, num_samples_inv, state.bounce, is_lamp);
 						}
 					}
 				}
diff --git a/intern/cycles/kernel/kernel_sse2.cpp b/intern/cycles/kernel/kernel_sse2.cpp
new file mode 100644
index 00000000000..e9d482ae5cf
--- /dev/null
+++ b/intern/cycles/kernel/kernel_sse2.cpp
@@ -0,0 +1,60 @@
+/*
+ * Copyright 2011, Blender Foundation.
+ *
+ * This program is free software; you can redistribute it and/or
+ * modify it under the terms of the GNU General Public License
+ * as published by the Free Software Foundation; either version 2
+ * of the License, or (at your option) any later version.
+ *
+ * This program is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+ * GNU General Public License for more details.
+ *
+ * You should have received a copy of the GNU General Public License
+ * along with this program; if not, write to the Free Software Foundation,
+ * Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
+ */
+
+/* Optimized CPU kernel entry points. This file is compiled with SSE2
+ * optimization flags and nearly all functions inlined, while kernel.cpp
+ * is compiled without for other CPU's. */
+
+#ifdef WITH_OPTIMIZED_KERNEL
+
+#include "kernel.h"
+#include "kernel_compat_cpu.h"
+#include "kernel_math.h"
+#include "kernel_types.h"
+#include "kernel_globals.h"
+#include "kernel_film.h"
+#include "kernel_path.h"
+#include "kernel_displace.h"
+
+CCL_NAMESPACE_BEGIN
+
+/* Path Tracing */
+
+void kernel_cpu_sse2_path_trace(KernelGlobals *kg, float *buffer, unsigned int *rng_state, int sample, int x, int y, int offset, int stride)
+{
+	kernel_path_trace(kg, buffer, rng_state, sample, x, y, offset, stride);
+}
+
+/* Tonemapping */
+
+void kernel_cpu_sse2_tonemap(KernelGlobals *kg, uchar4 *rgba, float *buffer, int sample, int resolution, int x, int y, int offset, int stride)
+{
+	kernel_film_tonemap(kg, rgba, buffer, sample, resolution, x, y, offset, stride);
+}
+
+/* Shader Evaluate */
+
+void kernel_cpu_sse2_shader(KernelGlobals *kg, uint4 *input, float4 *output, int type, int i)
+{
+	kernel_shader_evaluate(kg, input, output, (ShaderEvalType)type, i);
+}
+
+CCL_NAMESPACE_END
+
+#endif
+
diff --git a/intern/cycles/kernel/kernel_optimized.cpp b/intern/cycles/kernel/kernel_sse3.cpp
index 0b662095133..9a8b389cf68 100644
--- a/intern/cycles/kernel/kernel_optimized.cpp
+++ b/intern/cycles/kernel/kernel_sse3.cpp
@@ -35,21 +35,21 @@ CCL_NAMESPACE_BEGIN
 
 /* Path Tracing */
 
-void kernel_cpu_optimized_path_trace(KernelGlobals *kg, float *buffer, unsigned int *rng_state, int sample, int x, int y, int offset, int stride)
+void kernel_cpu_sse3_path_trace(KernelGlobals *kg, float *buffer, unsigned int *rng_state, int sample, int x, int y, int offset, int stride)
 {
 	kernel_path_trace(kg, buffer, rng_state, sample, x, y, offset, stride);
 }
 
 /* Tonemapping */
 
-void kernel_cpu_optimized_tonemap(KernelGlobals *kg, uchar4 *rgba, float *buffer, int sample, int resolution, int x, int y, int offset, int stride)
+void kernel_cpu_sse3_tonemap(KernelGlobals *kg, uchar4 *rgba, float *buffer, int sample, int resolution, int x, int y, int offset, int stride)
 {
 	kernel_film_tonemap(kg, rgba, buffer, sample, resolution, x, y, offset, stride);
 }
 
 /* Shader Evaluate */
 
-void kernel_cpu_optimized_shader(KernelGlobals *kg, uint4 *input, float4 *output, int type, int i)
+void kernel_cpu_sse3_shader(KernelGlobals *kg, uint4 *input, float4 *output, int type, int i)
 {
 	kernel_shader_evaluate(kg, input, output, (ShaderEvalType)type, i);
 }
diff --git a/intern/cycles/kernel/kernel_types.h b/intern/cycles/kernel/kernel_types.h
index 1236f43e018..52bdddc72fb 100644
--- a/intern/cycles/kernel/kernel_types.h
+++ b/intern/cycles/kernel/kernel_types.h
@@ -44,7 +44,6 @@ CCL_NAMESPACE_BEGIN
 #define __KERNEL_SHADING__
 #define __KERNEL_ADV_SHADING__
 #define __NON_PROGRESSIVE__
-#define __LAMP_MIS__
 #define __HAIR__
 #ifdef WITH_OSL
 #define __OSL__
@@ -56,9 +55,6 @@ CCL_NAMESPACE_BEGIN
 #if __CUDA_ARCH__ >= 200
 #define __KERNEL_ADV_SHADING__
 #endif
-#if __CUDA_ARCH__ >= 210
-#define __LAMP_MIS__
-#endif
 #endif
 
 #ifdef __KERNEL_OPENCL__
@@ -116,6 +112,7 @@ CCL_NAMESPACE_BEGIN
 #define __TRANSPARENT_SHADOWS__
 #define __PASSES__
 #define __BACKGROUND_MIS__
+#define __LAMP_MIS__
 #define __AO__
 #define __ANISOTROPIC__
 #define __CAMERA_MOTION__
@@ -288,8 +285,9 @@ typedef enum ShaderFlag {
 	SHADER_SMOOTH_NORMAL = (1 << 31),
 	SHADER_CAST_SHADOW = (1 << 30),
 	SHADER_AREA_LIGHT = (1 << 29),
+	SHADER_USE_MIS = (1 << 28),
 
-	SHADER_MASK = ~(SHADER_SMOOTH_NORMAL|SHADER_CAST_SHADOW|SHADER_AREA_LIGHT)
+	SHADER_MASK = ~(SHADER_SMOOTH_NORMAL|SHADER_CAST_SHADOW|SHADER_AREA_LIGHT|SHADER_USE_MIS)
 } ShaderFlag;
 
 /* Light Type */
@@ -680,7 +678,7 @@ typedef struct KernelIntegrator {
 	int transmission_samples;
 	int ao_samples;
 	int mesh_light_samples;
-	int pad1;
+	int use_lamp_mis;
 } KernelIntegrator;
 
 typedef struct KernelBVH {
diff --git a/intern/cycles/kernel/osl/SConscript b/intern/cycles/kernel/osl/SConscript
index fe7fec463a7..09899567128 100644
--- a/intern/cycles/kernel/osl/SConscript
+++ b/intern/cycles/kernel/osl/SConscript
@@ -46,6 +46,7 @@ defs.append('WITH_OSL')
 if env['OURPLATFORM'] in ('win32-vc', 'win64-vc'):
     cxxflags.append('-DBOOST_NO_RTTI -DBOOST_NO_TYPEID'.split())
     incs.append(env['BF_PTHREADS_INC'])
+    defs.append('OSL_STATIC_LIBRARY')
 else:
     cxxflags.append('-fno-rtti -DBOOST_NO_RTTI -DBOOST_NO_TYPEID'.split())
 
diff --git a/intern/cycles/kernel/shaders/node_musgrave_texture.osl b/intern/cycles/kernel/shaders/node_musgrave_texture.osl
index 38232ea0aeb..8675c23d0d8 100644
--- a/intern/cycles/kernel/shaders/node_musgrave_texture.osl
+++ b/intern/cycles/kernel/shaders/node_musgrave_texture.osl
@@ -37,14 +37,14 @@ float noise_musgrave_fBm(point p, string basis, float H, float lacunarity, float
 	int i;
 
 	for (i = 0; i < (int)octaves; i++) {
-		value += safe_noise(p) * pwr;
+		value += safe_noise(p, 0) * pwr;
 		pwr *= pwHL;
 		p *= lacunarity;
 	}
 
 	rmd = octaves - floor(octaves);
 	if (rmd != 0.0)
-		value += rmd * safe_noise(p) * pwr;
+		value += rmd * safe_noise(p, 0) * pwr;
 
 	return value;
 }
@@ -65,14 +65,14 @@ float noise_musgrave_multi_fractal(point p, string basis, float H, float lacunar
 	int i;
 
 	for (i = 0; i < (int)octaves; i++) {
-		value *= (pwr * safe_noise(p) + 1.0);
+		value *= (pwr * safe_noise(p, 0) + 1.0);
 		pwr *= pwHL;
 		p *= lacunarity;
 	}
 
 	rmd = octaves - floor(octaves);
 	if (rmd != 0.0)
-		value *= (rmd * pwr * safe_noise(p) + 1.0); /* correct? */
+		value *= (rmd * pwr * safe_noise(p, 0) + 1.0); /* correct? */
 
 	return value;
 }
@@ -93,11 +93,11 @@ float noise_musgrave_hetero_terrain(point p, string basis, float H, float lacuna
 	int i;
 
 	/* first unscaled octave of function; later octaves are scaled */
-	value = offset + safe_noise(p);
+	value = offset + safe_noise(p, 0);
 	p *= lacunarity;
 
 	for (i = 1; i < (int)octaves; i++) {
-		increment = (safe_noise(p) + offset) * pwr * value;
+		increment = (safe_noise(p, 0) + offset) * pwr * value;
 		value += increment;
 		pwr *= pwHL;
 		p *= lacunarity;
@@ -105,7 +105,7 @@ float noise_musgrave_hetero_terrain(point p, string basis, float H, float lacuna
 
 	rmd = octaves - floor(octaves);
 	if (rmd != 0.0) {
-		increment = (safe_noise(p) + offset) * pwr * value;
+		increment = (safe_noise(p, 0) + offset) * pwr * value;
 		value += rmd * increment;
 	}
 
@@ -128,7 +128,7 @@ float noise_musgrave_hybrid_multi_fractal(point p, string basis, float H,
 	float pwr = pwHL;
 	int i;
 
-	result = safe_noise(p) + offset;
+	result = safe_noise(p, 0) + offset;
 	weight = gain * result;
 	p *= lacunarity;
 
@@ -136,7 +136,7 @@ float noise_musgrave_hybrid_multi_fractal(point p, string basis, float H,
 		if (weight > 1.0)
 			weight = 1.0;
 
-		signal = (safe_noise(p) + offset) * pwr;
+		signal = (safe_noise(p, 0) + offset) * pwr;
 		pwr *= pwHL;
 		result += weight * signal;
 		weight *= gain * signal;
@@ -145,7 +145,7 @@ float noise_musgrave_hybrid_multi_fractal(point p, string basis, float H,
 
 	rmd = octaves - floor(octaves);
 	if (rmd != 0.0)
-		result += rmd * ((safe_noise(p) + offset) * pwr);
+		result += rmd * ((safe_noise(p, 0) + offset) * pwr);
 
 	return result;
 }
@@ -166,7 +166,7 @@ float noise_musgrave_ridged_multi_fractal(point p, string basis, float H,
 	float pwr = pwHL;
 	int i;
 
-	signal = offset - fabs(safe_noise(p));
+	signal = offset - fabs(safe_noise(p, 0));
 	signal *= signal;
 	result = signal;
 	weight = 1.0;
@@ -174,7 +174,7 @@ float noise_musgrave_ridged_multi_fractal(point p, string basis, float H,
 	for (i = 1; i < (int)octaves; i++) {
 		p *= lacunarity;
 		weight = clamp(signal * gain, 0.0, 1.0);
-		signal = offset - fabs(safe_noise(p));
+		signal = offset - fabs(safe_noise(p, 0));
 		signal *= signal;
 		signal *= weight;
 		result += signal * pwr;
diff --git a/intern/cycles/kernel/shaders/node_texture.h b/intern/cycles/kernel/shaders/node_texture.h
index 2de0fc0ea57..463c68a6c27 100644
--- a/intern/cycles/kernel/shaders/node_texture.h
+++ b/intern/cycles/kernel/shaders/node_texture.h
@@ -151,9 +151,17 @@ float voronoi_CrS(point p) { return 2.0 * voronoi_Cr(p) - 1.0; }
 
 /* Noise Bases */
 
-float safe_noise(point p)
+float safe_noise(point p, int type)
 {
-	float f = noise(p);
+	float f = 0.0;
+	
+	/* Perlin noise in range -1..1 */
+	if (type == 0)
+		f = noise("perlin", p);
+	
+	/* Perlin noise in range 0..1 */
+	else
+		f = noise(p);
 
 	/* can happen for big coordinates, things even out to 0.5 then anyway */
 	if(!isfinite(f))
@@ -167,7 +175,7 @@ float noise_basis(point p, string basis)
 	float result = 0.0;
 
 	if (basis == "Perlin")
-		result = safe_noise(p); /* returns perlin noise in range 0..1 */
+		result = safe_noise(p, 1);
 	if (basis == "Voronoi F1")
 		result = voronoi_F1S(p);
 	if (basis == "Voronoi F2")
diff --git a/intern/cycles/render/curves.h b/intern/cycles/render/curves.h
index 3a12d33d0a7..3527998339c 100644
--- a/intern/cycles/render/curves.h
+++ b/intern/cycles/render/curves.h
@@ -33,9 +33,11 @@ void curvebounds(float *lower, float *upper, float3 *p, int dim);
 
 typedef enum curve_presets {
 	CURVE_CUSTOM,
+	CURVE_FAST_PLANES,
 	CURVE_TANGENT_SHADING,
 	CURVE_TRUE_NORMAL,
-	CURVE_ACCURATE_PRESET
+	CURVE_ACCURATE_PRESET,
+	CURVE_SMOOTH_CURVES
 } curve_presets;
 
 typedef enum curve_primitives {
diff --git a/intern/cycles/render/image.cpp b/intern/cycles/render/image.cpp
index e6f8ab4a5d9..6bfaf48c0c9 100644
--- a/intern/cycles/render/image.cpp
+++ b/intern/cycles/render/image.cpp
@@ -85,14 +85,14 @@ bool ImageManager::set_animation_frame_update(int frame)
 	return false;
 }
 
-bool ImageManager::is_float_image(const string& filename, bool is_builtin)
+bool ImageManager::is_float_image(const string& filename, void *builtin_data)
 {
 	bool is_float = false;
 
-	if(is_builtin) {
+	if(builtin_data) {
 		if(builtin_image_info_cb) {
 			int width, height, channels;
-			builtin_image_info_cb(filename, is_float, width, height, channels);
+			builtin_image_info_cb(filename, builtin_data, is_float, width, height, channels);
 		}
 
 		return is_float;
@@ -123,13 +123,13 @@ bool ImageManager::is_float_image(const string& filename, bool is_builtin)
 	return is_float;
 }
 
-int ImageManager::add_image(const string& filename, bool is_builtin, bool animated, bool& is_float)
+int ImageManager::add_image(const string& filename, void *builtin_data, bool animated, bool& is_float)
 {
 	Image *img;
 	size_t slot;
 
 	/* load image info and find out if we need a float texture */
-	is_float = (pack_images)? false: is_float_image(filename, is_builtin);
+	is_float = (pack_images)? false: is_float_image(filename, builtin_data);
 
 	if(is_float) {
 		/* find existing image */
@@ -160,7 +160,7 @@ int ImageManager::add_image(const string& filename, bool is_builtin, bool animat
 		/* add new image */
 		img = new Image();
 		img->filename = filename;
-		img->is_builtin = is_builtin;
+		img->builtin_data = builtin_data;
 		img->need_load = true;
 		img->animated = animated;
 		img->users = 1;
@@ -195,7 +195,7 @@ int ImageManager::add_image(const string& filename, bool is_builtin, bool animat
 		/* add new image */
 		img = new Image();
 		img->filename = filename;
-		img->is_builtin = is_builtin;
+		img->builtin_data = builtin_data;
 		img->need_load = true;
 		img->animated = animated;
 		img->users = 1;
@@ -209,12 +209,12 @@ int ImageManager::add_image(const string& filename, bool is_builtin, bool animat
 	return slot;
 }
 
-void ImageManager::remove_image(const string& filename, bool is_builtin)
+void ImageManager::remove_image(const string& filename, void *builtin_data)
 {
 	size_t slot;
 
 	for(slot = 0; slot < images.size(); slot++) {
-		if(images[slot] && images[slot]->filename == filename && images[slot]->is_builtin == is_builtin) {
+		if(images[slot] && images[slot]->filename == filename && images[slot]->builtin_data == builtin_data) {
 			/* decrement user count */
 			images[slot]->users--;
 			assert(images[slot]->users >= 0);
@@ -232,7 +232,7 @@ void ImageManager::remove_image(const string& filename, bool is_builtin)
 	if(slot == images.size()) {
 		/* see if it's in a float texture slot */
 		for(slot = 0; slot < float_images.size(); slot++) {
-			if(float_images[slot] && float_images[slot]->filename == filename && float_images[slot]->is_builtin == is_builtin) {
+			if(float_images[slot] && float_images[slot]->filename == filename && float_images[slot]->builtin_data == builtin_data) {
 				/* decrement user count */
 				float_images[slot]->users--;
 				assert(float_images[slot]->users >= 0);
@@ -257,7 +257,7 @@ bool ImageManager::file_load_image(Image *img, device_vector<uchar4>& tex_img)
 	ImageInput *in = NULL;
 	int width, height, components;
 
-	if(!img->is_builtin) {
+	if(!img->builtin_data) {
 		/* load image from file through OIIO */
 		in = ImageInput::create(img->filename);
 
@@ -281,7 +281,7 @@ bool ImageManager::file_load_image(Image *img, device_vector<uchar4>& tex_img)
 			return false;
 
 		bool is_float;
-		builtin_image_info_cb(img->filename, is_float, width, height, components);
+		builtin_image_info_cb(img->filename, img->builtin_data, is_float, width, height, components);
 	}
 
 	/* we only handle certain number of components */
@@ -309,7 +309,7 @@ bool ImageManager::file_load_image(Image *img, device_vector<uchar4>& tex_img)
 		delete in;
 	}
 	else {
-		builtin_image_pixels_cb(img->filename, pixels);
+		builtin_image_pixels_cb(img->filename, img->builtin_data, pixels);
 	}
 
 	if(components == 3) {
@@ -340,7 +340,7 @@ bool ImageManager::file_load_float_image(Image *img, device_vector<float4>& tex_
 	ImageInput *in = NULL;
 	int width, height, components;
 
-	if(!img->is_builtin) {
+	if(!img->builtin_data) {
 		/* load image from file through OIIO */
 		in = ImageInput::create(img->filename);
 
@@ -365,7 +365,7 @@ bool ImageManager::file_load_float_image(Image *img, device_vector<float4>& tex_
 			return false;
 
 		bool is_float;
-		builtin_image_info_cb(img->filename, is_float, width, height, components);
+		builtin_image_info_cb(img->filename, img->builtin_data, is_float, width, height, components);
 	}
 
 	if(!(components == 1 || components == 3 || components == 4)) {
@@ -391,7 +391,7 @@ bool ImageManager::file_load_float_image(Image *img, device_vector<float4>& tex_
 		delete in;
 	}
 	else {
-		builtin_image_float_pixels_cb(img->filename, pixels);
+		builtin_image_float_pixels_cb(img->filename, img->builtin_data, pixels);
 	}
 
 	if(components == 3) {
diff --git a/intern/cycles/render/image.h b/intern/cycles/render/image.h
index e39ac14b60f..464b87ff530 100644
--- a/intern/cycles/render/image.h
+++ b/intern/cycles/render/image.h
@@ -51,9 +51,9 @@ public:
 	ImageManager();
 	~ImageManager();
 
-	int add_image(const string& filename, bool is_builtin, bool animated, bool& is_float);
-	void remove_image(const string& filename, bool is_builtin);
-	bool is_float_image(const string& filename, bool is_builtin);
+	int add_image(const string& filename, void *builtin_data, bool animated, bool& is_float);
+	void remove_image(const string& filename, void *builtin_data);
+	bool is_float_image(const string& filename, void *builtin_data);
 
 	void device_update(Device *device, DeviceScene *dscene, Progress& progress);
 	void device_free(Device *device, DeviceScene *dscene);
@@ -65,9 +65,9 @@ public:
 
 	bool need_update;
 
-	boost::function<void(const string &filename, bool &is_float, int &width, int &height, int &channels)> builtin_image_info_cb;
-	boost::function<bool(const string &filename, unsigned char *pixels)> builtin_image_pixels_cb;
-	boost::function<bool(const string &filename, float *pixels)> builtin_image_float_pixels_cb;
+	boost::function<void(const string &filename, void *data, bool &is_float, int &width, int &height, int &channels)> builtin_image_info_cb;
+	boost::function<bool(const string &filename, void *data, unsigned char *pixels)> builtin_image_pixels_cb;
+	boost::function<bool(const string &filename, void *data, float *pixels)> builtin_image_float_pixels_cb;
 private:
 	int tex_num_images;
 	int tex_num_float_images;
@@ -77,7 +77,7 @@ private:
 
 	struct Image {
 		string filename;
-		bool is_builtin;
+		void *builtin_data;
 
 		bool need_load;
 		bool animated;
diff --git a/intern/cycles/render/light.cpp b/intern/cycles/render/light.cpp
index 04fea1953e7..e7fb9514371 100644
--- a/intern/cycles/render/light.cpp
+++ b/intern/cycles/render/light.cpp
@@ -115,6 +115,8 @@ Light::Light()
 	spot_smooth = 0.0f;
 
 	cast_shadow = true;
+	use_mis = false;
+
 	shader = 0;
 	samples = 1;
 }
@@ -291,13 +293,19 @@ void LightManager::device_update_distribution(Device *device, DeviceScene *dscen
 
 	/* point lights */
 	float lightarea = (totarea > 0.0f)? totarea/scene->lights.size(): 1.0f;
+	bool use_lamp_mis = false;
 
 	for(int i = 0; i < scene->lights.size(); i++, offset++) {
+		Light *light = scene->lights[i];
+
 		distribution[offset].x = totarea;
 		distribution[offset].y = __int_as_float(~(int)i);
 		distribution[offset].z = 1.0f;
-		distribution[offset].w = scene->lights[i]->size;
+		distribution[offset].w = light->size;
 		totarea += lightarea;
+
+		if(light->size > 0.0f && light->use_mis)
+			use_lamp_mis = true;
 	}
 
 	/* normalize cumulative distribution functions */
@@ -344,6 +352,8 @@ void LightManager::device_update_distribution(Device *device, DeviceScene *dscen
 			kintegrator->inv_pdf_lights = 1.0f/kintegrator->pdf_lights;
 		}
 
+		kintegrator->use_lamp_mis = use_lamp_mis;
+
 		/* CDF */
 		device->tex_alloc("__light_distribution", dscene->light_distribution);
 	}
@@ -355,6 +365,7 @@ void LightManager::device_update_distribution(Device *device, DeviceScene *dscen
 		kintegrator->pdf_triangles = 0.0f;
 		kintegrator->pdf_lights = 0.0f;
 		kintegrator->inv_pdf_lights = 0.0f;
+		kintegrator->use_lamp_mis = false;
 	}
 }
 
@@ -484,6 +495,9 @@ void LightManager::device_update_points(Device *device, DeviceScene *dscene, Sce
 			float radius = light->size;
 			float invarea = (radius > 0.0f)? 1.0f/(M_PI_F*radius*radius): 1.0f;
 
+			if(light->use_mis && radius > 0.0f)
+				shader_id |= SHADER_USE_MIS;
+
 			light_data[i*LIGHT_SIZE + 0] = make_float4(__int_as_float(light->type), co.x, co.y, co.z);
 			light_data[i*LIGHT_SIZE + 1] = make_float4(__int_as_float(shader_id), radius, invarea, 0.0f);
 			light_data[i*LIGHT_SIZE + 2] = make_float4(0.0f, 0.0f, 0.0f, 0.0f);
@@ -498,6 +512,9 @@ void LightManager::device_update_points(Device *device, DeviceScene *dscene, Sce
 			float area = M_PI_F*radius*radius;
 			float invarea = (area > 0.0f)? 1.0f/area: 1.0f;
 
+			if(light->use_mis && area > 0.0f)
+				shader_id |= SHADER_USE_MIS;
+
 			light_data[i*LIGHT_SIZE + 0] = make_float4(__int_as_float(light->type), dir.x, dir.y, dir.z);
 			light_data[i*LIGHT_SIZE + 1] = make_float4(__int_as_float(shader_id), radius, cosangle, invarea);
 			light_data[i*LIGHT_SIZE + 2] = make_float4(0.0f, 0.0f, 0.0f, 0.0f);
@@ -505,6 +522,7 @@ void LightManager::device_update_points(Device *device, DeviceScene *dscene, Sce
 		}
 		else if(light->type == LIGHT_BACKGROUND) {
 			shader_id &= ~SHADER_AREA_LIGHT;
+			shader_id |= SHADER_USE_MIS;
 
 			light_data[i*LIGHT_SIZE + 0] = make_float4(__int_as_float(light->type), 0.0f, 0.0f, 0.0f);
 			light_data[i*LIGHT_SIZE + 1] = make_float4(__int_as_float(shader_id), 0.0f, 0.0f, 0.0f);
@@ -515,7 +533,10 @@ void LightManager::device_update_points(Device *device, DeviceScene *dscene, Sce
 			float3 axisu = light->axisu*(light->sizeu*light->size);
 			float3 axisv = light->axisv*(light->sizev*light->size);
 			float area = len(axisu)*len(axisv);
-			float invarea = (area > 0.0f)? 1.0f/area: 0.0f;
+			float invarea = (area > 0.0f)? 1.0f/area: 1.0f;
+
+			if(light->use_mis && area > 0.0f)
+				shader_id |= SHADER_USE_MIS;
 
 			light_data[i*LIGHT_SIZE + 0] = make_float4(__int_as_float(light->type), co.x, co.y, co.z);
 			light_data[i*LIGHT_SIZE + 1] = make_float4(__int_as_float(shader_id), axisu.x, axisu.y, axisu.z);
@@ -530,6 +551,9 @@ void LightManager::device_update_points(Device *device, DeviceScene *dscene, Sce
 			float spot_angle = cosf(light->spot_angle*0.5f);
 			float spot_smooth = (1.0f - spot_angle)*light->spot_smooth;
 
+			if(light->use_mis && radius > 0.0f)
+				shader_id |= SHADER_USE_MIS;
+
 			light_data[i*LIGHT_SIZE + 0] = make_float4(__int_as_float(light->type), co.x, co.y, co.z);
 			light_data[i*LIGHT_SIZE + 1] = make_float4(__int_as_float(shader_id), radius, invarea, spot_angle);
 			light_data[i*LIGHT_SIZE + 2] = make_float4(spot_smooth, dir.x, dir.y, dir.z);
diff --git a/intern/cycles/render/light.h b/intern/cycles/render/light.h
index 3cedde2596e..acd1692a41f 100644
--- a/intern/cycles/render/light.h
+++ b/intern/cycles/render/light.h
@@ -52,6 +52,7 @@ public:
 	float spot_smooth;
 
 	bool cast_shadow;
+	bool use_mis;
 
 	int shader;
 	int samples;
diff --git a/intern/cycles/render/nodes.cpp b/intern/cycles/render/nodes.cpp
index 5444299b948..a0a933ef682 100644
--- a/intern/cycles/render/nodes.cpp
+++ b/intern/cycles/render/nodes.cpp
@@ -142,7 +142,7 @@ ImageTextureNode::ImageTextureNode()
 	slot = -1;
 	is_float = -1;
 	filename = "";
-	is_builtin = false;
+	builtin_data = NULL;
 	color_space = ustring("Color");
 	projection = ustring("Flat");
 	projection_blend = 0.0f;
@@ -156,7 +156,7 @@ ImageTextureNode::ImageTextureNode()
 ImageTextureNode::~ImageTextureNode()
 {
 	if(image_manager)
-		image_manager->remove_image(filename, is_builtin);
+		image_manager->remove_image(filename, builtin_data);
 }
 
 ShaderNode *ImageTextureNode::clone() const
@@ -177,7 +177,7 @@ void ImageTextureNode::compile(SVMCompiler& compiler)
 	image_manager = compiler.image_manager;
 	if(is_float == -1) {
 		bool is_float_bool;
-		slot = image_manager->add_image(filename, is_builtin, animated, is_float_bool);
+		slot = image_manager->add_image(filename, builtin_data, animated, is_float_bool);
 		is_float = (int)is_float_bool;
 	}
 
@@ -238,7 +238,7 @@ void ImageTextureNode::compile(OSLCompiler& compiler)
 	tex_mapping.compile(compiler);
 
 	if(is_float == -1)
-		is_float = (int)image_manager->is_float_image(filename, false);
+		is_float = (int)image_manager->is_float_image(filename, NULL);
 
 	compiler.parameter("filename", filename.c_str());
 	if(is_float || color_space != "Color")
@@ -272,7 +272,7 @@ EnvironmentTextureNode::EnvironmentTextureNode()
 	slot = -1;
 	is_float = -1;
 	filename = "";
-	is_builtin = false;
+	builtin_data = NULL;
 	color_space = ustring("Color");
 	projection = ustring("Equirectangular");
 	animated = false;
@@ -285,7 +285,7 @@ EnvironmentTextureNode::EnvironmentTextureNode()
 EnvironmentTextureNode::~EnvironmentTextureNode()
 {
 	if(image_manager)
-		image_manager->remove_image(filename, is_builtin);
+		image_manager->remove_image(filename, builtin_data);
 }
 
 ShaderNode *EnvironmentTextureNode::clone() const
@@ -306,7 +306,7 @@ void EnvironmentTextureNode::compile(SVMCompiler& compiler)
 	image_manager = compiler.image_manager;
 	if(slot == -1) {
 		bool is_float_bool;
-		slot = image_manager->add_image(filename, is_builtin, animated, is_float_bool);
+		slot = image_manager->add_image(filename, builtin_data, animated, is_float_bool);
 		is_float = (int)is_float_bool;
 	}
 
@@ -356,7 +356,7 @@ void EnvironmentTextureNode::compile(OSLCompiler& compiler)
 	tex_mapping.compile(compiler);
 
 	if(is_float == -1)
-		is_float = (int)image_manager->is_float_image(filename, false);
+		is_float = (int)image_manager->is_float_image(filename, NULL);
 
 	compiler.parameter("filename", filename.c_str());
 	compiler.parameter("projection", projection);
diff --git a/intern/cycles/render/nodes.h b/intern/cycles/render/nodes.h
index 8b2d6a0e5c8..66be919d098 100644
--- a/intern/cycles/render/nodes.h
+++ b/intern/cycles/render/nodes.h
@@ -70,7 +70,7 @@ public:
 	int slot;
 	int is_float;
 	string filename;
-	bool is_builtin;
+	void *builtin_data;
 	ustring color_space;
 	ustring projection;
 	float projection_blend;
@@ -90,7 +90,7 @@ public:
 	int slot;
 	int is_float;
 	string filename;
-	bool is_builtin;
+	void *builtin_data;
 	ustring color_space;
 	ustring projection;
 	bool animated;
diff --git a/intern/cycles/render/object.cpp b/intern/cycles/render/object.cpp
index 8d8087266c1..b31650c160a 100644
--- a/intern/cycles/render/object.cpp
+++ b/intern/cycles/render/object.cpp
@@ -84,12 +84,20 @@ void Object::apply_transform()
 {
 	if(!mesh || tfm == transform_identity())
 		return;
-	
+
+	float3 c0 = transform_get_column(&tfm, 0);
+	float3 c1 = transform_get_column(&tfm, 1);
+	float3 c2 = transform_get_column(&tfm, 2);
+	float scalar = pow(fabsf(dot(cross(c0, c1), c2)), 1.0f/3.0f);
+
 	for(size_t i = 0; i < mesh->verts.size(); i++)
 		mesh->verts[i] = transform_point(&tfm, mesh->verts[i]);
 
-	for(size_t i = 0; i < mesh->curve_keys.size(); i++)
+	for(size_t i = 0; i < mesh->curve_keys.size(); i++) {
 		mesh->curve_keys[i].co = transform_point(&tfm, mesh->curve_keys[i].co);
+		/* scale for strand radius - only correct for uniform transforms*/
+		mesh->curve_keys[i].radius *= scalar;
+	}
 
 	/* store matrix to transform later. when accessing these as attributes we
 	 * do not want the transform to be applied for consistency between static
diff --git a/intern/cycles/render/session.cpp b/intern/cycles/render/session.cpp
index e49a4a3ed34..94cb69102d2 100644
--- a/intern/cycles/render/session.cpp
+++ b/intern/cycles/render/session.cpp
@@ -381,7 +381,7 @@ bool Session::acquire_tile(Device *tile_device, RenderTile& rtile)
 
 	buffer_params.get_offset_stride(rtile.offset, rtile.stride);
 
-	RenderBuffers *tilebuffers = new RenderBuffers(tile_device);
+	RenderBuffers *tilebuffers;
 
 	/* allocate buffers */
 	if(params.progressive_refine) {
@@ -552,7 +552,7 @@ void Session::run_cpu()
 			}
 			else if(need_tonemap) {
 				/* tonemap only if we do not reset, we don't we don't
-				 * wan't to show the result of an incomplete sample*/
+				 * want to show the result of an incomplete sample*/
 				tonemap();
 			}
 
diff --git a/intern/cycles/util/util_system.cpp b/intern/cycles/util/util_system.cpp
index 2d9f0fffae6..4fda090e09e 100644
--- a/intern/cycles/util/util_system.cpp
+++ b/intern/cycles/util/util_system.cpp
@@ -136,7 +136,7 @@ struct CPUCapabilities {
 	bool fma4;
 };
 
-bool system_cpu_support_optimized()
+static CPUCapabilities& system_cpu_capabilities()
 {
 	static CPUCapabilities caps;
 	static bool caps_init = false;
@@ -182,7 +182,18 @@ bool system_cpu_support_optimized()
 		caps_init = true;
 	}
 
-	/* optimization flags use these */
+	return caps;
+}
+
+bool system_cpu_support_sse2()
+{
+	CPUCapabilities& caps = system_cpu_capabilities();
+	return caps.sse && caps.sse2;
+}
+
+bool system_cpu_support_sse3()
+{
+	CPUCapabilities& caps = system_cpu_capabilities();
 	return caps.sse && caps.sse2 && caps.sse3;
 }
 
diff --git a/intern/cycles/util/util_system.h b/intern/cycles/util/util_system.h
index f25e009a250..257112883d1 100644
--- a/intern/cycles/util/util_system.h
+++ b/intern/cycles/util/util_system.h
@@ -26,7 +26,8 @@ CCL_NAMESPACE_BEGIN
 int system_cpu_thread_count();
 string system_cpu_brand_string();
 int system_cpu_bits();
-bool system_cpu_support_optimized();
+bool system_cpu_support_sse2();
+bool system_cpu_support_sse3();
 
 CCL_NAMESPACE_END
 
diff --git a/intern/ghost/CMakeLists.txt b/intern/ghost/CMakeLists.txt
index ea09987c564..12dd4c0d3eb 100644
--- a/intern/ghost/CMakeLists.txt
+++ b/intern/ghost/CMakeLists.txt
@@ -271,6 +271,9 @@ elseif(UNIX)
 
 	if(WITH_X11_XINPUT)
 		add_definitions(-DWITH_X11_XINPUT)
+		list(APPEND INC_SYS
+			${X11_Xinput_INCLUDE_PATH}
+		)
 	endif()
 
 elseif(WIN32)
diff --git a/intern/ghost/GHOST_C-api.h b/intern/ghost/GHOST_C-api.h
index ffdee6c1550..a92d0d33b65 100644
--- a/intern/ghost/GHOST_C-api.h
+++ b/intern/ghost/GHOST_C-api.h
@@ -149,6 +149,20 @@ extern void GHOST_GetMainDisplayDimensions(GHOST_SystemHandle systemhandle,
                                            GHOST_TUns32 *height);
 
 /**
+ * Returns the dimensions of all displays combine
+ * (the current workspace).
+ * No need to worrky about overlapping monitors.
+ * \param systemhandle The handle to the system
+ * \param width A pointer the width gets put in
+ * \param height A pointer the height gets put in
+ * \return void.
+ */
+extern void GHOST_GetAllDisplayDimensions(GHOST_SystemHandle systemhandle,
+                                           GHOST_TUns32 *width,
+                                           GHOST_TUns32 *height);
+
+
+/**
  * Create a new window.
  * The new window is added to the list of windows managed. 
  * Never explicitly delete the window, use disposeWindow() instead.
diff --git a/intern/ghost/GHOST_ISystem.h b/intern/ghost/GHOST_ISystem.h
index 1b3509c1ac3..841293e09e4 100644
--- a/intern/ghost/GHOST_ISystem.h
+++ b/intern/ghost/GHOST_ISystem.h
@@ -227,6 +227,12 @@ public:
 	virtual void getMainDisplayDimensions(GHOST_TUns32& width, GHOST_TUns32& height) const = 0;
 
 	/**
+	 * Returns the combine dimensions of all monitors.
+	 * \return The dimension of the workspace.
+	 */
+	virtual void getAllDisplayDimensions(GHOST_TUns32& width, GHOST_TUns32& height) const = 0;
+
+	/**
 	 * Create a new window.
 	 * The new window is added to the list of windows managed.
 	 * Never explicitly delete the window, use disposeWindow() instead.
diff --git a/intern/ghost/intern/GHOST_C-api.cpp b/intern/ghost/intern/GHOST_C-api.cpp
index be64acf8c94..7622a2a7330 100644
--- a/intern/ghost/intern/GHOST_C-api.cpp
+++ b/intern/ghost/intern/GHOST_C-api.cpp
@@ -123,7 +123,14 @@ void GHOST_GetMainDisplayDimensions(GHOST_SystemHandle systemhandle,
 	system->getMainDisplayDimensions(*width, *height);
 }
 
+void GHOST_GetAllDisplayDimensions(GHOST_SystemHandle systemhandle,
+                                    GHOST_TUns32 *width,
+                                    GHOST_TUns32 *height)
+{
+	GHOST_ISystem *system = (GHOST_ISystem *) systemhandle;
 
+	system->getAllDisplayDimensions(*width, *height);
+}
 
 GHOST_WindowHandle GHOST_CreateWindow(GHOST_SystemHandle systemhandle,
                                       const char *title,
diff --git a/intern/ghost/intern/GHOST_DisplayManagerCarbon.cpp b/intern/ghost/intern/GHOST_DisplayManagerCarbon.cpp
index 3f7d2102dac..25e9123dae6 100644
--- a/intern/ghost/intern/GHOST_DisplayManagerCarbon.cpp
+++ b/intern/ghost/intern/GHOST_DisplayManagerCarbon.cpp
@@ -150,7 +150,7 @@ GHOST_TSuccess GHOST_DisplayManagerCarbon::setCurrentDisplaySetting(GHOST_TUns8
 #endif // GHOST_DEBUG
 
 	CGDisplayErr err = ::CGDisplaySwitchToMode(m_displayIDs[display], displayModeValues);
-        
+
 	return err == CGDisplayNoErr ? GHOST_kSuccess : GHOST_kFailure;
 }
 
@@ -158,19 +158,18 @@ GHOST_TSuccess GHOST_DisplayManagerCarbon::setCurrentDisplaySetting(GHOST_TUns8
 long GHOST_DisplayManagerCarbon::getValue(CFDictionaryRef values, CFStringRef key) const
 {
 	CFNumberRef numberValue = (CFNumberRef) CFDictionaryGetValue(values, key);
-    
+
 	if (!numberValue)
 	{
 		return -1;
 	}
-    
+
 	long intValue;
-    
+
 	if (!CFNumberGetValue(numberValue, kCFNumberLongType, &intValue))
 	{
 		return -1;
 	}
-    
+
 	return intValue;
 }
-
diff --git a/intern/ghost/intern/GHOST_SystemCarbon.cpp b/intern/ghost/intern/GHOST_SystemCarbon.cpp
index f5784c7d451..bb6b1bfd4af 100644
--- a/intern/ghost/intern/GHOST_SystemCarbon.cpp
+++ b/intern/ghost/intern/GHOST_SystemCarbon.cpp
@@ -396,6 +396,11 @@ void GHOST_SystemCarbon::getMainDisplayDimensions(GHOST_TUns32& width, GHOST_TUn
 	height = bnds.bottom - bnds.top;
 }
 
+void GHOST_SystemCarbon::getAllDisplayDimensions(GHOST_TUns32& width, GHOST_TUns32& height) const
+{
+	/* TODO */
+	getMainDisplayDimensions(width, height);
+}
 
 GHOST_IWindow *GHOST_SystemCarbon::createWindow(
 		const STR_String& title,
diff --git a/intern/ghost/intern/GHOST_SystemCarbon.h b/intern/ghost/intern/GHOST_SystemCarbon.h
index 61e5a086a0f..b6821bc7419 100644
--- a/intern/ghost/intern/GHOST_SystemCarbon.h
+++ b/intern/ghost/intern/GHOST_SystemCarbon.h
@@ -91,6 +91,12 @@ public:
 	 * \return The dimension of the main display.
 	 */
 	virtual void getMainDisplayDimensions(GHOST_TUns32& width, GHOST_TUns32& height) const;
+
+	/**
+	 * Returns the combine dimensions of all monitors.
+	 * \return The dimension of the workspace.
+	 */
+	virtual void getAllDisplayDimensions(GHOST_TUns32& width, GHOST_TUns32& height) const;
 	
 	/**
 	 * Create a new window.
diff --git a/intern/ghost/intern/GHOST_SystemCocoa.h b/intern/ghost/intern/GHOST_SystemCocoa.h
index a1b372dac9a..257aeb52a37 100644
--- a/intern/ghost/intern/GHOST_SystemCocoa.h
+++ b/intern/ghost/intern/GHOST_SystemCocoa.h
@@ -89,6 +89,11 @@ public:
 	 */
 	virtual void getMainDisplayDimensions(GHOST_TUns32& width, GHOST_TUns32& height) const;
 	
+	/** Returns the combine dimensions of all monitors.
+	 * \return The dimension of the workspace.
+	 */
+	virtual void getAllDisplayDimensions(GHOST_TUns32& width, GHOST_TUns32& height) const;
+
 	/**
 	 * Create a new window.
 	 * The new window is added to the list of windows managed.
diff --git a/intern/ghost/intern/GHOST_SystemCocoa.mm b/intern/ghost/intern/GHOST_SystemCocoa.mm
index 1665180b687..e4245aaf095 100644
--- a/intern/ghost/intern/GHOST_SystemCocoa.mm
+++ b/intern/ghost/intern/GHOST_SystemCocoa.mm
@@ -719,6 +719,11 @@ void GHOST_SystemCocoa::getMainDisplayDimensions(GHOST_TUns32& width, GHOST_TUns
 	[pool drain];
 }
 
+void GHOST_SystemCocoa::getAllDisplayDimensions(GHOST_TUns32& width, GHOST_TUns32& height) const
+{
+	/* TODO! */
+	getMainDisplayDimensions(width, height);
+}
 
 GHOST_IWindow* GHOST_SystemCocoa::createWindow(
 	const STR_String& title, 
diff --git a/intern/ghost/intern/GHOST_SystemNULL.h b/intern/ghost/intern/GHOST_SystemNULL.h
index 63e3b5bc828..7021dea36dd 100644
--- a/intern/ghost/intern/GHOST_SystemNULL.h
+++ b/intern/ghost/intern/GHOST_SystemNULL.h
@@ -38,7 +38,7 @@ class GHOST_WindowNULL;
 class GHOST_SystemNULL : public GHOST_System {
 public:
 
-	GHOST_SystemNULL( ) : GHOST_System() { /* nop */ }
+	GHOST_SystemNULL() : GHOST_System() { /* nop */ }
 	~GHOST_SystemNULL() { /* nop */ }
 	bool processEvents(bool waitForEvent) { return false; }
 	int toggleConsole(int action) { return 0; }
@@ -46,11 +46,12 @@ public:
 	GHOST_TSuccess getButtons(GHOST_Buttons& buttons) const { return GHOST_kSuccess; }
 	GHOST_TUns8 *getClipboard(bool selection) const { return NULL; }
 	void putClipboard(GHOST_TInt8 *buffer, bool selection) const { /* nop */ }
-	GHOST_TUns64 getMilliSeconds( ) const { return 0; }
-	GHOST_TUns8 getNumDisplays( ) const { return GHOST_TUns8(1); }
-	GHOST_TSuccess getCursorPosition( GHOST_TInt32& x, GHOST_TInt32& y ) const { return GHOST_kFailure; }
-	GHOST_TSuccess setCursorPosition( GHOST_TInt32 x, GHOST_TInt32 y ) { return GHOST_kFailure; }
-	void getMainDisplayDimensions( GHOST_TUns32& width, GHOST_TUns32& height ) const { /* nop */ }
+	GHOST_TUns64 getMilliSeconds() const { return 0; }
+	GHOST_TUns8 getNumDisplays() const { return GHOST_TUns8(1); }
+	GHOST_TSuccess getCursorPosition(GHOST_TInt32& x, GHOST_TInt32& y) const { return GHOST_kFailure; }
+	GHOST_TSuccess setCursorPosition(GHOST_TInt32 x, GHOST_TInt32 y) { return GHOST_kFailure; }
+	void getMainDisplayDimensions(GHOST_TUns32& width, GHOST_TUns32& height) const { /* nop */ }
+	void getAllDisplayDimensions(GHOST_TUns32& width, GHOST_TUns32& height) const { /* nop */ }
 
 	GHOST_TSuccess init() {
 		GHOST_TSuccess success = GHOST_System::init();
@@ -82,11 +83,4 @@ public:
 	}
 };
 
-#endif
-
-
-
-
-
-
-
+#endif  /* __GHOST_SYSTEMNULL_H__ */
diff --git a/intern/ghost/intern/GHOST_SystemSDL.cpp b/intern/ghost/intern/GHOST_SystemSDL.cpp
index fdc4f33b784..da1836d88a0 100644
--- a/intern/ghost/intern/GHOST_SystemSDL.cpp
+++ b/intern/ghost/intern/GHOST_SystemSDL.cpp
@@ -114,6 +114,20 @@ GHOST_SystemSDL::init() {
 	return GHOST_kFailure;
 }
 
+/**
+ * Returns the dimensions of the main display on this system.
+ * \return The dimension of the main display.
+ */
+void
+GHOST_SystemSDL::getAllDisplayDimensions(GHOST_TUns32& width,
+                                         GHOST_TUns32& height) const
+{
+	SDL_DisplayMode mode;
+	SDL_GetDesktopDisplayMode(0, &mode); /* note, always 0 display */
+	width = mode.w;
+	height = mode.h;
+}
+
 void
 GHOST_SystemSDL::getMainDisplayDimensions(GHOST_TUns32& width,
                                           GHOST_TUns32& height) const
diff --git a/intern/ghost/intern/GHOST_SystemSDL.h b/intern/ghost/intern/GHOST_SystemSDL.h
index 43c9129fdd8..776dc2f66a5 100644
--- a/intern/ghost/intern/GHOST_SystemSDL.h
+++ b/intern/ghost/intern/GHOST_SystemSDL.h
@@ -88,6 +88,10 @@ public:
 	                  GHOST_TInt32 y);
 
 	void
+	getAllDisplayDimensions(GHOST_TUns32& width,
+	                        GHOST_TUns32& height) const;
+
+	void
 	getMainDisplayDimensions(GHOST_TUns32& width,
 	                         GHOST_TUns32& height) const;
 
diff --git a/intern/ghost/intern/GHOST_SystemWin32.cpp b/intern/ghost/intern/GHOST_SystemWin32.cpp
index 50d7b372dd6..138109ce48b 100644
--- a/intern/ghost/intern/GHOST_SystemWin32.cpp
+++ b/intern/ghost/intern/GHOST_SystemWin32.cpp
@@ -215,6 +215,11 @@ void GHOST_SystemWin32::getMainDisplayDimensions(GHOST_TUns32& width, GHOST_TUns
 	height = ::GetSystemMetrics(SM_CYSCREEN);
 }
 
+void GHOST_SystemWin32::getAllDisplayDimensions(GHOST_TUns32& width, GHOST_TUns32& height) const
+{
+	width = ::GetSystemMetrics(SM_XVIRTUALSCREEN);
+	height = ::GetSystemMetrics(SM_YVIRTUALSCREEN);
+}
 
 GHOST_IWindow *GHOST_SystemWin32::createWindow(
 		const STR_String& title,
diff --git a/intern/ghost/intern/GHOST_SystemWin32.h b/intern/ghost/intern/GHOST_SystemWin32.h
index 6a756d35872..e47dbda2e92 100644
--- a/intern/ghost/intern/GHOST_SystemWin32.h
+++ b/intern/ghost/intern/GHOST_SystemWin32.h
@@ -100,7 +100,13 @@ public:
 	 * \return The dimension of the main display.
 	 */
 	virtual void getMainDisplayDimensions(GHOST_TUns32& width, GHOST_TUns32& height) const;
-	
+
+	/**
+	 * Returns the dimensions of all displays on this system.
+	 * \return The dimension of the main display.
+	 */
+	virtual void getAllDisplayDimensions(GHOST_TUns32& width, GHOST_TUns32& height) const;
+
 	/**
 	 * Create a new window.
 	 * The new window is added to the list of windows managed. 
diff --git a/intern/ghost/intern/GHOST_SystemX11.cpp b/intern/ghost/intern/GHOST_SystemX11.cpp
index 28a228b2777..7ba8889fd1e 100644
--- a/intern/ghost/intern/GHOST_SystemX11.cpp
+++ b/intern/ghost/intern/GHOST_SystemX11.cpp
@@ -72,6 +72,9 @@
 /* for debugging - so we can breakpoint X11 errors */
 // #define USE_X11_ERROR_HANDLERS
 
+/* see [#34039] Fix Alt key glitch on Unity desktop */
+#define USE_UNITY_WORKAROUND
+
 static GHOST_TKey convertXKey(KeySym key);
 
 /* these are for copy and select copy */
@@ -234,6 +237,25 @@ getMainDisplayDimensions(
 		GHOST_TUns32& height) const
 {
 	if (m_display) {
+		/* note, for this to work as documented,
+		 * we would need to use Xinerama check r54370 for code that did thia,
+		 * we've since removed since its not worth the extra dep - campbell */
+		getAllDisplayDimensions(width, height);
+	}
+}
+
+
+/**
+ * Returns the dimensions of the main display on this system.
+ * \return The dimension of the main display.
+ */
+void
+GHOST_SystemX11::
+getAllDisplayDimensions(
+		GHOST_TUns32& width,
+		GHOST_TUns32& height) const
+{
+	if (m_display) {
 		width  = DisplayWidth(m_display, DefaultScreen(m_display));
 		height = DisplayHeight(m_display, DefaultScreen(m_display));
 	}
@@ -496,6 +518,54 @@ processEvents(
 
 			processEvent(&xevent);
 			anyProcessed = true;
+
+
+#ifdef USE_UNITY_WORKAROUND
+			/* note: processEvent() can't include this code because
+			 * KeymapNotify event have no valid window information. */
+
+			/* the X server generates KeymapNotify event immediately after
+			 * every EnterNotify and FocusIn event.  we handle this event
+			 * to correct modifier states. */
+			if (xevent.type == FocusIn) {
+				/* use previous event's window, because KeymapNotify event
+				 * has no window information. */
+				GHOST_WindowX11 *window = findGhostWindow(xevent.xany.window);
+				if (window && XPending(m_display) >= 2) {
+					XNextEvent(m_display, &xevent);
+
+					if (xevent.type == KeymapNotify) {
+						XEvent xev_next;
+
+						/* check if KeyPress or KeyRelease event was generated
+						 * in order to confirm the window is active. */
+						XPeekEvent(m_display, &xev_next);
+
+						if (xev_next.type == KeyPress || xev_next.type == KeyRelease) {
+							/* XK_Hyper_L/R currently unused */
+							const static KeySym modifiers[8] = {XK_Shift_L, XK_Shift_R,
+							                                    XK_Control_L, XK_Control_R,
+							                                    XK_Alt_L, XK_Alt_R,
+							                                    XK_Super_L, XK_Super_R};
+
+							for (int i = 0; i < (sizeof(modifiers) / sizeof(*modifiers)); i++) {
+								KeyCode kc = XKeysymToKeycode(m_display, modifiers[i]);
+								if (((xevent.xkeymap.key_vector[kc >> 3] >> (kc & 7)) & 1) != 0) {
+									pushEvent(new GHOST_EventKey(
+									              getMilliSeconds(),
+									              GHOST_kEventKeyDown,
+									              window,
+									              convertXKey(modifiers[i]),
+									              '\0',
+									              NULL));
+								}
+							}
+						}
+					}
+				}
+			}
+#endif  /* USE_UNITY_WORKAROUND */
+
 		}
 		
 		if (generateWindowExposeEvents()) {
diff --git a/intern/ghost/intern/GHOST_SystemX11.h b/intern/ghost/intern/GHOST_SystemX11.h
index a5d5d9b7a99..5fd494637f9 100644
--- a/intern/ghost/intern/GHOST_SystemX11.h
+++ b/intern/ghost/intern/GHOST_SystemX11.h
@@ -117,6 +117,16 @@ public:
 	    ) const;
 
 	/**
+	 * Returns the dimensions of all displays on this system.
+	 * \return The dimension of the main display.
+	 */
+	void
+	getAllDisplayDimensions(
+	    GHOST_TUns32& width,
+	    GHOST_TUns32& height
+	    ) const;
+
+	/**
 	 * Create a new window.
 	 * The new window is added to the list of windows managed. 
 	 * Never explicitly delete the window, use disposeWindow() instead.
diff --git a/intern/ghost/intern/GHOST_WindowCocoa.mm b/intern/ghost/intern/GHOST_WindowCocoa.mm
index ccd6cd93d4c..8d5f8bf89d5 100644
--- a/intern/ghost/intern/GHOST_WindowCocoa.mm
+++ b/intern/ghost/intern/GHOST_WindowCocoa.mm
@@ -58,9 +58,7 @@ extern "C" {
 	extern void wm_draw_update(bContext *C);
 };*/
 @interface CocoaWindowDelegate : NSObject
-#if MAC_OS_X_VERSION_MIN_REQUIRED >= 1060
 <NSWindowDelegate>
-#endif
 {
 	GHOST_SystemCocoa *systemCocoa;
 	GHOST_WindowCocoa *associatedWindow;
@@ -998,9 +996,7 @@ GHOST_TSuccess GHOST_WindowCocoa::setState(GHOST_TWindowState state)
 				//Show the new window
 				[tmpWindow makeKeyAndOrderFront:m_openGLView];
 				//Close and release old window
-				[m_window setDelegate:nil]; // To avoid the notification of "window closed" event
 				[m_window close];
-				[m_window release];
 				m_window = tmpWindow;
 #endif
 			
@@ -1057,9 +1053,7 @@ GHOST_TSuccess GHOST_WindowCocoa::setState(GHOST_TWindowState state)
 				//Show the new window
 				[tmpWindow makeKeyAndOrderFront:nil];
 				//Close and release old window
-				[m_window setDelegate:nil]; // To avoid the notification of "window closed" event
 				[m_window close];
-				[m_window release];
 				m_window = tmpWindow;
 #endif
 			
diff --git a/intern/ghost/intern/GHOST_WindowX11.cpp b/intern/ghost/intern/GHOST_WindowX11.cpp
index 2e9ac94b9c7..7ef3a739b7e 100644
--- a/intern/ghost/intern/GHOST_WindowX11.cpp
+++ b/intern/ghost/intern/GHOST_WindowX11.cpp
@@ -266,15 +266,14 @@ GHOST_WindowX11(
 	/* Specify which events we are interested in hearing. */
 
 	xattributes.event_mask =
-	    ExposureMask | StructureNotifyMask |
-	    KeyPressMask | KeyReleaseMask |
-	    EnterWindowMask | LeaveWindowMask |
-	    ButtonPressMask | ButtonReleaseMask |
-	    PointerMotionMask | FocusChangeMask | PropertyChangeMask;
+	        ExposureMask | StructureNotifyMask |
+	        KeyPressMask | KeyReleaseMask |
+	        EnterWindowMask | LeaveWindowMask |
+	        ButtonPressMask | ButtonReleaseMask |
+	        PointerMotionMask | FocusChangeMask |
+	        PropertyChangeMask | KeymapStateMask;
 
 	/* create the window! */
-
-	;
 	if (parentWindow == 0) {
 		m_window =  XCreateWindow(m_display,
 		                          RootWindow(m_display, m_visual->screen),
@@ -329,6 +328,26 @@ GHOST_WindowX11(
 	GHOST_PRINT("Set drop target\n");
 #endif
 
+	if (state == GHOST_kWindowStateMaximized || state == GHOST_kWindowStateFullScreen) {
+		Atom _NET_WM_STATE = XInternAtom(m_display, "_NET_WM_STATE", False);
+		Atom _NET_WM_STATE_MAXIMIZED_VERT = XInternAtom(m_display, "_NET_WM_STATE_MAXIMIZED_VERT", False);
+		Atom _NET_WM_STATE_MAXIMIZED_HORZ = XInternAtom(m_display, "_NET_WM_STATE_MAXIMIZED_HORZ", False);
+		Atom _NET_WM_STATE_FULLSCREEN     = XInternAtom(m_display, "_NET_WM_STATE_FULLSCREEN", False);
+		Atom atoms[2];
+		int count = 0;
+
+		if (state == GHOST_kWindowStateMaximized) {
+			atoms[count++] = _NET_WM_STATE_MAXIMIZED_VERT;
+			atoms[count++] = _NET_WM_STATE_MAXIMIZED_HORZ;
+		}
+		else {
+			atoms[count++] = _NET_WM_STATE_FULLSCREEN;
+		}
+
+		XChangeProperty(m_display, m_window, _NET_WM_STATE, XA_ATOM, 32,
+		                PropModeReplace, (unsigned char *)atoms, count);
+		m_post_init = False;
+	}
 	/*
 	 * One of the problem with WM-spec is that can't set a property
 	 * to a window that isn't mapped. That is why we can't "just
@@ -340,7 +359,7 @@ GHOST_WindowX11(
 	 * So, m_post_init indicate that we need wait for the MapNotify
 	 * event and then set the Window state to the m_post_state.
 	 */
-	if ((state != GHOST_kWindowStateNormal) && (state != GHOST_kWindowStateMinimized)) {
+	else if ((state != GHOST_kWindowStateNormal) && (state != GHOST_kWindowStateMinimized)) {
 		m_post_init = True;
 		m_post_state = state;
 	}
@@ -508,7 +527,7 @@ bool GHOST_WindowX11::createX11_XIC()
 	             EnterWindowMask | LeaveWindowMask |
 	             ButtonPressMask | ButtonReleaseMask |
 	             PointerMotionMask | FocusChangeMask |
-	             PropertyChangeMask | fevent);
+	             PropertyChangeMask | KeymapStateMask | fevent);
 	return true;
 }
 #endif
@@ -916,12 +935,12 @@ GHOST_TWindowState GHOST_WindowX11::getState() const
 	 */
 	if ((state == IconicState) || (state == WithdrawnState))
 		state_ret = GHOST_kWindowStateMinimized;
-	else if (netwmIsMaximized() == True)
-		state_ret = GHOST_kWindowStateMaximized;
 	else if (netwmIsFullScreen() == True)
 		state_ret = GHOST_kWindowStateFullScreen;
 	else if (motifIsFullScreen() == True)
 		state_ret = GHOST_kWindowStateFullScreen;
+	else if (netwmIsMaximized() == True)
+		state_ret = GHOST_kWindowStateMaximized;
 	return (state_ret);
 }
 
diff --git a/intern/opencl/CMakeLists.txt b/intern/opencl/CMakeLists.txt
new file mode 100644
index 00000000000..03855cfdf8b
--- /dev/null
+++ b/intern/opencl/CMakeLists.txt
@@ -0,0 +1,42 @@
+# ***** BEGIN GPL LICENSE BLOCK *****
+#
+# This program is free software; you can redistribute it and/or
+# modify it under the terms of the GNU General Public License
+# as published by the Free Software Foundation; either version 2
+# of the License, or (at your option) any later version.
+#
+# This program is distributed in the hope that it will be useful,
+# but WITHOUT ANY WARRANTY; without even the implied warranty of
+# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+# GNU General Public License for more details.
+#
+# You should have received a copy of the GNU General Public License
+# along with this program; if not, write to the Free Software Foundation,
+# Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
+#
+# The Original Code is Copyright (C) 2006, Blender Foundation
+# All rights reserved.
+#
+# The Original Code is: all of this file.
+#
+# Contributor(s): Jacques Beaurain.
+#
+# ***** END GPL LICENSE BLOCK *****
+
+set(INC
+	.
+)
+
+set(INC_SYS
+
+)
+
+set(SRC
+	OCL_opencl.h
+	intern/clew.h
+	intern/clew.c
+	intern/OCL_opencl.c
+)
+
+
+blender_add_lib(bf_intern_opencl "${SRC}" "${INC}" "${INC_SYS}")
diff --git a/intern/opencl/OCL_opencl.h b/intern/opencl/OCL_opencl.h
new file mode 100644
index 00000000000..4ee167b2fb4
--- /dev/null
+++ b/intern/opencl/OCL_opencl.h
@@ -0,0 +1,37 @@
+/*
+ * Copyright 2011, Blender Foundation.
+ *
+ * This program is free software; you can redistribute it and/or
+ * modify it under the terms of the GNU General Public License
+ * as published by the Free Software Foundation; either version 2
+ * of the License, or (at your option) any later version.
+ *
+ * This program is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+ * GNU General Public License for more details.
+ *
+ * You should have received a copy of the GNU General Public License
+ * along with this program; if not, write to the Free Software Foundation,
+ * Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
+ *
+ * Contributor: 
+ *		Jeroen Bakker 
+ *		Monique Dewanchand
+ */
+
+#ifndef OCL_OPENCL_H
+#define OCL_OPENCL_H
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+#include "intern/clew.h"
+void OCL_init(void);
+
+#ifdef __cplusplus
+}
+#endif
+
+#endif
diff --git a/intern/opencl/SConscript b/intern/opencl/SConscript
new file mode 100644
index 00000000000..41a6d720098
--- /dev/null
+++ b/intern/opencl/SConscript
@@ -0,0 +1,34 @@
+#!/usr/bin/env python
+#
+# ***** BEGIN GPL LICENSE BLOCK *****
+#
+# This program is free software; you can redistribute it and/or
+# modify it under the terms of the GNU General Public License
+# as published by the Free Software Foundation; either version 2
+# of the License, or (at your option) any later version.
+#
+# This program is distributed in the hope that it will be useful,
+# but WITHOUT ANY WARRANTY; without even the implied warranty of
+# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+# GNU General Public License for more details.
+#
+# You should have received a copy of the GNU General Public License
+# along with this program; if not, write to the Free Software Foundation,
+# Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
+#
+# The Original Code is Copyright (C) 2006, Blender Foundation
+# All rights reserved.
+#
+# The Original Code is: all of this file.
+#
+# Contributor(s): Nathan Letwory.
+#
+# ***** END GPL LICENSE BLOCK *****
+
+Import ('env')
+
+sources = env.Glob('intern/*.c')
+
+incs = '.'
+
+env.BlenderLib ( 'bf_intern_opencl', sources, Split(incs), libtype=['core','player'], priority = [192,192] )
diff --git a/intern/opencl/intern/OCL_opencl.c b/intern/opencl/intern/OCL_opencl.c
new file mode 100644
index 00000000000..e3130e16bde
--- /dev/null
+++ b/intern/opencl/intern/OCL_opencl.c
@@ -0,0 +1,37 @@
+/*
+ * Copyright 2011, Blender Foundation.
+ *
+ * This program is free software; you can redistribute it and/or
+ * modify it under the terms of the GNU General Public License
+ * as published by the Free Software Foundation; either version 2
+ * of the License, or (at your option) any later version.
+ *
+ * This program is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+ * GNU General Public License for more details.
+ *
+ * You should have received a copy of the GNU General Public License
+ * along with this program; if not, write to the Free Software Foundation,
+ * Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
+ *
+ * Contributor: 
+ *		Jeroen Bakker 
+ *		Monique Dewanchand
+ */
+
+#include "OCL_opencl.h"
+
+void OCL_init(void)
+{
+#ifdef _WIN32
+	const char *path = "OpenCL.dll";
+#elif defined(__APPLE__)
+	const char *path = "/Library/Frameworks/OpenCL.framework/OpenCL";
+#else
+	const char *path = "libOpenCL.so";
+#endif
+
+	clewInit(path);
+}
+
diff --git a/intern/opencl/intern/clew.c b/intern/opencl/intern/clew.c
new file mode 100644
index 00000000000..d68eb17288f
--- /dev/null
+++ b/intern/opencl/intern/clew.c
@@ -0,0 +1,311 @@
+//////////////////////////////////////////////////////////////////////////
+//  Copyright (c) 2009 Organic Vectory B.V.
+//  Written by George van Venrooij
+//
+//  Distributed under the Boost Software License, Version 1.0.
+//  (See accompanying file license.txt)
+//////////////////////////////////////////////////////////////////////////
+
+#include "clew.h"
+
+//! \file clew.c
+//! \brief OpenCL run-time loader source
+
+#ifndef CLCC_GENERATE_DOCUMENTATION
+#ifdef _WIN32
+    #define WIN32_LEAN_AND_MEAN
+    #define VC_EXTRALEAN
+    #include <windows.h>
+
+    typedef HMODULE             CLCC_DYNLIB_HANDLE;
+
+    #define CLCC_DYNLIB_OPEN    LoadLibrary
+    #define CLCC_DYNLIB_CLOSE   FreeLibrary
+    #define CLCC_DYNLIB_IMPORT  GetProcAddress
+#else
+    #include <dlfcn.h>
+    
+    typedef void*                   CLCC_DYNLIB_HANDLE;
+
+    #define CLCC_DYNLIB_OPEN(path)  dlopen(path, RTLD_NOW | RTLD_GLOBAL)
+    #define CLCC_DYNLIB_CLOSE       dlclose
+    #define CLCC_DYNLIB_IMPORT      dlsym
+#endif
+#else
+    //typedef implementation_defined  CLCC_DYNLIB_HANDLE;
+    //#define CLCC_DYNLIB_OPEN(path)  implementation_defined
+    //#define CLCC_DYNLIB_CLOSE       implementation_defined
+    //#define CLCC_DYNLIB_IMPORT      implementation_defined
+#endif
+
+#include <stdlib.h>
+
+//! \brief module handle
+static CLCC_DYNLIB_HANDLE module = NULL;
+
+//  Variables holding function entry points
+#ifndef CLCC_GENERATE_DOCUMENTATION
+PFNCLGETPLATFORMIDS                 __oclGetPlatformIDs                = NULL;
+PFNCLGETPLATFORMINFO                __oclGetPlatformInfo               = NULL;
+PFNCLGETDEVICEIDS                   __oclGetDeviceIDs                  = NULL;
+PFNCLGETDEVICEINFO                  __oclGetDeviceInfo                 = NULL;
+PFNCLCREATECONTEXT                  __oclCreateContext                 = NULL;
+PFNCLCREATECONTEXTFROMTYPE          __oclCreateContextFromType         = NULL;
+PFNCLRETAINCONTEXT                  __oclRetainContext                 = NULL;
+PFNCLRELEASECONTEXT                 __oclReleaseContext                = NULL;
+PFNCLGETCONTEXTINFO                 __oclGetContextInfo                = NULL;
+PFNCLCREATECOMMANDQUEUE             __oclCreateCommandQueue            = NULL;
+PFNCLRETAINCOMMANDQUEUE             __oclRetainCommandQueue            = NULL;
+PFNCLRELEASECOMMANDQUEUE            __oclReleaseCommandQueue           = NULL;
+PFNCLGETCOMMANDQUEUEINFO            __oclGetCommandQueueInfo           = NULL;
+PFNCLSETCOMMANDQUEUEPROPERTY        __oclSetCommandQueueProperty       = NULL;
+PFNCLCREATEBUFFER                   __oclCreateBuffer                  = NULL;
+PFNCLCREATEIMAGE2D                  __oclCreateImage2D                 = NULL;
+PFNCLCREATEIMAGE3D                  __oclCreateImage3D                 = NULL;
+PFNCLRETAINMEMOBJECT                __oclRetainMemObject               = NULL;
+PFNCLRELEASEMEMOBJECT               __oclReleaseMemObject              = NULL;
+PFNCLGETSUPPORTEDIMAGEFORMATS       __oclGetSupportedImageFormats      = NULL;
+PFNCLGETMEMOBJECTINFO               __oclGetMemObjectInfo              = NULL;
+PFNCLGETIMAGEINFO                   __oclGetImageInfo                  = NULL;
+PFNCLCREATESAMPLER                  __oclCreateSampler                 = NULL;
+PFNCLRETAINSAMPLER                  __oclRetainSampler                 = NULL;
+PFNCLRELEASESAMPLER                 __oclReleaseSampler                = NULL;
+PFNCLGETSAMPLERINFO                 __oclGetSamplerInfo                = NULL;
+PFNCLCREATEPROGRAMWITHSOURCE        __oclCreateProgramWithSource       = NULL;
+PFNCLCREATEPROGRAMWITHBINARY        __oclCreateProgramWithBinary       = NULL;
+PFNCLRETAINPROGRAM                  __oclRetainProgram                 = NULL;
+PFNCLRELEASEPROGRAM                 __oclReleaseProgram                = NULL;
+PFNCLBUILDPROGRAM                   __oclBuildProgram                  = NULL;
+PFNCLUNLOADCOMPILER                 __oclUnloadCompiler                = NULL;
+PFNCLGETPROGRAMINFO                 __oclGetProgramInfo                = NULL;
+PFNCLGETPROGRAMBUILDINFO            __oclGetProgramBuildInfo           = NULL;
+PFNCLCREATEKERNEL                   __oclCreateKernel                  = NULL;
+PFNCLCREATEKERNELSINPROGRAM         __oclCreateKernelsInProgram        = NULL;
+PFNCLRETAINKERNEL                   __oclRetainKernel                  = NULL;
+PFNCLRELEASEKERNEL                  __oclReleaseKernel                 = NULL;
+PFNCLSETKERNELARG                   __oclSetKernelArg                  = NULL;
+PFNCLGETKERNELINFO                  __oclGetKernelInfo                 = NULL;
+PFNCLGETKERNELWORKGROUPINFO         __oclGetKernelWorkGroupInfo        = NULL;
+PFNCLWAITFOREVENTS                  __oclWaitForEvents                 = NULL;
+PFNCLGETEVENTINFO                   __oclGetEventInfo                  = NULL;
+PFNCLRETAINEVENT                    __oclRetainEvent                   = NULL;
+PFNCLRELEASEEVENT                   __oclReleaseEvent                  = NULL;
+PFNCLGETEVENTPROFILINGINFO          __oclGetEventProfilingInfo         = NULL;
+PFNCLFLUSH                          __oclFlush                         = NULL;
+PFNCLFINISH                         __oclFinish                        = NULL;
+PFNCLENQUEUEREADBUFFER              __oclEnqueueReadBuffer             = NULL;
+PFNCLENQUEUEWRITEBUFFER             __oclEnqueueWriteBuffer            = NULL;
+PFNCLENQUEUECOPYBUFFER              __oclEnqueueCopyBuffer             = NULL;
+PFNCLENQUEUEREADIMAGE               __oclEnqueueReadImage              = NULL;
+PFNCLENQUEUEWRITEIMAGE              __oclEnqueueWriteImage             = NULL;
+PFNCLENQUEUECOPYIMAGE               __oclEnqueueCopyImage              = NULL;
+PFNCLENQUEUECOPYIMAGETOBUFFER       __oclEnqueueCopyImageToBuffer      = NULL;
+PFNCLENQUEUECOPYBUFFERTOIMAGE       __oclEnqueueCopyBufferToImage      = NULL;
+PFNCLENQUEUEMAPBUFFER               __oclEnqueueMapBuffer              = NULL;
+PFNCLENQUEUEMAPIMAGE                __oclEnqueueMapImage               = NULL;
+PFNCLENQUEUEUNMAPMEMOBJECT          __oclEnqueueUnmapMemObject         = NULL;
+PFNCLENQUEUENDRANGEKERNEL           __oclEnqueueNDRangeKernel          = NULL;
+PFNCLENQUEUETASK                    __oclEnqueueTask                   = NULL;
+PFNCLENQUEUENATIVEKERNEL            __oclEnqueueNativeKernel           = NULL;
+PFNCLENQUEUEMARKER                  __oclEnqueueMarker                 = NULL;
+PFNCLENQUEUEWAITFOREVENTS           __oclEnqueueWaitForEvents          = NULL;
+PFNCLENQUEUEBARRIER                 __oclEnqueueBarrier                = NULL;
+PFNCLGETEXTENSIONFUNCTIONADDRESS    __oclGetExtensionFunctionAddress   = NULL;
+#endif  //  CLCC_GENERATE_DOCUMENTATION
+
+
+//! \brief Unloads OpenCL dynamic library, should not be called directly
+static void clewExit(void)
+{
+    if (module != NULL)
+    {
+        //  Ignore errors
+        CLCC_DYNLIB_CLOSE(module);
+        module = NULL;
+    }
+}
+
+//! \param path path to dynamic library to load
+//! \return CLEW_ERROR_OPEN_FAILED if the library could not be opened
+//! CLEW_ERROR_ATEXIT_FAILED if atexit(clewExit) failed
+//! CLEW_SUCCESS when the library was succesfully loaded
+int clewInit(const char* path)
+{
+    int error = 0;
+
+    //  Check if already initialized
+    if (module != NULL)
+    {
+        return CLEW_SUCCESS;
+    }
+
+    //  Load library
+    module = CLCC_DYNLIB_OPEN(path);
+
+    //  Check for errors
+    if (module == NULL)
+    {
+        return CLEW_ERROR_OPEN_FAILED;
+    }
+
+    //  Set unloading
+    error = atexit(clewExit);
+
+    if (error)
+    {
+        //  Failure queing atexit, shutdown with error
+        CLCC_DYNLIB_CLOSE(module);
+        module = NULL;
+
+        return CLEW_ERROR_ATEXIT_FAILED;
+    }
+
+    //  Determine function entry-points
+    __oclGetPlatformIDs                = (PFNCLGETPLATFORMIDS              )CLCC_DYNLIB_IMPORT(module, "clGetPlatformIDs");
+    __oclGetPlatformInfo               = (PFNCLGETPLATFORMINFO             )CLCC_DYNLIB_IMPORT(module, "clGetPlatformInfo");
+    __oclGetDeviceIDs                  = (PFNCLGETDEVICEIDS                )CLCC_DYNLIB_IMPORT(module, "clGetDeviceIDs");
+    __oclGetDeviceInfo                 = (PFNCLGETDEVICEINFO               )CLCC_DYNLIB_IMPORT(module, "clGetDeviceInfo");
+    __oclCreateContext                 = (PFNCLCREATECONTEXT               )CLCC_DYNLIB_IMPORT(module, "clCreateContext");
+    __oclCreateContextFromType         = (PFNCLCREATECONTEXTFROMTYPE       )CLCC_DYNLIB_IMPORT(module, "clCreateContextFromType");
+    __oclRetainContext                 = (PFNCLRETAINCONTEXT               )CLCC_DYNLIB_IMPORT(module, "clRetainContext");
+    __oclReleaseContext                = (PFNCLRELEASECONTEXT              )CLCC_DYNLIB_IMPORT(module, "clReleaseContext");
+    __oclGetContextInfo                = (PFNCLGETCONTEXTINFO              )CLCC_DYNLIB_IMPORT(module, "clGetContextInfo");
+    __oclCreateCommandQueue            = (PFNCLCREATECOMMANDQUEUE          )CLCC_DYNLIB_IMPORT(module, "clCreateCommandQueue");
+    __oclRetainCommandQueue            = (PFNCLRETAINCOMMANDQUEUE          )CLCC_DYNLIB_IMPORT(module, "clRetainCommandQueue");
+    __oclReleaseCommandQueue           = (PFNCLRELEASECOMMANDQUEUE         )CLCC_DYNLIB_IMPORT(module, "clReleaseCommandQueue");
+    __oclGetCommandQueueInfo           = (PFNCLGETCOMMANDQUEUEINFO         )CLCC_DYNLIB_IMPORT(module, "clGetCommandQueueInfo");
+    __oclSetCommandQueueProperty       = (PFNCLSETCOMMANDQUEUEPROPERTY     )CLCC_DYNLIB_IMPORT(module, "clSetCommandQueueProperty");
+    __oclCreateBuffer                  = (PFNCLCREATEBUFFER                )CLCC_DYNLIB_IMPORT(module, "clCreateBuffer");
+    __oclCreateImage2D                 = (PFNCLCREATEIMAGE2D               )CLCC_DYNLIB_IMPORT(module, "clCreateImage2D");
+    __oclCreateImage3D                 = (PFNCLCREATEIMAGE3D               )CLCC_DYNLIB_IMPORT(module, "clCreateImage3D");
+    __oclRetainMemObject               = (PFNCLRETAINMEMOBJECT             )CLCC_DYNLIB_IMPORT(module, "clRetainMemObject");
+    __oclReleaseMemObject              = (PFNCLRELEASEMEMOBJECT            )CLCC_DYNLIB_IMPORT(module, "clReleaseMemObject");
+    __oclGetSupportedImageFormats      = (PFNCLGETSUPPORTEDIMAGEFORMATS    )CLCC_DYNLIB_IMPORT(module, "clGetSupportedImageFormats");
+    __oclGetMemObjectInfo              = (PFNCLGETMEMOBJECTINFO            )CLCC_DYNLIB_IMPORT(module, "clGetMemObjectInfo");
+    __oclGetImageInfo                  = (PFNCLGETIMAGEINFO                )CLCC_DYNLIB_IMPORT(module, "clGetImageInfo");
+    __oclCreateSampler                 = (PFNCLCREATESAMPLER               )CLCC_DYNLIB_IMPORT(module, "clCreateSampler");
+    __oclRetainSampler                 = (PFNCLRETAINSAMPLER               )CLCC_DYNLIB_IMPORT(module, "clRetainSampler");
+    __oclReleaseSampler                = (PFNCLRELEASESAMPLER              )CLCC_DYNLIB_IMPORT(module, "clReleaseSampler");
+    __oclGetSamplerInfo                = (PFNCLGETSAMPLERINFO              )CLCC_DYNLIB_IMPORT(module, "clGetSamplerInfo");
+    __oclCreateProgramWithSource       = (PFNCLCREATEPROGRAMWITHSOURCE     )CLCC_DYNLIB_IMPORT(module, "clCreateProgramWithSource");
+    __oclCreateProgramWithBinary       = (PFNCLCREATEPROGRAMWITHBINARY     )CLCC_DYNLIB_IMPORT(module, "clCreateProgramWithBinary");
+    __oclRetainProgram                 = (PFNCLRETAINPROGRAM               )CLCC_DYNLIB_IMPORT(module, "clRetainProgram");
+    __oclReleaseProgram                = (PFNCLRELEASEPROGRAM              )CLCC_DYNLIB_IMPORT(module, "clReleaseProgram");
+    __oclBuildProgram                  = (PFNCLBUILDPROGRAM                )CLCC_DYNLIB_IMPORT(module, "clBuildProgram");
+    __oclUnloadCompiler                = (PFNCLUNLOADCOMPILER              )CLCC_DYNLIB_IMPORT(module, "clUnloadCompiler");
+    __oclGetProgramInfo                = (PFNCLGETPROGRAMINFO              )CLCC_DYNLIB_IMPORT(module, "clGetProgramInfo");
+    __oclGetProgramBuildInfo           = (PFNCLGETPROGRAMBUILDINFO         )CLCC_DYNLIB_IMPORT(module, "clGetProgramBuildInfo");
+    __oclCreateKernel                  = (PFNCLCREATEKERNEL                )CLCC_DYNLIB_IMPORT(module, "clCreateKernel");
+    __oclCreateKernelsInProgram        = (PFNCLCREATEKERNELSINPROGRAM      )CLCC_DYNLIB_IMPORT(module, "clCreateKernelsInProgram");
+    __oclRetainKernel                  = (PFNCLRETAINKERNEL                )CLCC_DYNLIB_IMPORT(module, "clRetainKernel");
+    __oclReleaseKernel                 = (PFNCLRELEASEKERNEL               )CLCC_DYNLIB_IMPORT(module, "clReleaseKernel");
+    __oclSetKernelArg                  = (PFNCLSETKERNELARG                )CLCC_DYNLIB_IMPORT(module, "clSetKernelArg");
+    __oclGetKernelInfo                 = (PFNCLGETKERNELINFO               )CLCC_DYNLIB_IMPORT(module, "clGetKernelInfo");
+    __oclGetKernelWorkGroupInfo        = (PFNCLGETKERNELWORKGROUPINFO      )CLCC_DYNLIB_IMPORT(module, "clGetKernelWorkGroupInfo");
+    __oclWaitForEvents                 = (PFNCLWAITFOREVENTS               )CLCC_DYNLIB_IMPORT(module, "clWaitForEvents");
+    __oclGetEventInfo                  = (PFNCLGETEVENTINFO                )CLCC_DYNLIB_IMPORT(module, "clGetEventInfo");
+    __oclRetainEvent                   = (PFNCLRETAINEVENT                 )CLCC_DYNLIB_IMPORT(module, "clRetainEvent");
+    __oclReleaseEvent                  = (PFNCLRELEASEEVENT                )CLCC_DYNLIB_IMPORT(module, "clReleaseEvent");
+    __oclGetEventProfilingInfo         = (PFNCLGETEVENTPROFILINGINFO       )CLCC_DYNLIB_IMPORT(module, "clGetEventProfilingInfo");
+    __oclFlush                         = (PFNCLFLUSH                       )CLCC_DYNLIB_IMPORT(module, "clFlush");
+    __oclFinish                        = (PFNCLFINISH                      )CLCC_DYNLIB_IMPORT(module, "clFinish");
+    __oclEnqueueReadBuffer             = (PFNCLENQUEUEREADBUFFER           )CLCC_DYNLIB_IMPORT(module, "clEnqueueReadBuffer");
+    __oclEnqueueWriteBuffer            = (PFNCLENQUEUEWRITEBUFFER          )CLCC_DYNLIB_IMPORT(module, "clEnqueueWriteBuffer");
+    __oclEnqueueCopyBuffer             = (PFNCLENQUEUECOPYBUFFER           )CLCC_DYNLIB_IMPORT(module, "clEnqueueCopyBuffer");
+    __oclEnqueueReadImage              = (PFNCLENQUEUEREADIMAGE            )CLCC_DYNLIB_IMPORT(module, "clEnqueueReadImage");
+    __oclEnqueueWriteImage             = (PFNCLENQUEUEWRITEIMAGE           )CLCC_DYNLIB_IMPORT(module, "clEnqueueWriteImage");
+    __oclEnqueueCopyImage              = (PFNCLENQUEUECOPYIMAGE            )CLCC_DYNLIB_IMPORT(module, "clEnqueueCopyImage");
+    __oclEnqueueCopyImageToBuffer      = (PFNCLENQUEUECOPYIMAGETOBUFFER    )CLCC_DYNLIB_IMPORT(module, "clEnqueueCopyImageToBuffer");
+    __oclEnqueueCopyBufferToImage      = (PFNCLENQUEUECOPYBUFFERTOIMAGE    )CLCC_DYNLIB_IMPORT(module, "clEnqueueCopyBufferToImage");
+    __oclEnqueueMapBuffer              = (PFNCLENQUEUEMAPBUFFER            )CLCC_DYNLIB_IMPORT(module, "clEnqueueMapBuffer");
+    __oclEnqueueMapImage               = (PFNCLENQUEUEMAPIMAGE             )CLCC_DYNLIB_IMPORT(module, "clEnqueueMapImage");
+    __oclEnqueueUnmapMemObject         = (PFNCLENQUEUEUNMAPMEMOBJECT       )CLCC_DYNLIB_IMPORT(module, "clEnqueueUnmapMemObject");
+    __oclEnqueueNDRangeKernel          = (PFNCLENQUEUENDRANGEKERNEL        )CLCC_DYNLIB_IMPORT(module, "clEnqueueNDRangeKernel");
+    __oclEnqueueTask                   = (PFNCLENQUEUETASK                 )CLCC_DYNLIB_IMPORT(module, "clEnqueueTask");
+    __oclEnqueueNativeKernel           = (PFNCLENQUEUENATIVEKERNEL         )CLCC_DYNLIB_IMPORT(module, "clEnqueueNativeKernel");
+    __oclEnqueueMarker                 = (PFNCLENQUEUEMARKER               )CLCC_DYNLIB_IMPORT(module, "clEnqueueMarker");
+    __oclEnqueueWaitForEvents          = (PFNCLENQUEUEWAITFOREVENTS        )CLCC_DYNLIB_IMPORT(module, "clEnqueueWaitForEvents");
+    __oclEnqueueBarrier                = (PFNCLENQUEUEBARRIER              )CLCC_DYNLIB_IMPORT(module, "clEnqueueBarrier");
+    __oclGetExtensionFunctionAddress   = (PFNCLGETEXTENSIONFUNCTIONADDRESS )CLCC_DYNLIB_IMPORT(module, "clGetExtensionFunctionAddress");
+
+    return CLEW_SUCCESS;
+}
+
+//! \param error CL error code
+//! \return a string representation of the error code
+const char* clewErrorString(cl_int error)
+{
+    static const char* strings[] =
+    {
+        // Error Codes
+          "CL_SUCCESS"                                  //   0
+        , "CL_DEVICE_NOT_FOUND"                         //  -1
+        , "CL_DEVICE_NOT_AVAILABLE"                     //  -2
+        , "CL_COMPILER_NOT_AVAILABLE"                   //  -3
+        , "CL_MEM_OBJECT_ALLOCATION_FAILURE"            //  -4
+        , "CL_OUT_OF_RESOURCES"                         //  -5
+        , "CL_OUT_OF_HOST_MEMORY"                       //  -6
+        , "CL_PROFILING_INFO_NOT_AVAILABLE"             //  -7
+        , "CL_MEM_COPY_OVERLAP"                         //  -8
+        , "CL_IMAGE_FORMAT_MISMATCH"                    //  -9
+        , "CL_IMAGE_FORMAT_NOT_SUPPORTED"               //  -10
+        , "CL_BUILD_PROGRAM_FAILURE"                    //  -11
+        , "CL_MAP_FAILURE"                              //  -12
+
+        , ""    //  -13
+        , ""    //  -14
+        , ""    //  -15
+        , ""    //  -16
+        , ""    //  -17
+        , ""    //  -18
+        , ""    //  -19
+
+        , ""    //  -20
+        , ""    //  -21
+        , ""    //  -22
+        , ""    //  -23
+        , ""    //  -24
+        , ""    //  -25
+        , ""    //  -26
+        , ""    //  -27
+        , ""    //  -28
+        , ""    //  -29
+
+        , "CL_INVALID_VALUE"                            //  -30
+        , "CL_INVALID_DEVICE_TYPE"                      //  -31
+        , "CL_INVALID_PLATFORM"                         //  -32
+        , "CL_INVALID_DEVICE"                           //  -33
+        , "CL_INVALID_CONTEXT"                          //  -34
+        , "CL_INVALID_QUEUE_PROPERTIES"                 //  -35
+        , "CL_INVALID_COMMAND_QUEUE"                    //  -36
+        , "CL_INVALID_HOST_PTR"                         //  -37
+        , "CL_INVALID_MEM_OBJECT"                       //  -38
+        , "CL_INVALID_IMAGE_FORMAT_DESCRIPTOR"          //  -39
+        , "CL_INVALID_IMAGE_SIZE"                       //  -40
+        , "CL_INVALID_SAMPLER"                          //  -41
+        , "CL_INVALID_BINARY"                           //  -42
+        , "CL_INVALID_BUILD_OPTIONS"                    //  -43
+        , "CL_INVALID_PROGRAM"                          //  -44
+        , "CL_INVALID_PROGRAM_EXECUTABLE"               //  -45
+        , "CL_INVALID_KERNEL_NAME"                      //  -46
+        , "CL_INVALID_KERNEL_DEFINITION"                //  -47
+        , "CL_INVALID_KERNEL"                           //  -48
+        , "CL_INVALID_ARG_INDEX"                        //  -49
+        , "CL_INVALID_ARG_VALUE"                        //  -50
+        , "CL_INVALID_ARG_SIZE"                         //  -51
+        , "CL_INVALID_KERNEL_ARGS"                      //  -52
+        , "CL_INVALID_WORK_DIMENSION"                   //  -53
+        , "CL_INVALID_WORK_GROUP_SIZE"                  //  -54
+        , "CL_INVALID_WORK_ITEM_SIZE"                   //  -55
+        , "CL_INVALID_GLOBAL_OFFSET"                    //  -56
+        , "CL_INVALID_EVENT_WAIT_LIST"                  //  -57
+        , "CL_INVALID_EVENT"                            //  -58
+        , "CL_INVALID_OPERATION"                        //  -59
+        , "CL_INVALID_GL_OBJECT"                        //  -60
+        , "CL_INVALID_BUFFER_SIZE"                      //  -61
+        , "CL_INVALID_MIP_LEVEL"                        //  -62
+        , "CL_INVALID_GLOBAL_WORK_SIZE"                 //  -63
+    };
+
+    return strings[-error];
+}
diff --git a/intern/opencl/intern/clew.h b/intern/opencl/intern/clew.h
new file mode 100644
index 00000000000..bb7e0134dcf
--- /dev/null
+++ b/intern/opencl/intern/clew.h
@@ -0,0 +1,1317 @@
+#ifndef CLCC_CLEW_HPP_INCLUDED
+#define CLCC_CLEW_HPP_INCLUDED
+
+//////////////////////////////////////////////////////////////////////////
+//  Copyright (c) 2009 Organic Vectory B.V.
+//  Written by George van Venrooij
+//
+//  Distributed under the Boost Software License, Version 1.0.
+//  (See accompanying file license.txt)
+//////////////////////////////////////////////////////////////////////////
+
+//! \file clew.h
+//! \brief OpenCL run-time loader header
+//!
+//! This file contains a copy of the contents of CL.H and CL_PLATFORM.H from the 
+//! official OpenCL spec. The purpose of this code is to load the OpenCL dynamic
+//! library at run-time and thus allow the executable to function on many
+//! platforms regardless of the vendor of the OpenCL driver actually installed.
+//! Some of the techniques used here were inspired by work done in the GLEW
+//! library (http://glew.sourceforge.net/)
+
+//  Run-time dynamic linking functionality based on concepts used in GLEW
+#ifdef  __OPENCL_CL_H
+#error cl.h included before clew.h
+#endif
+
+#ifdef  __OPENCL_CL_PLATFORM_H
+#error cl_platform.h included before clew.h
+#endif
+
+#ifndef CLCC_GENERATE_DOCUMENTATION
+//  Prevent cl.h inclusion
+#define __OPENCL_CL_H
+//  Prevent cl_platform.h inclusion
+#define __CL_PLATFORM_H
+#endif  //  CLCC_GENERATE_DOCUMENTATION
+
+/*******************************************************************************
+* Copyright (c) 2008-2009 The Khronos Group Inc.
+*
+* Permission is hereby granted, free of charge, to any person obtaining a
+* copy of this software and/or associated documentation files (the
+* "Materials"), to deal in the Materials without restriction, including
+* without limitation the rights to use, copy, modify, merge, publish,
+* distribute, sublicense, and/or sell copies of the Materials, and to
+* permit persons to whom the Materials are furnished to do so, subject to
+* the following conditions:
+*
+* The above copyright notice and this permission notice shall be included
+* in all copies or substantial portions of the Materials.
+*
+* THE MATERIALS ARE PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
+* EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF
+* MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.
+* IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY
+* CLAIM, DAMAGES OR OTHER LIABILITY, WHETHER IN AN ACTION OF CONTRACT,
+* TORT OR OTHERWISE, ARISING FROM, OUT OF OR IN CONNECTION WITH THE
+* MATERIALS OR THE USE OR OTHER DEALINGS IN THE MATERIALS.
+******************************************************************************/
+#ifdef __APPLE__
+/* Contains #defines for AVAILABLE_MAC_OS_X_VERSION_10_6_AND_LATER below */
+#include <AvailabilityMacros.h>
+#endif
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+#ifndef CLCC_GENERATE_DOCUMENTATION
+
+#if defined(_WIN32)
+#define CL_API_ENTRY
+#define CL_API_CALL __stdcall
+#else
+#define CL_API_ENTRY
+#define CL_API_CALL
+#endif
+
+#if defined(__APPLE__)
+//JBKK removed for compatibility with blender trunk #define CL_API_SUFFIX__VERSION_1_0   AVAILABLE_MAC_OS_X_VERSION_10_6_AND_LATER
+#define CL_API_SUFFIX__VERSION_1_0
+#define CL_EXTENSION_WEAK_LINK       __attribute__((weak_import))
+#else
+#define CL_API_SUFFIX__VERSION_1_0
+#define CL_EXTENSION_WEAK_LINK                         
+#endif
+
+#if defined(_WIN32) && defined(_MSC_VER)
+
+/* scalar types  */
+typedef signed   __int8         cl_char;
+typedef unsigned __int8         cl_uchar;
+typedef signed   __int16        cl_short;
+typedef unsigned __int16        cl_ushort;
+typedef signed   __int32        cl_int;
+typedef unsigned __int32        cl_uint;
+typedef signed   __int64        cl_long;
+typedef unsigned __int64        cl_ulong;
+
+typedef unsigned __int16        cl_half;
+typedef float                   cl_float;
+typedef double                  cl_double;
+
+
+/*
+* Vector types 
+*
+*  Note:   OpenCL requires that all types be naturally aligned. 
+*          This means that vector types must be naturally aligned.
+*          For example, a vector of four floats must be aligned to
+*          a 16 byte boundary (calculated as 4 * the natural 4-byte 
+*          alignment of the float).  The alignment qualifiers here
+*          will only function properly if your compiler supports them
+*          and if you don't actively work to defeat them.  For example,
+*          in order for a cl_float4 to be 16 byte aligned in a struct,
+*          the start of the struct must itself be 16-byte aligned. 
+*
+*          Maintaining proper alignment is the user's responsibility.
+*/
+typedef signed   __int8          cl_char2[2];
+typedef signed   __int8          cl_char4[4];
+typedef signed   __int8          cl_char8[8];
+typedef signed   __int8          cl_char16[16];
+typedef unsigned __int8         cl_uchar2[2];
+typedef unsigned __int8         cl_uchar4[4];
+typedef unsigned __int8         cl_uchar8[8];
+typedef unsigned __int8         cl_uchar16[16];
+
+typedef signed   __int16         cl_short2[2];
+typedef signed   __int16         cl_short4[4];
+typedef signed   __int16         cl_short8[8];
+typedef signed   __int16         cl_short16[16];
+typedef unsigned __int16        cl_ushort2[2];
+typedef unsigned __int16        cl_ushort4[4];
+typedef unsigned __int16        cl_ushort8[8];
+typedef unsigned __int16        cl_ushort16[16];
+
+typedef signed   __int32         cl_int2[2];
+typedef signed   __int32         cl_int4[4];
+typedef signed   __int32         cl_int8[8];
+typedef signed   __int32         cl_int16[16];
+typedef unsigned __int32        cl_uint2[2];
+typedef unsigned __int32        cl_uint4[4];
+typedef unsigned __int32        cl_uint8[8];
+typedef unsigned __int32        cl_uint16[16];
+
+typedef signed   __int64         cl_long2[2];
+typedef signed   __int64         cl_long4[4];
+typedef signed   __int64         cl_long8[8];
+typedef signed   __int64         cl_long16[16];
+typedef unsigned __int64        cl_ulong2[2];
+typedef unsigned __int64        cl_ulong4[4];
+typedef unsigned __int64        cl_ulong8[8];
+typedef unsigned __int64        cl_ulong16[16];
+
+typedef float           cl_float2[2];
+typedef float           cl_float4[4];
+typedef float           cl_float8[8];
+typedef float           cl_float16[16];
+
+typedef double          cl_double2[2];
+typedef double          cl_double4[4];
+typedef double          cl_double8[8];
+typedef double          cl_double16[16];
+/* There are no vector types for half */
+
+#else
+
+#include <stdint.h>
+
+/* scalar types  */
+typedef int8_t          cl_char;
+typedef uint8_t         cl_uchar;
+typedef int16_t         cl_short    __attribute__((aligned(2)));
+typedef uint16_t        cl_ushort   __attribute__((aligned(2)));
+typedef int32_t         cl_int      __attribute__((aligned(4)));
+typedef uint32_t        cl_uint     __attribute__((aligned(4)));
+typedef int64_t         cl_long     __attribute__((aligned(8)));
+typedef uint64_t        cl_ulong    __attribute__((aligned(8)));
+
+typedef uint16_t        cl_half     __attribute__((aligned(2)));
+typedef float           cl_float    __attribute__((aligned(4)));
+typedef double          cl_double   __attribute__((aligned(8)));
+
+/*
+* Vector types 
+*
+*  Note:   OpenCL requires that all types be naturally aligned. 
+*          This means that vector types must be naturally aligned.
+*          For example, a vector of four floats must be aligned to
+*          a 16 byte boundary (calculated as 4 * the natural 4-byte 
+*          alignment of the float).  The alignment qualifiers here
+*          will only function properly if your compiler supports them
+*          and if you don't actively work to defeat them.  For example,
+*          in order for a cl_float4 to be 16 byte aligned in a struct,
+*          the start of the struct must itself be 16-byte aligned. 
+*
+*          Maintaining proper alignment is the user's responsibility.
+*/
+typedef int8_t          cl_char2[2]     __attribute__((aligned(2)));
+typedef int8_t          cl_char4[4]     __attribute__((aligned(4)));
+typedef int8_t          cl_char8[8]     __attribute__((aligned(8)));
+typedef int8_t          cl_char16[16]   __attribute__((aligned(16)));
+typedef uint8_t         cl_uchar2[2]    __attribute__((aligned(2)));
+typedef uint8_t         cl_uchar4[4]    __attribute__((aligned(4)));
+typedef uint8_t         cl_uchar8[8]    __attribute__((aligned(8)));
+typedef uint8_t         cl_uchar16[16]  __attribute__((aligned(16)));
+
+typedef int16_t         cl_short2[2]     __attribute__((aligned(4)));
+typedef int16_t         cl_short4[4]     __attribute__((aligned(8)));
+typedef int16_t         cl_short8[8]     __attribute__((aligned(16)));
+typedef int16_t         cl_short16[16]   __attribute__((aligned(32)));
+typedef uint16_t        cl_ushort2[2]    __attribute__((aligned(4)));
+typedef uint16_t        cl_ushort4[4]    __attribute__((aligned(8)));
+typedef uint16_t        cl_ushort8[8]    __attribute__((aligned(16)));
+typedef uint16_t        cl_ushort16[16]  __attribute__((aligned(32)));
+
+typedef int32_t         cl_int2[2]      __attribute__((aligned(8)));
+typedef int32_t         cl_int4[4]      __attribute__((aligned(16)));
+typedef int32_t         cl_int8[8]      __attribute__((aligned(32)));
+typedef int32_t         cl_int16[16]    __attribute__((aligned(64)));
+typedef uint32_t        cl_uint2[2]     __attribute__((aligned(8)));
+typedef uint32_t        cl_uint4[4]     __attribute__((aligned(16)));
+typedef uint32_t        cl_uint8[8]     __attribute__((aligned(32)));
+typedef uint32_t        cl_uint16[16]   __attribute__((aligned(64)));
+
+typedef int64_t         cl_long2[2]     __attribute__((aligned(16)));
+typedef int64_t         cl_long4[4]     __attribute__((aligned(32)));
+typedef int64_t         cl_long8[8]     __attribute__((aligned(64)));
+typedef int64_t         cl_long16[16]   __attribute__((aligned(128)));
+typedef uint64_t        cl_ulong2[2]    __attribute__((aligned(16)));
+typedef uint64_t        cl_ulong4[4]    __attribute__((aligned(32)));
+typedef uint64_t        cl_ulong8[8]    __attribute__((aligned(64)));
+typedef uint64_t        cl_ulong16[16]  __attribute__((aligned(128)));
+
+typedef float           cl_float2[2]    __attribute__((aligned(8)));
+typedef float           cl_float4[4]    __attribute__((aligned(16)));
+typedef float           cl_float8[8]    __attribute__((aligned(32)));
+typedef float           cl_float16[16]  __attribute__((aligned(64)));
+
+typedef double          cl_double2[2]   __attribute__((aligned(16)));
+typedef double          cl_double4[4]   __attribute__((aligned(32)));
+typedef double          cl_double8[8]   __attribute__((aligned(64)));
+typedef double          cl_double16[16] __attribute__((aligned(128)));
+
+/* There are no vector types for half */
+
+#endif
+
+/******************************************************************************/
+
+// Macro names and corresponding values defined by OpenCL
+
+#define CL_CHAR_BIT         8
+#define CL_SCHAR_MAX        127
+#define CL_SCHAR_MIN        (-127-1)
+#define CL_CHAR_MAX         CL_SCHAR_MAX
+#define CL_CHAR_MIN         CL_SCHAR_MIN
+#define CL_UCHAR_MAX        255
+#define CL_SHRT_MAX         32767
+#define CL_SHRT_MIN         (-32767-1)
+#define CL_USHRT_MAX        65535
+#define CL_INT_MAX          2147483647
+#define CL_INT_MIN          (-2147483647-1)
+#define CL_UINT_MAX         0xffffffffU
+#define CL_LONG_MAX         ((cl_long) 0x7FFFFFFFFFFFFFFFLL)
+#define CL_LONG_MIN         ((cl_long) -0x7FFFFFFFFFFFFFFFLL - 1LL)
+#define CL_ULONG_MAX        ((cl_ulong) 0xFFFFFFFFFFFFFFFFULL)
+
+#define CL_FLT_DIG          6
+#define CL_FLT_MANT_DIG     24
+#define CL_FLT_MAX_10_EXP   +38
+#define CL_FLT_MAX_EXP      +128
+#define CL_FLT_MIN_10_EXP   -37
+#define CL_FLT_MIN_EXP      -125
+#define CL_FLT_RADIX        2
+#if defined(_MSC_VER)
+// MSVC doesn't understand hex floats
+#define CL_FLT_MAX          3.402823466e+38F
+#define CL_FLT_MIN          1.175494351e-38F
+#define CL_FLT_EPSILON      1.192092896e-07F
+#else
+#define CL_FLT_MAX          0x1.fffffep127f
+#define CL_FLT_MIN          0x1.0p-126f
+#define CL_FLT_EPSILON      0x1.0p-23f
+#endif
+
+#define CL_DBL_DIG          15
+#define CL_DBL_MANT_DIG     53
+#define CL_DBL_MAX_10_EXP   +308
+#define CL_DBL_MAX_EXP      +1024
+#define CL_DBL_MIN_10_EXP   -307
+#define CL_DBL_MIN_EXP      -1021
+#define CL_DBL_RADIX        2
+#if defined(_MSC_VER)
+// MSVC doesn't understand hex floats
+#define CL_DBL_MAX          1.7976931348623158e+308
+#define CL_DBL_MIN          2.2250738585072014e-308
+#define CL_DBL_EPSILON      2.2204460492503131e-016
+#else
+#define CL_DBL_MAX          0x1.fffffffffffffp1023
+#define CL_DBL_MIN          0x1.0p-1022
+#define CL_DBL_EPSILON      0x1.0p-52
+#endif
+
+#include <stddef.h>
+
+
+//  CL.h contents
+/******************************************************************************/
+
+typedef struct _cl_platform_id *    cl_platform_id;
+typedef struct _cl_device_id *      cl_device_id;
+typedef struct _cl_context *        cl_context;
+typedef struct _cl_command_queue *  cl_command_queue;
+typedef struct _cl_mem *            cl_mem;
+typedef struct _cl_program *        cl_program;
+typedef struct _cl_kernel *         cl_kernel;
+typedef struct _cl_event *          cl_event;
+typedef struct _cl_sampler *        cl_sampler;
+
+typedef cl_uint             cl_bool;                     /* WARNING!  Unlike cl_ types in cl_platform.h, cl_bool is not guaranteed to be the same size as the bool in kernels. */ 
+typedef cl_ulong            cl_bitfield;
+typedef cl_bitfield         cl_device_type;
+typedef cl_uint             cl_platform_info;
+typedef cl_uint             cl_device_info;
+typedef cl_bitfield         cl_device_address_info;
+typedef cl_bitfield         cl_device_fp_config;
+typedef cl_uint             cl_device_mem_cache_type;
+typedef cl_uint             cl_device_local_mem_type;
+typedef cl_bitfield         cl_device_exec_capabilities;
+typedef cl_bitfield         cl_command_queue_properties;
+
+typedef intptr_t			cl_context_properties;
+typedef cl_uint             cl_context_info;
+typedef cl_uint             cl_command_queue_info;
+typedef cl_uint             cl_channel_order;
+typedef cl_uint             cl_channel_type;
+typedef cl_bitfield         cl_mem_flags;
+typedef cl_uint             cl_mem_object_type;
+typedef cl_uint             cl_mem_info;
+typedef cl_uint             cl_image_info;
+typedef cl_uint             cl_addressing_mode;
+typedef cl_uint             cl_filter_mode;
+typedef cl_uint             cl_sampler_info;
+typedef cl_bitfield         cl_map_flags;
+typedef cl_uint             cl_program_info;
+typedef cl_uint             cl_program_build_info;
+typedef cl_int              cl_build_status;
+typedef cl_uint             cl_kernel_info;
+typedef cl_uint             cl_kernel_work_group_info;
+typedef cl_uint             cl_event_info;
+typedef cl_uint             cl_command_type;
+typedef cl_uint             cl_profiling_info;
+
+typedef struct _cl_image_format {
+    cl_channel_order        image_channel_order;
+    cl_channel_type         image_channel_data_type;
+} cl_image_format;
+
+
+
+/******************************************************************************/
+
+// Error Codes
+#define CL_SUCCESS                                  0
+#define CL_DEVICE_NOT_FOUND                         -1
+#define CL_DEVICE_NOT_AVAILABLE                     -2
+#define CL_COMPILER_NOT_AVAILABLE                   -3
+#define CL_MEM_OBJECT_ALLOCATION_FAILURE            -4
+#define CL_OUT_OF_RESOURCES                         -5
+#define CL_OUT_OF_HOST_MEMORY                       -6
+#define CL_PROFILING_INFO_NOT_AVAILABLE             -7
+#define CL_MEM_COPY_OVERLAP                         -8
+#define CL_IMAGE_FORMAT_MISMATCH                    -9
+#define CL_IMAGE_FORMAT_NOT_SUPPORTED               -10
+#define CL_BUILD_PROGRAM_FAILURE                    -11
+#define CL_MAP_FAILURE                              -12
+
+#define CL_INVALID_VALUE                            -30
+#define CL_INVALID_DEVICE_TYPE                      -31
+#define CL_INVALID_PLATFORM                         -32
+#define CL_INVALID_DEVICE                           -33
+#define CL_INVALID_CONTEXT                          -34
+#define CL_INVALID_QUEUE_PROPERTIES                 -35
+#define CL_INVALID_COMMAND_QUEUE                    -36
+#define CL_INVALID_HOST_PTR                         -37
+#define CL_INVALID_MEM_OBJECT                       -38
+#define CL_INVALID_IMAGE_FORMAT_DESCRIPTOR          -39
+#define CL_INVALID_IMAGE_SIZE                       -40
+#define CL_INVALID_SAMPLER                          -41
+#define CL_INVALID_BINARY                           -42
+#define CL_INVALID_BUILD_OPTIONS                    -43
+#define CL_INVALID_PROGRAM                          -44
+#define CL_INVALID_PROGRAM_EXECUTABLE               -45
+#define CL_INVALID_KERNEL_NAME                      -46
+#define CL_INVALID_KERNEL_DEFINITION                -47
+#define CL_INVALID_KERNEL                           -48
+#define CL_INVALID_ARG_INDEX                        -49
+#define CL_INVALID_ARG_VALUE                        -50
+#define CL_INVALID_ARG_SIZE                         -51
+#define CL_INVALID_KERNEL_ARGS                      -52
+#define CL_INVALID_WORK_DIMENSION                   -53
+#define CL_INVALID_WORK_GROUP_SIZE                  -54
+#define CL_INVALID_WORK_ITEM_SIZE                   -55
+#define CL_INVALID_GLOBAL_OFFSET                    -56
+#define CL_INVALID_EVENT_WAIT_LIST                  -57
+#define CL_INVALID_EVENT                            -58
+#define CL_INVALID_OPERATION                        -59
+#define CL_INVALID_GL_OBJECT                        -60
+#define CL_INVALID_BUFFER_SIZE                      -61
+#define CL_INVALID_MIP_LEVEL                        -62
+#define CL_INVALID_GLOBAL_WORK_SIZE                 -63
+
+// OpenCL Version
+#define CL_VERSION_1_0                              1
+
+// cl_bool
+#define CL_FALSE                                    0
+#define CL_TRUE                                     1
+
+// cl_platform_info
+#define CL_PLATFORM_PROFILE                         0x0900
+#define CL_PLATFORM_VERSION                         0x0901
+#define CL_PLATFORM_NAME                            0x0902
+#define CL_PLATFORM_VENDOR                          0x0903
+#define CL_PLATFORM_EXTENSIONS                      0x0904
+
+// cl_device_type - bitfield
+#define CL_DEVICE_TYPE_DEFAULT                      (1 << 0)
+#define CL_DEVICE_TYPE_CPU                          (1 << 1)
+#define CL_DEVICE_TYPE_GPU                          (1 << 2)
+#define CL_DEVICE_TYPE_ACCELERATOR                  (1 << 3)
+#define CL_DEVICE_TYPE_ALL                          0xFFFFFFFF
+
+// cl_device_info
+#define CL_DEVICE_TYPE                              0x1000
+#define CL_DEVICE_VENDOR_ID                         0x1001
+#define CL_DEVICE_MAX_COMPUTE_UNITS                 0x1002
+#define CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS          0x1003
+#define CL_DEVICE_MAX_WORK_GROUP_SIZE               0x1004
+#define CL_DEVICE_MAX_WORK_ITEM_SIZES               0x1005
+#define CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR       0x1006
+#define CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT      0x1007
+#define CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT        0x1008
+#define CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG       0x1009
+#define CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT      0x100A
+#define CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE     0x100B
+#define CL_DEVICE_MAX_CLOCK_FREQUENCY               0x100C
+#define CL_DEVICE_ADDRESS_BITS                      0x100D
+#define CL_DEVICE_MAX_READ_IMAGE_ARGS               0x100E
+#define CL_DEVICE_MAX_WRITE_IMAGE_ARGS              0x100F
+#define CL_DEVICE_MAX_MEM_ALLOC_SIZE                0x1010
+#define CL_DEVICE_IMAGE2D_MAX_WIDTH                 0x1011
+#define CL_DEVICE_IMAGE2D_MAX_HEIGHT                0x1012
+#define CL_DEVICE_IMAGE3D_MAX_WIDTH                 0x1013
+#define CL_DEVICE_IMAGE3D_MAX_HEIGHT                0x1014
+#define CL_DEVICE_IMAGE3D_MAX_DEPTH                 0x1015
+#define CL_DEVICE_IMAGE_SUPPORT                     0x1016
+#define CL_DEVICE_MAX_PARAMETER_SIZE                0x1017
+#define CL_DEVICE_MAX_SAMPLERS                      0x1018
+#define CL_DEVICE_MEM_BASE_ADDR_ALIGN               0x1019
+#define CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE          0x101A
+#define CL_DEVICE_SINGLE_FP_CONFIG                  0x101B
+#define CL_DEVICE_GLOBAL_MEM_CACHE_TYPE             0x101C
+#define CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE         0x101D
+#define CL_DEVICE_GLOBAL_MEM_CACHE_SIZE             0x101E
+#define CL_DEVICE_GLOBAL_MEM_SIZE                   0x101F
+#define CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE          0x1020
+#define CL_DEVICE_MAX_CONSTANT_ARGS                 0x1021
+#define CL_DEVICE_LOCAL_MEM_TYPE                    0x1022
+#define CL_DEVICE_LOCAL_MEM_SIZE                    0x1023
+#define CL_DEVICE_ERROR_CORRECTION_SUPPORT          0x1024
+#define CL_DEVICE_PROFILING_TIMER_RESOLUTION        0x1025
+#define CL_DEVICE_ENDIAN_LITTLE                     0x1026
+#define CL_DEVICE_AVAILABLE                         0x1027
+#define CL_DEVICE_COMPILER_AVAILABLE                0x1028
+#define CL_DEVICE_EXECUTION_CAPABILITIES            0x1029
+#define CL_DEVICE_QUEUE_PROPERTIES                  0x102A
+#define CL_DEVICE_NAME                              0x102B
+#define CL_DEVICE_VENDOR                            0x102C
+#define CL_DRIVER_VERSION                           0x102D
+#define CL_DEVICE_PROFILE                           0x102E
+#define CL_DEVICE_VERSION                           0x102F
+#define CL_DEVICE_EXTENSIONS                        0x1030
+#define CL_DEVICE_PLATFORM                          0x1031
+
+// cl_device_fp_config - bitfield
+#define CL_FP_DENORM                                (1 << 0)
+#define CL_FP_INF_NAN                               (1 << 1)
+#define CL_FP_ROUND_TO_NEAREST                      (1 << 2)
+#define CL_FP_ROUND_TO_ZERO                         (1 << 3)
+#define CL_FP_ROUND_TO_INF                          (1 << 4)
+#define CL_FP_FMA                                   (1 << 5)
+
+// cl_device_mem_cache_type
+#define CL_NONE                                     0x0
+#define CL_READ_ONLY_CACHE                          0x1
+#define CL_READ_WRITE_CACHE                         0x2
+
+// cl_device_local_mem_type
+#define CL_LOCAL                                    0x1
+#define CL_GLOBAL                                   0x2
+
+// cl_device_exec_capabilities - bitfield
+#define CL_EXEC_KERNEL                              (1 << 0)
+#define CL_EXEC_NATIVE_KERNEL                       (1 << 1)
+
+// cl_command_queue_properties - bitfield
+#define CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE      (1 << 0)
+#define CL_QUEUE_PROFILING_ENABLE                   (1 << 1)
+
+// cl_context_info
+#define CL_CONTEXT_REFERENCE_COUNT                  0x1080
+#define CL_CONTEXT_DEVICES                          0x1081
+#define CL_CONTEXT_PROPERTIES                       0x1082
+
+// cl_context_properties
+#define CL_CONTEXT_PLATFORM                         0x1084
+
+// cl_command_queue_info
+#define CL_QUEUE_CONTEXT                            0x1090
+#define CL_QUEUE_DEVICE                             0x1091
+#define CL_QUEUE_REFERENCE_COUNT                    0x1092
+#define CL_QUEUE_PROPERTIES                         0x1093
+
+// cl_mem_flags - bitfield
+#define CL_MEM_READ_WRITE                           (1 << 0)
+#define CL_MEM_WRITE_ONLY                           (1 << 1)
+#define CL_MEM_READ_ONLY                            (1 << 2)
+#define CL_MEM_USE_HOST_PTR                         (1 << 3)
+#define CL_MEM_ALLOC_HOST_PTR                       (1 << 4)
+#define CL_MEM_COPY_HOST_PTR                        (1 << 5)
+
+// cl_channel_order
+#define CL_R                                        0x10B0
+#define CL_A                                        0x10B1
+#define CL_RG                                       0x10B2
+#define CL_RA                                       0x10B3
+#define CL_RGB                                      0x10B4
+#define CL_RGBA                                     0x10B5
+#define CL_BGRA                                     0x10B6
+#define CL_ARGB                                     0x10B7
+#define CL_INTENSITY                                0x10B8
+#define CL_LUMINANCE                                0x10B9
+
+// cl_channel_type
+#define CL_SNORM_INT8                               0x10D0
+#define CL_SNORM_INT16                              0x10D1
+#define CL_UNORM_INT8                               0x10D2
+#define CL_UNORM_INT16                              0x10D3
+#define CL_UNORM_SHORT_565                          0x10D4
+#define CL_UNORM_SHORT_555                          0x10D5
+#define CL_UNORM_INT_101010                         0x10D6
+#define CL_SIGNED_INT8                              0x10D7
+#define CL_SIGNED_INT16                             0x10D8
+#define CL_SIGNED_INT32                             0x10D9
+#define CL_UNSIGNED_INT8                            0x10DA
+#define CL_UNSIGNED_INT16                           0x10DB
+#define CL_UNSIGNED_INT32                           0x10DC
+#define CL_HALF_FLOAT                               0x10DD
+#define CL_FLOAT                                    0x10DE
+
+// cl_mem_object_type
+#define CL_MEM_OBJECT_BUFFER                        0x10F0
+#define CL_MEM_OBJECT_IMAGE2D                       0x10F1
+#define CL_MEM_OBJECT_IMAGE3D                       0x10F2
+
+// cl_mem_info
+#define CL_MEM_TYPE                                 0x1100
+#define CL_MEM_FLAGS                                0x1101
+#define CL_MEM_SIZE                                 0x1102
+#define CL_MEM_HOST_PTR                             0x1103
+#define CL_MEM_MAP_COUNT                            0x1104
+#define CL_MEM_REFERENCE_COUNT                      0x1105
+#define CL_MEM_CONTEXT                              0x1106
+
+// cl_image_info
+#define CL_IMAGE_FORMAT                             0x1110
+#define CL_IMAGE_ELEMENT_SIZE                       0x1111
+#define CL_IMAGE_ROW_PITCH                          0x1112
+#define CL_IMAGE_SLICE_PITCH                        0x1113
+#define CL_IMAGE_WIDTH                              0x1114
+#define CL_IMAGE_HEIGHT                             0x1115
+#define CL_IMAGE_DEPTH                              0x1116
+
+// cl_addressing_mode
+#define CL_ADDRESS_NONE                             0x1130
+#define CL_ADDRESS_CLAMP_TO_EDGE                    0x1131
+#define CL_ADDRESS_CLAMP                            0x1132
+#define CL_ADDRESS_REPEAT                           0x1133
+
+// cl_filter_mode
+#define CL_FILTER_NEAREST                           0x1140
+#define CL_FILTER_LINEAR                            0x1141
+
+// cl_sampler_info
+#define CL_SAMPLER_REFERENCE_COUNT                  0x1150
+#define CL_SAMPLER_CONTEXT                          0x1151
+#define CL_SAMPLER_NORMALIZED_COORDS                0x1152
+#define CL_SAMPLER_ADDRESSING_MODE                  0x1153
+#define CL_SAMPLER_FILTER_MODE                      0x1154
+
+// cl_map_flags - bitfield
+#define CL_MAP_READ                                 (1 << 0)
+#define CL_MAP_WRITE                                (1 << 1)
+
+// cl_program_info
+#define CL_PROGRAM_REFERENCE_COUNT                  0x1160
+#define CL_PROGRAM_CONTEXT                          0x1161
+#define CL_PROGRAM_NUM_DEVICES                      0x1162
+#define CL_PROGRAM_DEVICES                          0x1163
+#define CL_PROGRAM_SOURCE                           0x1164
+#define CL_PROGRAM_BINARY_SIZES                     0x1165
+#define CL_PROGRAM_BINARIES                         0x1166
+
+// cl_program_build_info
+#define CL_PROGRAM_BUILD_STATUS                     0x1181
+#define CL_PROGRAM_BUILD_OPTIONS                    0x1182
+#define CL_PROGRAM_BUILD_LOG                        0x1183
+
+// cl_build_status
+#define CL_BUILD_SUCCESS                            0
+#define CL_BUILD_NONE                               -1
+#define CL_BUILD_ERROR                              -2
+#define CL_BUILD_IN_PROGRESS                        -3
+
+// cl_kernel_info
+#define CL_KERNEL_FUNCTION_NAME                     0x1190
+#define CL_KERNEL_NUM_ARGS                          0x1191
+#define CL_KERNEL_REFERENCE_COUNT                   0x1192
+#define CL_KERNEL_CONTEXT                           0x1193
+#define CL_KERNEL_PROGRAM                           0x1194
+
+// cl_kernel_work_group_info
+#define CL_KERNEL_WORK_GROUP_SIZE                   0x11B0
+#define CL_KERNEL_COMPILE_WORK_GROUP_SIZE           0x11B1
+#define CL_KERNEL_LOCAL_MEM_SIZE                    0x11B2
+
+// cl_event_info
+#define CL_EVENT_COMMAND_QUEUE                      0x11D0
+#define CL_EVENT_COMMAND_TYPE                       0x11D1
+#define CL_EVENT_REFERENCE_COUNT                    0x11D2
+#define CL_EVENT_COMMAND_EXECUTION_STATUS           0x11D3
+
+// cl_command_type
+#define CL_COMMAND_NDRANGE_KERNEL                   0x11F0
+#define CL_COMMAND_TASK                             0x11F1
+#define CL_COMMAND_NATIVE_KERNEL                    0x11F2
+#define CL_COMMAND_READ_BUFFER                      0x11F3
+#define CL_COMMAND_WRITE_BUFFER                     0x11F4
+#define CL_COMMAND_COPY_BUFFER                      0x11F5
+#define CL_COMMAND_READ_IMAGE                       0x11F6
+#define CL_COMMAND_WRITE_IMAGE                      0x11F7
+#define CL_COMMAND_COPY_IMAGE                       0x11F8
+#define CL_COMMAND_COPY_IMAGE_TO_BUFFER             0x11F9
+#define CL_COMMAND_COPY_BUFFER_TO_IMAGE             0x11FA
+#define CL_COMMAND_MAP_BUFFER                       0x11FB
+#define CL_COMMAND_MAP_IMAGE                        0x11FC
+#define CL_COMMAND_UNMAP_MEM_OBJECT                 0x11FD
+#define CL_COMMAND_MARKER                           0x11FE
+#define CL_COMMAND_ACQUIRE_GL_OBJECTS               0x11FF
+#define CL_COMMAND_RELEASE_GL_OBJECTS               0x1200
+
+// command execution status
+#define CL_COMPLETE                                 0x0
+#define CL_RUNNING                                  0x1
+#define CL_SUBMITTED                                0x2
+#define CL_QUEUED                                   0x3
+
+// cl_profiling_info
+#define CL_PROFILING_COMMAND_QUEUED                 0x1280
+#define CL_PROFILING_COMMAND_SUBMIT                 0x1281
+#define CL_PROFILING_COMMAND_START                  0x1282
+#define CL_PROFILING_COMMAND_END                    0x1283
+
+/********************************************************************************************************/
+
+/********************************************************************************************************/
+
+//  Function signature typedef's
+
+// Platform API
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLGETPLATFORMIDS)(cl_uint          /* num_entries */,
+                 cl_platform_id * /* platforms */,
+                 cl_uint *        /* num_platforms */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL * 
+PFNCLGETPLATFORMINFO)(cl_platform_id   /* platform */, 
+                  cl_platform_info /* param_name */,
+                  size_t           /* param_value_size */, 
+                  void *           /* param_value */,
+                  size_t *         /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+// Device APIs
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLGETDEVICEIDS)(cl_platform_id   /* platform */,
+               cl_device_type   /* device_type */, 
+               cl_uint          /* num_entries */, 
+               cl_device_id *   /* devices */, 
+               cl_uint *        /* num_devices */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLGETDEVICEINFO)(cl_device_id    /* device */,
+                cl_device_info  /* param_name */, 
+                size_t          /* param_value_size */, 
+                void *          /* param_value */,
+                size_t *        /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+// Context APIs  
+typedef CL_API_ENTRY cl_context (CL_API_CALL *
+PFNCLCREATECONTEXT)(const cl_context_properties * /* properties */,
+                cl_uint                       /* num_devices */,
+                const cl_device_id *          /* devices */,
+                void (*pfn_notify)(const char *, const void *, size_t, void *) /* pfn_notify */,
+                void *                        /* user_data */,
+                cl_int *                      /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_context (CL_API_CALL *
+PFNCLCREATECONTEXTFROMTYPE)(const cl_context_properties * /* properties */,
+                        cl_device_type                /* device_type */,
+                        void (*pfn_notify)(const char *, const void *, size_t, void *) /* pfn_notify */,
+                        void *                        /* user_data */,
+                        cl_int *                      /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLRETAINCONTEXT)(cl_context /* context */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLRELEASECONTEXT)(cl_context /* context */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLGETCONTEXTINFO)(cl_context         /* context */, 
+                 cl_context_info    /* param_name */, 
+                 size_t             /* param_value_size */, 
+                 void *             /* param_value */, 
+                 size_t *           /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+// Command Queue APIs
+typedef CL_API_ENTRY cl_command_queue (CL_API_CALL *
+PFNCLCREATECOMMANDQUEUE)(cl_context                     /* context */, 
+                     cl_device_id                   /* device */, 
+                     cl_command_queue_properties    /* properties */,
+                     cl_int *                       /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLRETAINCOMMANDQUEUE)(cl_command_queue /* command_queue */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLRELEASECOMMANDQUEUE)(cl_command_queue /* command_queue */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLGETCOMMANDQUEUEINFO)(cl_command_queue      /* command_queue */,
+                      cl_command_queue_info /* param_name */,
+                      size_t                /* param_value_size */,
+                      void *                /* param_value */,
+                      size_t *              /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLSETCOMMANDQUEUEPROPERTY)(cl_command_queue              /* command_queue */,
+                          cl_command_queue_properties   /* properties */, 
+                          cl_bool                        /* enable */,
+                          cl_command_queue_properties * /* old_properties */) CL_API_SUFFIX__VERSION_1_0;
+
+// Memory Object APIs
+typedef CL_API_ENTRY cl_mem (CL_API_CALL *
+PFNCLCREATEBUFFER)(cl_context   /* context */,
+               cl_mem_flags /* flags */,
+               size_t       /* size */,
+               void *       /* host_ptr */,
+               cl_int *     /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_mem (CL_API_CALL *
+PFNCLCREATEIMAGE2D)(cl_context              /* context */,
+                cl_mem_flags            /* flags */,
+                const cl_image_format * /* image_format */,
+                size_t                  /* image_width */,
+                size_t                  /* image_height */,
+                size_t                  /* image_row_pitch */, 
+                void *                  /* host_ptr */,
+                cl_int *                /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_mem (CL_API_CALL *
+PFNCLCREATEIMAGE3D)(cl_context              /* context */,
+                cl_mem_flags            /* flags */,
+                const cl_image_format * /* image_format */,
+                size_t                  /* image_width */, 
+                size_t                  /* image_height */,
+                size_t                  /* image_depth */, 
+                size_t                  /* image_row_pitch */, 
+                size_t                  /* image_slice_pitch */, 
+                void *                  /* host_ptr */,
+                cl_int *                /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLRETAINMEMOBJECT)(cl_mem /* memobj */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLRELEASEMEMOBJECT)(cl_mem /* memobj */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLGETSUPPORTEDIMAGEFORMATS)(cl_context           /* context */,
+                           cl_mem_flags         /* flags */,
+                           cl_mem_object_type   /* image_type */,
+                           cl_uint              /* num_entries */,
+                           cl_image_format *    /* image_formats */,
+                           cl_uint *            /* num_image_formats */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLGETMEMOBJECTINFO)(cl_mem           /* memobj */,
+                   cl_mem_info      /* param_name */, 
+                   size_t           /* param_value_size */,
+                   void *           /* param_value */,
+                   size_t *         /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLGETIMAGEINFO)(cl_mem           /* image */,
+               cl_image_info    /* param_name */, 
+               size_t           /* param_value_size */,
+               void *           /* param_value */,
+               size_t *         /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+// Sampler APIs
+typedef CL_API_ENTRY cl_sampler (CL_API_CALL *
+PFNCLCREATESAMPLER)(cl_context          /* context */,
+                cl_bool             /* normalized_coords */, 
+                cl_addressing_mode  /* addressing_mode */, 
+                cl_filter_mode      /* filter_mode */,
+                cl_int *            /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLRETAINSAMPLER)(cl_sampler /* sampler */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLRELEASESAMPLER)(cl_sampler /* sampler */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLGETSAMPLERINFO)(cl_sampler         /* sampler */,
+                 cl_sampler_info    /* param_name */,
+                 size_t             /* param_value_size */,
+                 void *             /* param_value */,
+                 size_t *           /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+// Program Object APIs
+typedef CL_API_ENTRY cl_program (CL_API_CALL *
+PFNCLCREATEPROGRAMWITHSOURCE)(cl_context        /* context */,
+                          cl_uint           /* count */,
+                          const char **     /* strings */,
+                          const size_t *    /* lengths */,
+                          cl_int *          /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_program (CL_API_CALL *
+PFNCLCREATEPROGRAMWITHBINARY)(cl_context                     /* context */,
+                          cl_uint                        /* num_devices */,
+                          const cl_device_id *           /* device_list */,
+                          const size_t *                 /* lengths */,
+                          const unsigned char **         /* binaries */,
+                          cl_int *                       /* binary_status */,
+                          cl_int *                       /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLRETAINPROGRAM)(cl_program /* program */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLRELEASEPROGRAM)(cl_program /* program */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLBUILDPROGRAM)(cl_program           /* program */,
+               cl_uint              /* num_devices */,
+               const cl_device_id * /* device_list */,
+               const char *         /* options */, 
+               void (*pfn_notify)(cl_program /* program */, void * /* user_data */),
+               void *               /* user_data */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLUNLOADCOMPILER)(void) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLGETPROGRAMINFO)(cl_program         /* program */,
+                 cl_program_info    /* param_name */,
+                 size_t             /* param_value_size */,
+                 void *             /* param_value */,
+                 size_t *           /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLGETPROGRAMBUILDINFO)(cl_program            /* program */,
+                      cl_device_id          /* device */,
+                      cl_program_build_info /* param_name */,
+                      size_t                /* param_value_size */,
+                      void *                /* param_value */,
+                      size_t *              /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+// Kernel Object APIs
+typedef CL_API_ENTRY cl_kernel (CL_API_CALL *
+PFNCLCREATEKERNEL)(cl_program      /* program */,
+               const char *    /* kernel_name */,
+               cl_int *        /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLCREATEKERNELSINPROGRAM)(cl_program     /* program */,
+                         cl_uint        /* num_kernels */,
+                         cl_kernel *    /* kernels */,
+                         cl_uint *      /* num_kernels_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLRETAINKERNEL)(cl_kernel    /* kernel */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLRELEASEKERNEL)(cl_kernel   /* kernel */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLSETKERNELARG)(cl_kernel    /* kernel */,
+               cl_uint      /* arg_index */,
+               size_t       /* arg_size */,
+               const void * /* arg_value */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLGETKERNELINFO)(cl_kernel       /* kernel */,
+                cl_kernel_info  /* param_name */,
+                size_t          /* param_value_size */,
+                void *          /* param_value */,
+                size_t *        /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLGETKERNELWORKGROUPINFO)(cl_kernel                  /* kernel */,
+                         cl_device_id               /* device */,
+                         cl_kernel_work_group_info  /* param_name */,
+                         size_t                     /* param_value_size */,
+                         void *                     /* param_value */,
+                         size_t *                   /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+// Event Object APIs
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLWAITFOREVENTS)(cl_uint             /* num_events */,
+                const cl_event *    /* event_list */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLGETEVENTINFO)(cl_event         /* event */,
+               cl_event_info    /* param_name */,
+               size_t           /* param_value_size */,
+               void *           /* param_value */,
+               size_t *         /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLRETAINEVENT)(cl_event /* event */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLRELEASEEVENT)(cl_event /* event */) CL_API_SUFFIX__VERSION_1_0;
+
+// Profiling APIs
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLGETEVENTPROFILINGINFO)(cl_event            /* event */,
+                        cl_profiling_info   /* param_name */,
+                        size_t              /* param_value_size */,
+                        void *              /* param_value */,
+                        size_t *            /* param_value_size_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+// Flush and Finish APIs
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLFLUSH)(cl_command_queue /* command_queue */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLFINISH)(cl_command_queue /* command_queue */) CL_API_SUFFIX__VERSION_1_0;
+
+// Enqueued Commands APIs
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLENQUEUEREADBUFFER)(cl_command_queue    /* command_queue */,
+                    cl_mem              /* buffer */,
+                    cl_bool             /* blocking_read */,
+                    size_t              /* offset */,
+                    size_t              /* cb */, 
+                    void *              /* ptr */,
+                    cl_uint             /* num_events_in_wait_list */,
+                    const cl_event *    /* event_wait_list */,
+                    cl_event *          /* event */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLENQUEUEWRITEBUFFER)(cl_command_queue   /* command_queue */, 
+                     cl_mem             /* buffer */, 
+                     cl_bool            /* blocking_write */, 
+                     size_t             /* offset */, 
+                     size_t             /* cb */, 
+                     const void *       /* ptr */, 
+                     cl_uint            /* num_events_in_wait_list */, 
+                     const cl_event *   /* event_wait_list */, 
+                     cl_event *         /* event */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLENQUEUECOPYBUFFER)(cl_command_queue    /* command_queue */, 
+                    cl_mem              /* src_buffer */,
+                    cl_mem              /* dst_buffer */, 
+                    size_t              /* src_offset */,
+                    size_t              /* dst_offset */,
+                    size_t              /* cb */, 
+                    cl_uint             /* num_events_in_wait_list */,
+                    const cl_event *    /* event_wait_list */,
+                    cl_event *          /* event */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLENQUEUEREADIMAGE)(cl_command_queue     /* command_queue */,
+                   cl_mem               /* image */,
+                   cl_bool              /* blocking_read */, 
+                   const size_t *       /* origin[3] */,
+                   const size_t *       /* region[3] */,
+                   size_t               /* row_pitch */,
+                   size_t               /* slice_pitch */, 
+                   void *               /* ptr */,
+                   cl_uint              /* num_events_in_wait_list */,
+                   const cl_event *     /* event_wait_list */,
+                   cl_event *           /* event */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLENQUEUEWRITEIMAGE)(cl_command_queue    /* command_queue */,
+                    cl_mem              /* image */,
+                    cl_bool             /* blocking_write */, 
+                    const size_t *      /* origin[3] */,
+                    const size_t *      /* region[3] */,
+                    size_t              /* input_row_pitch */,
+                    size_t              /* input_slice_pitch */, 
+                    const void *        /* ptr */,
+                    cl_uint             /* num_events_in_wait_list */,
+                    const cl_event *    /* event_wait_list */,
+                    cl_event *          /* event */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLENQUEUECOPYIMAGE)(cl_command_queue     /* command_queue */,
+                   cl_mem               /* src_image */,
+                   cl_mem               /* dst_image */, 
+                   const size_t *       /* src_origin[3] */,
+                   const size_t *       /* dst_origin[3] */,
+                   const size_t *       /* region[3] */, 
+                   cl_uint              /* num_events_in_wait_list */,
+                   const cl_event *     /* event_wait_list */,
+                   cl_event *           /* event */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLENQUEUECOPYIMAGETOBUFFER)(cl_command_queue /* command_queue */,
+                           cl_mem           /* src_image */,
+                           cl_mem           /* dst_buffer */, 
+                           const size_t *   /* src_origin[3] */,
+                           const size_t *   /* region[3] */, 
+                           size_t           /* dst_offset */,
+                           cl_uint          /* num_events_in_wait_list */,
+                           const cl_event * /* event_wait_list */,
+                           cl_event *       /* event */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLENQUEUECOPYBUFFERTOIMAGE)(cl_command_queue /* command_queue */,
+                           cl_mem           /* src_buffer */,
+                           cl_mem           /* dst_image */, 
+                           size_t           /* src_offset */,
+                           const size_t *   /* dst_origin[3] */,
+                           const size_t *   /* region[3] */, 
+                           cl_uint          /* num_events_in_wait_list */,
+                           const cl_event * /* event_wait_list */,
+                           cl_event *       /* event */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY void * (CL_API_CALL *
+PFNCLENQUEUEMAPBUFFER)(cl_command_queue /* command_queue */,
+                   cl_mem           /* buffer */,
+                   cl_bool          /* blocking_map */, 
+                   cl_map_flags     /* map_flags */,
+                   size_t           /* offset */,
+                   size_t           /* cb */,
+                   cl_uint          /* num_events_in_wait_list */,
+                   const cl_event * /* event_wait_list */,
+                   cl_event *       /* event */,
+                   cl_int *         /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY void * (CL_API_CALL *
+PFNCLENQUEUEMAPIMAGE)(cl_command_queue  /* command_queue */,
+                  cl_mem            /* image */, 
+                  cl_bool           /* blocking_map */, 
+                  cl_map_flags      /* map_flags */, 
+                  const size_t *    /* origin[3] */,
+                  const size_t *    /* region[3] */,
+                  size_t *          /* image_row_pitch */,
+                  size_t *          /* image_slice_pitch */,
+                  cl_uint           /* num_events_in_wait_list */,
+                  const cl_event *  /* event_wait_list */,
+                  cl_event *        /* event */,
+                  cl_int *          /* errcode_ret */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLENQUEUEUNMAPMEMOBJECT)(cl_command_queue /* command_queue */,
+                        cl_mem           /* memobj */,
+                        void *           /* mapped_ptr */,
+                        cl_uint          /* num_events_in_wait_list */,
+                        const cl_event *  /* event_wait_list */,
+                        cl_event *        /* event */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLENQUEUENDRANGEKERNEL)(cl_command_queue /* command_queue */,
+                       cl_kernel        /* kernel */,
+                       cl_uint          /* work_dim */,
+                       const size_t *   /* global_work_offset */,
+                       const size_t *   /* global_work_size */,
+                       const size_t *   /* local_work_size */,
+                       cl_uint          /* num_events_in_wait_list */,
+                       const cl_event * /* event_wait_list */,
+                       cl_event *       /* event */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLENQUEUETASK)(cl_command_queue  /* command_queue */,
+              cl_kernel         /* kernel */,
+              cl_uint           /* num_events_in_wait_list */,
+              const cl_event *  /* event_wait_list */,
+              cl_event *        /* event */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLENQUEUENATIVEKERNEL)(cl_command_queue  /* command_queue */,
+                      void (*user_func)(void *), 
+                      void *            /* args */,
+                      size_t            /* cb_args */, 
+                      cl_uint           /* num_mem_objects */,
+                      const cl_mem *    /* mem_list */,
+                      const void **     /* args_mem_loc */,
+                      cl_uint           /* num_events_in_wait_list */,
+                      const cl_event *  /* event_wait_list */,
+                      cl_event *        /* event */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLENQUEUEMARKER)(cl_command_queue    /* command_queue */,
+                cl_event *          /* event */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLENQUEUEWAITFOREVENTS)(cl_command_queue /* command_queue */,
+                       cl_uint          /* num_events */,
+                       const cl_event * /* event_list */) CL_API_SUFFIX__VERSION_1_0;
+
+typedef CL_API_ENTRY cl_int (CL_API_CALL *
+PFNCLENQUEUEBARRIER)(cl_command_queue /* command_queue */) CL_API_SUFFIX__VERSION_1_0;
+
+// Extension function access
+//
+// Returns the extension function address for the given function name,
+// or NULL if a valid function can not be found.  The client must
+// check to make sure the address is not NULL, before using or 
+// calling the returned function address.
+//
+typedef CL_API_ENTRY void * (CL_API_CALL * PFNCLGETEXTENSIONFUNCTIONADDRESS)(const char * /* func_name */) CL_API_SUFFIX__VERSION_1_0;
+
+
+#define CLEW_STATIC
+
+#ifdef CLEW_STATIC
+#  define CLEWAPI extern
+#else
+#  ifdef CLEW_BUILD
+#    define CLEWAPI extern __declspec(dllexport)
+#  else
+#    define CLEWAPI extern __declspec(dllimport)
+#  endif
+#endif
+
+#if defined(_WIN32)
+#define CLEW_FUN_EXPORT extern
+#else
+#define CLEW_FUN_EXPORT CLEWAPI
+#endif
+
+#define CLEW_GET_FUN(x) x
+
+
+//  Variables holding function entry points
+CLEW_FUN_EXPORT     PFNCLGETPLATFORMIDS                 __oclGetPlatformIDs                ;
+CLEW_FUN_EXPORT     PFNCLGETPLATFORMINFO                __oclGetPlatformInfo               ;
+CLEW_FUN_EXPORT     PFNCLGETDEVICEIDS                   __oclGetDeviceIDs                  ;
+CLEW_FUN_EXPORT     PFNCLGETDEVICEINFO                  __oclGetDeviceInfo                 ;
+CLEW_FUN_EXPORT     PFNCLCREATECONTEXT                  __oclCreateContext                 ;
+CLEW_FUN_EXPORT     PFNCLCREATECONTEXTFROMTYPE          __oclCreateContextFromType         ;
+CLEW_FUN_EXPORT     PFNCLRETAINCONTEXT                  __oclRetainContext                 ;
+CLEW_FUN_EXPORT     PFNCLRELEASECONTEXT                 __oclReleaseContext                ;
+CLEW_FUN_EXPORT     PFNCLGETCONTEXTINFO                 __oclGetContextInfo                ;
+CLEW_FUN_EXPORT     PFNCLCREATECOMMANDQUEUE             __oclCreateCommandQueue            ;
+CLEW_FUN_EXPORT     PFNCLRETAINCOMMANDQUEUE             __oclRetainCommandQueue            ;
+CLEW_FUN_EXPORT     PFNCLRELEASECOMMANDQUEUE            __oclReleaseCommandQueue           ;
+CLEW_FUN_EXPORT     PFNCLGETCOMMANDQUEUEINFO            __oclGetCommandQueueInfo           ;
+CLEW_FUN_EXPORT     PFNCLSETCOMMANDQUEUEPROPERTY        __oclSetCommandQueueProperty       ;
+CLEW_FUN_EXPORT     PFNCLCREATEBUFFER                   __oclCreateBuffer                  ;
+CLEW_FUN_EXPORT     PFNCLCREATEIMAGE2D                  __oclCreateImage2D                 ;
+CLEW_FUN_EXPORT     PFNCLCREATEIMAGE3D                  __oclCreateImage3D                 ;
+CLEW_FUN_EXPORT     PFNCLRETAINMEMOBJECT                __oclRetainMemObject               ;
+CLEW_FUN_EXPORT     PFNCLRELEASEMEMOBJECT               __oclReleaseMemObject              ;
+CLEW_FUN_EXPORT     PFNCLGETSUPPORTEDIMAGEFORMATS       __oclGetSupportedImageFormats      ;
+CLEW_FUN_EXPORT     PFNCLGETMEMOBJECTINFO               __oclGetMemObjectInfo              ;
+CLEW_FUN_EXPORT     PFNCLGETIMAGEINFO                   __oclGetImageInfo                  ;
+CLEW_FUN_EXPORT     PFNCLCREATESAMPLER                  __oclCreateSampler                 ;
+CLEW_FUN_EXPORT     PFNCLRETAINSAMPLER                  __oclRetainSampler                 ;
+CLEW_FUN_EXPORT     PFNCLRELEASESAMPLER                 __oclReleaseSampler                ;
+CLEW_FUN_EXPORT     PFNCLGETSAMPLERINFO                 __oclGetSamplerInfo                ;
+CLEW_FUN_EXPORT     PFNCLCREATEPROGRAMWITHSOURCE        __oclCreateProgramWithSource       ;
+CLEW_FUN_EXPORT     PFNCLCREATEPROGRAMWITHBINARY        __oclCreateProgramWithBinary       ;
+CLEW_FUN_EXPORT     PFNCLRETAINPROGRAM                  __oclRetainProgram                 ;
+CLEW_FUN_EXPORT     PFNCLRELEASEPROGRAM                 __oclReleaseProgram                ;
+CLEW_FUN_EXPORT     PFNCLBUILDPROGRAM                   __oclBuildProgram                  ;
+CLEW_FUN_EXPORT     PFNCLUNLOADCOMPILER                 __oclUnloadCompiler                ;
+CLEW_FUN_EXPORT     PFNCLGETPROGRAMINFO                 __oclGetProgramInfo                ;
+CLEW_FUN_EXPORT     PFNCLGETPROGRAMBUILDINFO            __oclGetProgramBuildInfo           ;
+CLEW_FUN_EXPORT     PFNCLCREATEKERNEL                   __oclCreateKernel                  ;
+CLEW_FUN_EXPORT     PFNCLCREATEKERNELSINPROGRAM         __oclCreateKernelsInProgram        ;
+CLEW_FUN_EXPORT     PFNCLRETAINKERNEL                   __oclRetainKernel                  ;
+CLEW_FUN_EXPORT     PFNCLRELEASEKERNEL                  __oclReleaseKernel                 ;
+CLEW_FUN_EXPORT     PFNCLSETKERNELARG                   __oclSetKernelArg                  ;
+CLEW_FUN_EXPORT     PFNCLGETKERNELINFO                  __oclGetKernelInfo                 ;
+CLEW_FUN_EXPORT     PFNCLGETKERNELWORKGROUPINFO         __oclGetKernelWorkGroupInfo        ;
+CLEW_FUN_EXPORT     PFNCLWAITFOREVENTS                  __oclWaitForEvents                 ;
+CLEW_FUN_EXPORT     PFNCLGETEVENTINFO                   __oclGetEventInfo                  ;
+CLEW_FUN_EXPORT     PFNCLRETAINEVENT                    __oclRetainEvent                   ;
+CLEW_FUN_EXPORT     PFNCLRELEASEEVENT                   __oclReleaseEvent                  ;
+CLEW_FUN_EXPORT     PFNCLGETEVENTPROFILINGINFO          __oclGetEventProfilingInfo         ;
+CLEW_FUN_EXPORT     PFNCLFLUSH                          __oclFlush                         ;
+CLEW_FUN_EXPORT     PFNCLFINISH                         __oclFinish                        ;
+CLEW_FUN_EXPORT     PFNCLENQUEUEREADBUFFER              __oclEnqueueReadBuffer             ;
+CLEW_FUN_EXPORT     PFNCLENQUEUEWRITEBUFFER             __oclEnqueueWriteBuffer            ;
+CLEW_FUN_EXPORT     PFNCLENQUEUECOPYBUFFER              __oclEnqueueCopyBuffer             ;
+CLEW_FUN_EXPORT     PFNCLENQUEUEREADIMAGE               __oclEnqueueReadImage              ;
+CLEW_FUN_EXPORT     PFNCLENQUEUEWRITEIMAGE              __oclEnqueueWriteImage             ;
+CLEW_FUN_EXPORT     PFNCLENQUEUECOPYIMAGE               __oclEnqueueCopyImage              ;
+CLEW_FUN_EXPORT     PFNCLENQUEUECOPYIMAGETOBUFFER       __oclEnqueueCopyImageToBuffer      ;
+CLEW_FUN_EXPORT     PFNCLENQUEUECOPYBUFFERTOIMAGE       __oclEnqueueCopyBufferToImage      ;
+CLEW_FUN_EXPORT     PFNCLENQUEUEMAPBUFFER               __oclEnqueueMapBuffer              ;
+CLEW_FUN_EXPORT     PFNCLENQUEUEMAPIMAGE                __oclEnqueueMapImage               ;
+CLEW_FUN_EXPORT     PFNCLENQUEUEUNMAPMEMOBJECT          __oclEnqueueUnmapMemObject         ;
+CLEW_FUN_EXPORT     PFNCLENQUEUENDRANGEKERNEL           __oclEnqueueNDRangeKernel          ;
+CLEW_FUN_EXPORT     PFNCLENQUEUETASK                    __oclEnqueueTask                   ;
+CLEW_FUN_EXPORT     PFNCLENQUEUENATIVEKERNEL            __oclEnqueueNativeKernel           ;
+CLEW_FUN_EXPORT     PFNCLENQUEUEMARKER                  __oclEnqueueMarker                 ;
+CLEW_FUN_EXPORT     PFNCLENQUEUEWAITFOREVENTS           __oclEnqueueWaitForEvents          ;
+CLEW_FUN_EXPORT     PFNCLENQUEUEBARRIER                 __oclEnqueueBarrier                ;
+CLEW_FUN_EXPORT     PFNCLGETEXTENSIONFUNCTIONADDRESS    __oclGetExtensionFunctionAddress   ;
+
+
+#define	clGetPlatformIDs                CLEW_GET_FUN(__oclGetPlatformIDs                )
+#define	clGetPlatformInfo               CLEW_GET_FUN(__oclGetPlatformInfo               )
+#define	clGetDeviceIDs                  CLEW_GET_FUN(__oclGetDeviceIDs                  )
+#define	clGetDeviceInfo                 CLEW_GET_FUN(__oclGetDeviceInfo                 )
+#define	clCreateContext                 CLEW_GET_FUN(__oclCreateContext                 )
+#define	clCreateContextFromType         CLEW_GET_FUN(__oclCreateContextFromType         )
+#define	clRetainContext                 CLEW_GET_FUN(__oclRetainContext                 )
+#define	clReleaseContext                CLEW_GET_FUN(__oclReleaseContext                )
+#define	clGetContextInfo                CLEW_GET_FUN(__oclGetContextInfo                )
+#define	clCreateCommandQueue            CLEW_GET_FUN(__oclCreateCommandQueue            )
+#define	clRetainCommandQueue            CLEW_GET_FUN(__oclRetainCommandQueue            )
+#define	clReleaseCommandQueue           CLEW_GET_FUN(__oclReleaseCommandQueue           )
+#define	clGetCommandQueueInfo           CLEW_GET_FUN(__oclGetCommandQueueInfo           )
+#define	clSetCommandQueueProperty       CLEW_GET_FUN(__oclSetCommandQueueProperty       )
+#define	clCreateBuffer                  CLEW_GET_FUN(__oclCreateBuffer                  )
+#define	clCreateImage2D                 CLEW_GET_FUN(__oclCreateImage2D                 )
+#define	clCreateImage3D                 CLEW_GET_FUN(__oclCreateImage3D                 )
+#define	clRetainMemObject               CLEW_GET_FUN(__oclRetainMemObject               )
+#define	clReleaseMemObject              CLEW_GET_FUN(__oclReleaseMemObject              )
+#define	clGetSupportedImageFormats      CLEW_GET_FUN(__oclGetSupportedImageFormats      )
+#define	clGetMemObjectInfo              CLEW_GET_FUN(__oclGetMemObjectInfo              )
+#define	clGetImageInfo                  CLEW_GET_FUN(__oclGetImageInfo                  )
+#define	clCreateSampler                 CLEW_GET_FUN(__oclCreateSampler                 )
+#define	clRetainSampler                 CLEW_GET_FUN(__oclRetainSampler                 )
+#define	clReleaseSampler                CLEW_GET_FUN(__oclReleaseSampler                )
+#define	clGetSamplerInfo                CLEW_GET_FUN(__oclGetSamplerInfo                )
+#define	clCreateProgramWithSource       CLEW_GET_FUN(__oclCreateProgramWithSource       )
+#define	clCreateProgramWithBinary       CLEW_GET_FUN(__oclCreateProgramWithBinary       )
+#define	clRetainProgram                 CLEW_GET_FUN(__oclRetainProgram                 )
+#define	clReleaseProgram                CLEW_GET_FUN(__oclReleaseProgram                )
+#define	clBuildProgram                  CLEW_GET_FUN(__oclBuildProgram                  )
+#define	clUnloadCompiler                CLEW_GET_FUN(__oclUnloadCompiler                )
+#define	clGetProgramInfo                CLEW_GET_FUN(__oclGetProgramInfo                )
+#define	clGetProgramBuildInfo           CLEW_GET_FUN(__oclGetProgramBuildInfo           )
+#define	clCreateKernel                  CLEW_GET_FUN(__oclCreateKernel                  )
+#define	clCreateKernelsInProgram        CLEW_GET_FUN(__oclCreateKernelsInProgram        )
+#define	clRetainKernel                  CLEW_GET_FUN(__oclRetainKernel                  )
+#define	clReleaseKernel                 CLEW_GET_FUN(__oclReleaseKernel                 )
+#define	clSetKernelArg                  CLEW_GET_FUN(__oclSetKernelArg                  )
+#define	clGetKernelInfo                 CLEW_GET_FUN(__oclGetKernelInfo                 )
+#define	clGetKernelWorkGroupInfo        CLEW_GET_FUN(__oclGetKernelWorkGroupInfo        )
+#define	clWaitForEvents                 CLEW_GET_FUN(__oclWaitForEvents                 )
+#define	clGetEventInfo                  CLEW_GET_FUN(__oclGetEventInfo                  )
+#define	clRetainEvent                   CLEW_GET_FUN(__oclRetainEvent                   )
+#define	clReleaseEvent                  CLEW_GET_FUN(__oclReleaseEvent                  )
+#define	clGetEventProfilingInfo         CLEW_GET_FUN(__oclGetEventProfilingInfo         )
+#define	clFlush                         CLEW_GET_FUN(__oclFlush                         )
+#define	clFinish                        CLEW_GET_FUN(__oclFinish                        )
+#define	clEnqueueReadBuffer             CLEW_GET_FUN(__oclEnqueueReadBuffer             )
+#define	clEnqueueWriteBuffer            CLEW_GET_FUN(__oclEnqueueWriteBuffer            )
+#define	clEnqueueCopyBuffer             CLEW_GET_FUN(__oclEnqueueCopyBuffer             )
+#define	clEnqueueReadImage              CLEW_GET_FUN(__oclEnqueueReadImage              )
+#define	clEnqueueWriteImage             CLEW_GET_FUN(__oclEnqueueWriteImage             )
+#define	clEnqueueCopyImage              CLEW_GET_FUN(__oclEnqueueCopyImage              )
+#define	clEnqueueCopyImageToBuffer      CLEW_GET_FUN(__oclEnqueueCopyImageToBuffer      )
+#define	clEnqueueCopyBufferToImage      CLEW_GET_FUN(__oclEnqueueCopyBufferToImage      )
+#define	clEnqueueMapBuffer              CLEW_GET_FUN(__oclEnqueueMapBuffer              )
+#define	clEnqueueMapImage               CLEW_GET_FUN(__oclEnqueueMapImage               )
+#define	clEnqueueUnmapMemObject         CLEW_GET_FUN(__oclEnqueueUnmapMemObject         )
+#define	clEnqueueNDRangeKernel          CLEW_GET_FUN(__oclEnqueueNDRangeKernel          )
+#define	clEnqueueTask                   CLEW_GET_FUN(__oclEnqueueTask                   )
+#define	clEnqueueNativeKernel           CLEW_GET_FUN(__oclEnqueueNativeKernel           )
+#define	clEnqueueMarker                 CLEW_GET_FUN(__oclEnqueueMarker                 )
+#define	clEnqueueWaitForEvents          CLEW_GET_FUN(__oclEnqueueWaitForEvents          )
+#define	clEnqueueBarrier                CLEW_GET_FUN(__oclEnqueueBarrier                )
+#define	clGetExtensionFunctionAddress   CLEW_GET_FUN(__oclGetExtensionFunctionAddress   )
+
+#endif  //  CLCC_GENERATE_DOCUMENTATION
+
+#define CLEW_SUCCESS                0       //!<    Success error code
+#define CLEW_ERROR_OPEN_FAILED      -1      //!<    Error code for failing to open the dynamic library
+#define CLEW_ERROR_ATEXIT_FAILED    -2      //!<    Error code for failing to queue the closing of the dynamic library to atexit()
+
+//! \brief Load OpenCL dynamic library and set function entry points
+int         clewInit        (const char*);
+//! \brief Convert an OpenCL error code to its string equivalent
+const char* clewErrorString (cl_int error);
+
+#ifdef __cplusplus
+}
+#endif
+
+#endif  //  CLCC_CLEW_HPP_INCLUDED
diff --git a/intern/rigidbody/CMakeLists.txt b/intern/rigidbody/CMakeLists.txt
new file mode 100644
index 00000000000..a323e55e570
--- /dev/null
+++ b/intern/rigidbody/CMakeLists.txt
@@ -0,0 +1,35 @@
+# ***** BEGIN GPL LICENSE BLOCK *****
+#
+# This program is free software; you can redistribute it and/or
+# modify it under the terms of the GNU General Public License
+# as published by the Free Software Foundation; either version 2
+# of the License, or (at your option) any later version.
+#
+# This program is distributed in the hope that it will be useful,
+# but WITHOUT ANY WARRANTY; without even the implied warranty of
+# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+# GNU General Public License for more details.
+#
+# You should have received a copy of the GNU General Public License
+# along with this program; if not, write to the Free Software Foundation,
+# Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
+#
+# The Original Code is Copyright (C) 2006, Blender Foundation
+# All rights reserved.
+#
+# The Original Code is: all of this file.
+#
+# ***** END GPL LICENSE BLOCK *****
+
+SET(INC
+	.
+	../../extern/bullet2/src
+)
+
+set(SRC
+	rb_bullet_api.cpp
+	
+	RBI_api.h
+)
+
+blender_add_lib(bf_intern_rigidbody "${SRC}" "${INC}" "${INC_SYS}")
diff --git a/intern/rigidbody/RBI_api.h b/intern/rigidbody/RBI_api.h
new file mode 100644
index 00000000000..9dde2cc3049
--- /dev/null
+++ b/intern/rigidbody/RBI_api.h
@@ -0,0 +1,309 @@
+/*
+ * ***** BEGIN GPL LICENSE BLOCK *****
+ *
+ * This program is free software; you can redistribute it and/or
+ * modify it under the terms of the GNU General Public License
+ * as published by the Free Software Foundation; either version 2
+ * of the License, or (at your option) any later version.
+ *
+ * This program is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+ * GNU General Public License for more details.
+ *
+ * You should have received a copy of the GNU General Public License
+ * along with this program; if not, write to the Free Software Foundation,
+ * Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
+ *
+ * The Original Code is Copyright (C) 2013 Blender Foundation,
+ * All rights reserved.
+ *
+ * The Original Code is: all of this file.
+ *
+ * Contributor(s): Joshua Leung, Sergej Reich
+ *
+ * ***** END GPL LICENSE BLOCK *****
+ */
+
+/** \file RBI_api.h
+ *  \ingroup RigidBody
+ *  \brief Rigid Body API for interfacing with external Physics Engines
+ */
+
+#ifndef __RB_API_H__
+#define __RB_API_H__
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+/* API Notes:
+ * Currently, this API is optimised for Bullet RigidBodies, and doesn't
+ * take into account other Physics Engines. Some tweaking may be necessary
+ * to allow other systems to be used, in particular there may be references
+ * to datatypes that aren't used here...
+ *
+ * -- Joshua Leung (22 June 2010)
+ */
+
+/* ********************************** */
+/* Partial Type Defines - Aliases for the type of data we store */
+
+// ----------
+
+/* Dynamics World */
+typedef struct rbDynamicsWorld rbDynamicsWorld;
+
+/* Rigid Body */
+typedef struct rbRigidBody rbRigidBody;
+
+/* Collision Shape */
+typedef struct rbCollisionShape rbCollisionShape;
+
+/* Mesh Data (for Collision Shapes of Meshes) */
+typedef struct rbMeshData rbMeshData;
+
+/* Constraint */
+typedef struct rbConstraint rbConstraint;
+
+/* ********************************** */
+/* Dynamics World Methods */
+
+/* Setup ---------------------------- */
+
+/* Create a new dynamics world instance */
+// TODO: add args to set the type of constraint solvers, etc.
+extern rbDynamicsWorld *RB_dworld_new(const float gravity[3]);
+
+/* Delete the given dynamics world, and free any extra data it may require */
+extern void RB_dworld_delete(rbDynamicsWorld *world);
+
+/* Settings ------------------------- */
+
+/* Gravity */
+extern void RB_dworld_get_gravity(rbDynamicsWorld *world, float g_out[3]);
+extern void RB_dworld_set_gravity(rbDynamicsWorld *world, const float g_in[3]);
+
+/* Constraint Solver */
+extern void RB_dworld_set_solver_iterations(rbDynamicsWorld *world, int num_solver_iterations);
+/* Split Impulse */
+extern void RB_dworld_set_split_impulse(rbDynamicsWorld *world, int split_impulse);
+
+/* Simulation ----------------------- */
+
+/* Step the simulation by the desired amount (in seconds) with extra controls on substep sizes and maximum substeps */
+extern void RB_dworld_step_simulation(rbDynamicsWorld *world, float timeStep, int maxSubSteps, float timeSubStep);
+
+/* Export -------------------------- */
+
+/* Exports the dynamics world to physics simulator's serialisation format */
+void RB_dworld_export(rbDynamicsWorld *world, const char *filename);
+
+/* ********************************** */
+/* Rigid Body Methods */
+
+/* Setup ---------------------------- */
+
+/* Add RigidBody to dynamics world */
+extern void RB_dworld_add_body(rbDynamicsWorld *world, rbRigidBody *body, int col_groups);
+
+/* Remove RigidBody from dynamics world */
+extern void RB_dworld_remove_body(rbDynamicsWorld *world, rbRigidBody *body);
+
+/* ............ */
+
+/* Create new RigidBody instance */
+extern rbRigidBody *RB_body_new(rbCollisionShape *shape, const float loc[3], const float rot[4]);
+
+/* Delete the given RigidBody instance */
+extern void RB_body_delete(rbRigidBody *body);
+
+/* Settings ------------------------- */
+
+/* 'Type' */
+extern void RB_body_set_type(rbRigidBody *body, int type, float mass);
+
+/* ............ */
+
+/* Collision Shape */
+extern void RB_body_set_collision_shape(rbRigidBody *body, rbCollisionShape *shape);
+
+/* ............ */
+
+/* Mass */
+extern float RB_body_get_mass(rbRigidBody *body);
+extern void RB_body_set_mass(rbRigidBody *body, float value);
+
+/* Friction */
+extern float RB_body_get_friction(rbRigidBody *body);
+extern void RB_body_set_friction(rbRigidBody *body, float value);
+
+/* Restitution */
+extern float RB_body_get_restitution(rbRigidBody *body);
+extern void RB_body_set_restitution(rbRigidBody *body, float value);
+
+/* Damping */
+extern float RB_body_get_linear_damping(rbRigidBody *body);
+extern void RB_body_set_linear_damping(rbRigidBody *body, float value);
+
+extern float RB_body_get_angular_damping(rbRigidBody *body);
+extern void RB_body_set_angular_damping(rbRigidBody *body, float value);
+
+extern void RB_body_set_damping(rbRigidBody *object, float linear, float angular);
+
+/* Sleeping Thresholds */
+extern float RB_body_get_linear_sleep_thresh(rbRigidBody *body);
+extern void RB_body_set_linear_sleep_thresh(rbRigidBody *body, float value);
+
+extern float RB_body_get_angular_sleep_thresh(rbRigidBody *body);
+extern void RB_body_set_angular_sleep_thresh(rbRigidBody *body, float value);
+
+extern void RB_body_set_sleep_thresh(rbRigidBody *body, float linear, float angular);
+
+/* Linear Velocity */
+extern void RB_body_get_linear_velocity(rbRigidBody *body, float v_out[3]);
+extern void RB_body_set_linear_velocity(rbRigidBody *body, const float v_in[3]);
+
+/* Angular Velocity */
+extern void RB_body_get_angular_velocity(rbRigidBody *body, float v_out[3]);
+extern void RB_body_set_angular_velocity(rbRigidBody *body, const float v_in[3]);
+
+/* Linear/Angular Factor, used to lock translation/roation axes */
+extern void RB_body_set_linear_factor(rbRigidBody *object, float x, float y, float z);
+extern void RB_body_set_angular_factor(rbRigidBody *object, float x, float y, float z);
+
+/* Kinematic State */
+extern void RB_body_set_kinematic_state(rbRigidBody *body, int kinematic);
+
+/* RigidBody Interface - Rigid Body Activation States */
+extern int RB_body_get_activation_state(rbRigidBody *body);
+extern void RB_body_set_activation_state(rbRigidBody *body, int use_deactivation);
+extern void RB_body_activate(rbRigidBody *body);
+extern void RB_body_deactivate(rbRigidBody *body);
+
+
+/* Simulation ----------------------- */
+
+/* Get current transform matrix of RigidBody to use in Blender (OpenGL format) */
+extern void RB_body_get_transform_matrix(rbRigidBody *body, float m_out[4][4]);
+
+/* Set RigidBody's location and rotation */
+extern void RB_body_set_loc_rot(rbRigidBody *body, const float loc[3], const float rot[4]);
+/* Set RigidBody's local scaling */
+extern void RB_body_set_scale(rbRigidBody *body, const float scale[3]);
+
+/* ............ */
+
+/* Get RigidBody's position as vector */
+void RB_body_get_position(rbRigidBody *body, float v_out[3]);
+/* Get RigidBody's orientation as quaternion */
+void RB_body_get_orientation(rbRigidBody *body, float v_out[4]);
+
+/* ............ */
+
+extern void RB_body_apply_central_force(rbRigidBody *body, const float v_in[3]);
+
+/* ********************************** */
+/* Collision Shape Methods */
+
+/* Setup (Standard Shapes) ----------- */
+
+extern rbCollisionShape *RB_shape_new_box(float x, float y, float z);
+extern rbCollisionShape *RB_shape_new_sphere(float radius);
+extern rbCollisionShape *RB_shape_new_capsule(float radius, float height);
+extern rbCollisionShape *RB_shape_new_cone(float radius, float height);
+extern rbCollisionShape *RB_shape_new_cylinder(float radius, float height);
+
+/* Setup (Convex Hull) ------------ */
+
+extern rbCollisionShape *RB_shape_new_convex_hull(float *verts, int stride, int count, float margin, bool *can_embed);
+
+/* Setup (Triangle Mesh) ---------- */
+
+/* 1 */
+extern rbMeshData *RB_trimesh_data_new(void);
+extern void RB_trimesh_add_triangle(rbMeshData *mesh, const float v1[3], const float v2[3], const float v3[3]);
+/* 2a - Triangle Meshes */
+extern rbCollisionShape *RB_shape_new_trimesh(rbMeshData *mesh);
+/* 2b - GImpact Meshes */
+extern rbCollisionShape *RB_shape_new_gimpact_mesh(rbMeshData *mesh);
+
+
+/* Cleanup --------------------------- */
+
+extern void RB_shape_delete(rbCollisionShape *shape);
+
+/* Settings --------------------------- */
+
+/* Collision Margin */
+extern float RB_shape_get_margin(rbCollisionShape *shape);
+extern void RB_shape_set_margin(rbCollisionShape *shape, float value);
+
+/* ********************************** */
+/* Constraints */
+
+/* Setup ----------------------------- */
+
+/* Add Rigid Body Constraint to simulation world */
+extern void RB_dworld_add_constraint(rbDynamicsWorld *world, rbConstraint *con, int disable_collisions);
+
+/* Remove Rigid Body Constraint from simulation world */
+extern void RB_dworld_remove_constraint(rbDynamicsWorld *world, rbConstraint *con);
+
+extern rbConstraint *RB_constraint_new_point(float pivot[3], rbRigidBody *rb1, rbRigidBody *rb2);
+extern rbConstraint *RB_constraint_new_fixed(float pivot[3], float orn[4], rbRigidBody *rb1, rbRigidBody *rb2);
+extern rbConstraint *RB_constraint_new_hinge(float pivot[3], float orn[4], rbRigidBody *rb1, rbRigidBody *rb2);
+extern rbConstraint *RB_constraint_new_slider(float pivot[3], float orn[4], rbRigidBody *rb1, rbRigidBody *rb2);
+extern rbConstraint *RB_constraint_new_piston(float pivot[3], float orn[4], rbRigidBody *rb1, rbRigidBody *rb2);
+extern rbConstraint *RB_constraint_new_6dof(float pivot[3], float orn[4], rbRigidBody *rb1, rbRigidBody *rb2);
+extern rbConstraint *RB_constraint_new_6dof_spring(float pivot[3], float orn[4], rbRigidBody *rb1, rbRigidBody *rb2);
+
+/* ............ */
+
+/* Cleanup --------------------------- */
+
+extern void RB_constraint_delete(rbConstraint *con);
+
+/* Settings --------------------------- */
+
+/* Enable or disable constraint */
+extern void RB_constraint_set_enabled(rbConstraint *con, int enabled);
+
+/* Limits */
+#define RB_LIMIT_LIN_X 0
+#define RB_LIMIT_LIN_Y 1
+#define RB_LIMIT_LIN_Z 2
+#define RB_LIMIT_ANG_X 3
+#define RB_LIMIT_ANG_Y 4
+#define RB_LIMIT_ANG_Z 5
+/* Bullet uses the following convention:
+ * - lower limit == upper limit -> axis is locked
+ * - lower limit > upper limit -> axis is free
+ * - lower limit < upper limit -> axis is limited in given range
+ */
+extern void RB_constraint_set_limits_hinge(rbConstraint *con, float lower, float upper);
+extern void RB_constraint_set_limits_slider(rbConstraint *con, float lower, float upper);
+extern void RB_constraint_set_limits_piston(rbConstraint *con, float lin_lower, float lin_upper, float ang_lower, float ang_upper);
+extern void RB_constraint_set_limits_6dof(rbConstraint *con, int axis, float lower, float upper);
+
+/* 6dof spring specific */
+extern void RB_constraint_set_stiffness_6dof_spring(rbConstraint *con, int axis, float stiffness);
+extern void RB_constraint_set_damping_6dof_spring(rbConstraint *con, int axis, float damping);
+extern void RB_constraint_set_spring_6dof_spring(rbConstraint *con, int axis, int enable);
+extern void RB_constraint_set_equilibrium_6dof_spring(rbConstraint *con);
+
+/* Set number of constraint solver iterations made per step, this overrided world setting
+ * To use default set it to -1 */
+extern void RB_constraint_set_solver_iterations(rbConstraint *con, int num_solver_iterations);
+
+/* Set breaking impulse threshold, if constraint shouldn't break it can be set to FLT_MAX */
+extern void RB_constraint_set_breaking_threshold(rbConstraint *con, float threshold);
+
+/* ********************************** */
+
+#ifdef __cplusplus
+}
+#endif
+
+#endif /* __RB_API_H__ */
+ 
diff --git a/intern/rigidbody/SConscript b/intern/rigidbody/SConscript
new file mode 100644
index 00000000000..977281f8eef
--- /dev/null
+++ b/intern/rigidbody/SConscript
@@ -0,0 +1,42 @@
+#!/usr/bin/python
+# $Id: SConscript $
+# ***** BEGIN GPL LICENSE BLOCK *****
+#
+# This program is free software; you can redistribute it and/or
+# modify it under the terms of the GNU General Public License
+# as published by the Free Software Foundation; either version 2
+# of the License, or (at your option) any later version.
+#
+# This program is distributed in the hope that it will be useful,
+# but WITHOUT ANY WARRANTY; without even the implied warranty of
+# MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+# GNU General Public License for more details.
+#
+# You should have received a copy of the GNU General Public License
+# along with this program; if not, write to the Free Software Foundation,
+# Inc., 59 Temple Place - Suite 330, Boston, MA  02111-1307, USA.
+#
+# The Original Code is Copyright (C) 2010, Blender Foundation
+# All rights reserved.
+#
+# The Original Code is: all of this file.
+#
+# Contributor(s): Joshua Leung
+#
+# ***** END GPL LICENSE BLOCK *****
+
+Import('env')
+
+# XXX: we need a contingency plan for when not compiling with Bullet,
+# since this module will always get included... 
+# This problem will also apply to other engines at a later date too...
+sources = env.Glob('*.cpp')
+
+incs = [
+    '.',
+    '../../extern/bullet2/src',
+    ]
+
+env.BlenderLib('bf_intern_rigidbody', sources=sources,
+               includes=incs, defines=[],
+               libtype=['core', 'player'], priority=[180, 30])
diff --git a/intern/rigidbody/rb_bullet_api.cpp b/intern/rigidbody/rb_bullet_api.cpp
new file mode 100644
index 00000000000..58345d4e08c
--- /dev/null
+++ b/intern/rigidbody/rb_bullet_api.cpp
@@ -0,0 +1,950 @@
+/*
+ * ***** BEGIN GPL LICENSE BLOCK *****
+ *
+ * This program is free software; you can redistribute it and/or
+ * modify it under the terms of the GNU General Public License
+ * as published by the Free Software Foundation; either version 2
+ * of the License, or (at your option) any later version.
+ *
+ * This program is distributed in the hope that it will be useful,
+ * but WITHOUT ANY WARRANTY; without even the implied warranty of
+ * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+ * GNU General Public License for more details.
+ *
+ * You should have received a copy of the GNU General Public License
+ * along with this program; if not, write to the Free Software Foundation,
+ * Inc., 51 Franklin Street, Fifth Floor, Boston, MA 02110-1301, USA.
+ *
+ * The Original Code is Copyright (C) 2013 Blender Foundation
+ * All rights reserved.
+ *
+ * The Original Code is: all of this file.
+ *
+ * Contributor(s): Joshua Leung, Sergej Reich
+ *
+ * ***** END GPL LICENSE BLOCK *****
+ */
+
+/** \file rb_bullet_api.cpp
+ *  \ingroup RigidBody
+ *  \brief Rigid Body API implementation for Bullet
+ */
+
+/*
+Bullet Continuous Collision Detection and Physics Library
+Copyright (c) 2003-2006 Erwin Coumans  http://continuousphysics.com/Bullet/
+
+This software is provided 'as-is', without any express or implied warranty.
+In no event will the authors be held liable for any damages arising from the use of this software.
+Permission is granted to anyone to use this software for any purpose, 
+including commercial applications, and to alter it and redistribute it freely, 
+subject to the following restrictions:
+
+1. The origin of this software must not be misrepresented; you must not claim that you wrote the original software. If you use this software in a product, an acknowledgment in the product documentation would be appreciated but is not required.
+2. Altered source versions must be plainly marked as such, and must not be misrepresented as being the original software.
+3. This notice may not be removed or altered from any source distribution.
+*/
+ 
+/* This file defines the "RigidBody interface" for the 
+ * Bullet Physics Engine. This API is designed to be used
+ * from C-code in Blender as part of the Rigid Body simulation
+ * system.
+ *
+ * It is based on the Bullet C-API, but is heavily modified to 
+ * give access to more data types and to offer a nicer interface.
+ *
+ * -- Joshua Leung, June 2010
+ */
+
+#include <stdio.h>
+ 
+#include "RBI_api.h"
+
+#include "btBulletDynamicsCommon.h"
+
+#include "LinearMath/btVector3.h"
+#include "LinearMath/btScalar.h"	
+#include "LinearMath/btMatrix3x3.h"
+#include "LinearMath/btTransform.h"
+#include "LinearMath/btConvexHullComputer.h"
+
+#include "BulletCollision/Gimpact/btGImpactShape.h"
+#include "BulletCollision/Gimpact/btGImpactCollisionAlgorithm.h"
+#include "BulletCollision/CollisionShapes/btScaledBvhTriangleMeshShape.h"
+
+struct rbDynamicsWorld {
+	btDiscreteDynamicsWorld *dynamicsWorld;
+	btDefaultCollisionConfiguration *collisionConfiguration;
+	btDispatcher *dispatcher;
+	btBroadphaseInterface *pairCache;
+	btConstraintSolver *constraintSolver;
+	btOverlapFilterCallback *filterCallback;
+};
+struct rbRigidBody {
+	btRigidBody *body;
+	int col_groups;
+};
+
+struct rbCollisionShape {
+	btCollisionShape *cshape;
+	btTriangleMesh *mesh;
+};
+
+struct rbFilterCallback : public btOverlapFilterCallback
+{
+	virtual bool needBroadphaseCollision(btBroadphaseProxy *proxy0, btBroadphaseProxy *proxy1) const
+	{
+		rbRigidBody *rb0 = (rbRigidBody *)((btRigidBody *)proxy0->m_clientObject)->getUserPointer();
+		rbRigidBody *rb1 = (rbRigidBody *)((btRigidBody *)proxy1->m_clientObject)->getUserPointer();
+		
+		bool collides;
+		collides = (proxy0->m_collisionFilterGroup & proxy1->m_collisionFilterMask) != 0;
+		collides = collides && (proxy1->m_collisionFilterGroup & proxy0->m_collisionFilterMask);
+		collides = collides && (rb0->col_groups & rb1->col_groups);
+		
+		return collides;
+	}
+};
+
+static inline void copy_v3_btvec3(float vec[3], const btVector3 &btvec)
+{
+	vec[0] = (float)btvec[0];
+	vec[1] = (float)btvec[1];
+	vec[2] = (float)btvec[2];
+}
+static inline void copy_quat_btquat(float quat[3], const btQuaternion &btquat)
+{
+	quat[0] = btquat.getW();
+	quat[1] = btquat.getX();
+	quat[2] = btquat.getY();
+	quat[3] = btquat.getZ();
+}
+
+/* ********************************** */
+/* Dynamics World Methods */
+
+/* Setup ---------------------------- */
+
+rbDynamicsWorld *RB_dworld_new(const float gravity[3])
+{
+	rbDynamicsWorld *world = new rbDynamicsWorld;
+	
+	/* collision detection/handling */
+	world->collisionConfiguration = new btDefaultCollisionConfiguration();
+	
+	world->dispatcher = new btCollisionDispatcher(world->collisionConfiguration);
+	btGImpactCollisionAlgorithm::registerAlgorithm((btCollisionDispatcher *)world->dispatcher); // XXX: experimental
+	
+	world->pairCache = new btDbvtBroadphase();
+	
+	world->filterCallback = new rbFilterCallback();
+	world->pairCache->getOverlappingPairCache()->setOverlapFilterCallback(world->filterCallback);
+
+	/* constraint solving */
+	world->constraintSolver = new btSequentialImpulseConstraintSolver();
+
+	/* world */
+	world->dynamicsWorld = new btDiscreteDynamicsWorld(world->dispatcher,
+	                                                   world->pairCache,
+	                                                   world->constraintSolver,
+	                                                   world->collisionConfiguration);
+
+	RB_dworld_set_gravity(world, gravity);
+	
+	return world;
+}
+
+void RB_dworld_delete(rbDynamicsWorld *world)
+{
+	/* bullet doesn't like if we free these in a different order */
+	delete world->dynamicsWorld;
+	delete world->constraintSolver;
+	delete world->pairCache;
+	delete world->dispatcher;
+	delete world->collisionConfiguration;
+	delete world->filterCallback;
+	delete world;
+}
+
+/* Settings ------------------------- */
+
+/* Gravity */
+void RB_dworld_get_gravity(rbDynamicsWorld *world, float g_out[3])
+{
+	copy_v3_btvec3(g_out, world->dynamicsWorld->getGravity());
+}
+
+void RB_dworld_set_gravity(rbDynamicsWorld *world, const float g_in[3])
+{
+	world->dynamicsWorld->setGravity(btVector3(g_in[0], g_in[1], g_in[2]));
+}
+
+/* Constraint Solver */
+void RB_dworld_set_solver_iterations(rbDynamicsWorld *world, int num_solver_iterations)
+{
+	btContactSolverInfo& info = world->dynamicsWorld->getSolverInfo();
+	
+	info.m_numIterations = num_solver_iterations;
+}
+
+/* Split Impulse */
+void RB_dworld_set_split_impulse(rbDynamicsWorld *world, int split_impulse)
+{
+	btContactSolverInfo& info = world->dynamicsWorld->getSolverInfo();
+	
+	info.m_splitImpulse = split_impulse;
+}
+
+/* Simulation ----------------------- */
+
+void RB_dworld_step_simulation(rbDynamicsWorld *world, float timeStep, int maxSubSteps, float timeSubStep)
+{
+	world->dynamicsWorld->stepSimulation(timeStep, maxSubSteps, timeSubStep);
+}
+
+/* Export -------------------------- */
+
+/* Exports entire dynamics world to Bullet's "*.bullet" binary format
+ * which is similar to Blender's SDNA system...
+ * < rbDynamicsWorld: dynamics world to write to file
+ * < filename: assumed to be a valid filename, with .bullet extension 
+ */
+void RB_dworld_export(rbDynamicsWorld *world, const char *filename)
+{
+	//create a large enough buffer. There is no method to pre-calculate the buffer size yet.
+	int maxSerializeBufferSize = 1024 * 1024 * 5;
+	
+	btDefaultSerializer *serializer = new btDefaultSerializer(maxSerializeBufferSize);
+	world->dynamicsWorld->serialize(serializer);
+	
+	FILE *file = fopen(filename, "wb");
+	fwrite(serializer->getBufferPointer(), serializer->getCurrentBufferSize(), 1, file);
+	fclose(file);
+}
+
+/* ********************************** */
+/* Rigid Body Methods */
+
+/* Setup ---------------------------- */
+
+void RB_dworld_add_body(rbDynamicsWorld *world, rbRigidBody *object, int col_groups)
+{
+	btRigidBody *body = object->body;
+	object->col_groups = col_groups;
+	
+	world->dynamicsWorld->addRigidBody(body);
+}
+
+void RB_dworld_remove_body(rbDynamicsWorld *world, rbRigidBody *object)
+{
+	btRigidBody *body = object->body;
+	
+	world->dynamicsWorld->removeRigidBody(body);
+}
+
+/* ............ */
+
+rbRigidBody *RB_body_new(rbCollisionShape *shape, const float loc[3], const float rot[4])
+{
+	rbRigidBody *object = new rbRigidBody;
+	/* current transform */
+	btTransform trans;
+	trans.setOrigin(btVector3(loc[0], loc[1], loc[2]));
+	trans.setRotation(btQuaternion(rot[1], rot[2], rot[3], rot[0]));
+	
+	/* create motionstate, which is necessary for interpolation (includes reverse playback) */
+	btDefaultMotionState *motionState = new btDefaultMotionState(trans);
+	
+	/* make rigidbody */
+	btRigidBody::btRigidBodyConstructionInfo rbInfo(1.0f, motionState, shape->cshape);
+	
+	object->body = new btRigidBody(rbInfo);
+	
+	object->body->setUserPointer(object);
+	
+	return object;
+}
+
+void RB_body_delete(rbRigidBody *object)
+{
+	btRigidBody *body = object->body;
+	
+	/* motion state */
+	btMotionState *ms = body->getMotionState();
+	if (ms)
+		delete ms;
+	
+	/* collision shape is done elsewhere... */
+	
+	/* body itself */
+	
+	/* manually remove constraint refs of the rigid body, normally this happens when removing constraints from the world
+	 * but since we delete everything when the world is rebult, we need to do it manually here */
+	for (int i = body->getNumConstraintRefs() - 1; i >= 0; i--) {
+		btTypedConstraint *con = body->getConstraintRef(i);
+		body->removeConstraintRef(con);
+	}
+	
+	delete body;
+	delete object;
+}
+
+/* Settings ------------------------- */
+
+void RB_body_set_collision_shape(rbRigidBody *object, rbCollisionShape *shape)
+{
+	btRigidBody *body = object->body;
+	
+	/* set new collision shape */
+	body->setCollisionShape(shape->cshape);
+	
+	/* recalculate inertia, since that depends on the collision shape... */
+	RB_body_set_mass(object, RB_body_get_mass(object));
+}
+
+/* ............ */
+
+float RB_body_get_mass(rbRigidBody *object)
+{
+	btRigidBody *body = object->body;
+	
+	/* there isn't really a mass setting, but rather 'inverse mass'  
+	 * which we convert back to mass by taking the reciprocal again 
+	 */
+	float value = (float)body->getInvMass();
+	
+	if (value)
+		value = 1.0 / value;
+		
+	return value;
+}
+
+void RB_body_set_mass(rbRigidBody *object, float value)
+{
+	btRigidBody *body = object->body;
+	btVector3 localInertia(0, 0, 0);
+	
+	/* calculate new inertia if non-zero mass */
+	if (value) {
+		btCollisionShape *shape = body->getCollisionShape();
+		shape->calculateLocalInertia(value, localInertia);
+	}
+	
+	body->setMassProps(value, localInertia);
+	body->updateInertiaTensor();
+}
+
+
+float RB_body_get_friction(rbRigidBody *object)
+{
+	btRigidBody *body = object->body;
+	return body->getFriction();
+}
+
+void RB_body_set_friction(rbRigidBody *object, float value)
+{
+	btRigidBody *body = object->body;
+	body->setFriction(value);
+}
+
+
+float RB_body_get_restitution(rbRigidBody *object)
+{
+	btRigidBody *body = object->body;
+	return body->getRestitution();
+}
+
+void RB_body_set_restitution(rbRigidBody *object, float value)
+{
+	btRigidBody *body = object->body;
+	body->setRestitution(value);
+}
+
+
+float RB_body_get_linear_damping(rbRigidBody *object)
+{
+	btRigidBody *body = object->body;
+	return body->getLinearDamping();
+}
+
+void RB_body_set_linear_damping(rbRigidBody *object, float value)
+{
+	RB_body_set_damping(object, value, RB_body_get_linear_damping(object));
+}
+
+float RB_body_get_angular_damping(rbRigidBody *object)
+{
+	btRigidBody *body = object->body;
+	return body->getAngularDamping();
+}
+
+void RB_body_set_angular_damping(rbRigidBody *object, float value)
+{
+	RB_body_set_damping(object, RB_body_get_linear_damping(object), value);
+}
+
+void RB_body_set_damping(rbRigidBody *object, float linear, float angular)
+{
+	btRigidBody *body = object->body;
+	body->setDamping(linear, angular);
+}
+
+
+float RB_body_get_linear_sleep_thresh(rbRigidBody *object)
+{
+	btRigidBody *body = object->body;
+	return body->getLinearSleepingThreshold();
+}
+
+void RB_body_set_linear_sleep_thresh(rbRigidBody *object, float value)
+{
+	RB_body_set_sleep_thresh(object, value, RB_body_get_angular_sleep_thresh(object));
+}
+
+float RB_body_get_angular_sleep_thresh(rbRigidBody *object)
+{
+	btRigidBody *body = object->body;
+	return body->getAngularSleepingThreshold();
+}
+
+void RB_body_set_angular_sleep_thresh(rbRigidBody *object, float value)
+{
+	RB_body_set_sleep_thresh(object, RB_body_get_linear_sleep_thresh(object), value);
+}
+
+void RB_body_set_sleep_thresh(rbRigidBody *object, float linear, float angular)
+{
+	btRigidBody *body = object->body;
+	body->setSleepingThresholds(linear, angular);
+}
+
+/* ............ */
+
+void RB_body_get_linear_velocity(rbRigidBody *object, float v_out[3])
+{
+	btRigidBody *body = object->body;
+	
+	copy_v3_btvec3(v_out, body->getLinearVelocity());
+}
+
+void RB_body_set_linear_velocity(rbRigidBody *object, const float v_in[3])
+{
+	btRigidBody *body = object->body;
+	
+	body->setLinearVelocity(btVector3(v_in[0], v_in[1], v_in[2]));
+}
+
+
+void RB_body_get_angular_velocity(rbRigidBody *object, float v_out[3])
+{
+	btRigidBody *body = object->body;
+	
+	copy_v3_btvec3(v_out, body->getAngularVelocity());
+}
+
+void RB_body_set_angular_velocity(rbRigidBody *object, const float v_in[3])
+{
+	btRigidBody *body = object->body;
+	
+	body->setAngularVelocity(btVector3(v_in[0], v_in[1], v_in[2]));
+}
+
+void RB_body_set_linear_factor(rbRigidBody *object, float x, float y, float z)
+{
+	btRigidBody *body = object->body;
+	body->setLinearFactor(btVector3(x, y, z));
+}
+
+void RB_body_set_angular_factor(rbRigidBody *object, float x, float y, float z)
+{
+	btRigidBody *body = object->body;
+	body->setAngularFactor(btVector3(x, y, z));
+}
+
+/* ............ */
+
+void RB_body_set_kinematic_state(rbRigidBody *object, int kinematic)
+{
+	btRigidBody *body = object->body;
+	if (kinematic)
+		body->setCollisionFlags(body->getCollisionFlags() | btCollisionObject::CF_KINEMATIC_OBJECT);
+	else
+		body->setCollisionFlags(body->getCollisionFlags() & ~btCollisionObject::CF_KINEMATIC_OBJECT);
+}
+
+/* ............ */
+
+void RB_body_set_activation_state(rbRigidBody *object, int use_deactivation)
+{
+	btRigidBody *body = object->body;
+	if (use_deactivation)
+		body->forceActivationState(ACTIVE_TAG);
+	else
+		body->setActivationState(DISABLE_DEACTIVATION);
+}
+void RB_body_activate(rbRigidBody *object)
+{
+	btRigidBody *body = object->body;
+	body->setActivationState(ACTIVE_TAG);
+}
+void RB_body_deactivate(rbRigidBody *object)
+{
+	btRigidBody *body = object->body;
+	body->setActivationState(ISLAND_SLEEPING);
+}
+
+/* ............ */
+
+
+
+/* Simulation ----------------------- */
+
+/* The transform matrices Blender uses are OpenGL-style matrices, 
+ * while Bullet uses the Right-Handed coordinate system style instead.
+ */
+
+void RB_body_get_transform_matrix(rbRigidBody *object, float m_out[4][4])
+{
+	btRigidBody *body = object->body;
+	btMotionState *ms = body->getMotionState();
+	
+	btTransform trans;
+	ms->getWorldTransform(trans);
+	
+	trans.getOpenGLMatrix((btScalar *)m_out);
+}
+
+void RB_body_set_loc_rot(rbRigidBody *object, const float loc[3], const float rot[4])
+{
+	btRigidBody *body = object->body;
+	btMotionState *ms = body->getMotionState();
+	
+	/* set transform matrix */
+	btTransform trans;
+	trans.setOrigin(btVector3(loc[0], loc[1], loc[2]));
+	trans.setRotation(btQuaternion(rot[1], rot[2], rot[3], rot[0]));
+	
+	ms->setWorldTransform(trans);
+}
+
+void RB_body_set_scale(rbRigidBody *object, const float scale[3])
+{
+	btRigidBody *body = object->body;
+	
+	/* apply scaling factor from matrix above to the collision shape */
+	btCollisionShape *cshape = body->getCollisionShape();
+	if (cshape) {
+		cshape->setLocalScaling(btVector3(scale[0], scale[1], scale[2]));
+		
+		/* GIimpact shapes have to be updated to take scaling into account */
+		if (cshape->getShapeType() == GIMPACT_SHAPE_PROXYTYPE)
+			((btGImpactMeshShape *)cshape)->updateBound();
+	}
+}
+
+/* ............ */
+/* Read-only state info about status of simulation */
+
+void RB_body_get_position(rbRigidBody *object, float v_out[3])
+{
+	btRigidBody *body = object->body;
+	
+	copy_v3_btvec3(v_out, body->getWorldTransform().getOrigin());
+}
+
+void RB_body_get_orientation(rbRigidBody *object, float v_out[4])
+{
+	btRigidBody *body = object->body;
+	
+	copy_quat_btquat(v_out, body->getWorldTransform().getRotation());
+}
+
+/* ............ */
+/* Overrides for simulation */
+
+void RB_body_apply_central_force(rbRigidBody *object, const float v_in[3])
+{
+	btRigidBody *body = object->body;
+	
+	body->applyCentralForce(btVector3(v_in[0], v_in[1], v_in[2]));
+}
+
+/* ********************************** */
+/* Collision Shape Methods */
+
+/* Setup (Standard Shapes) ----------- */
+
+rbCollisionShape *RB_shape_new_box(float x, float y, float z)
+{
+	rbCollisionShape *shape = new rbCollisionShape;
+	shape->cshape = new btBoxShape(btVector3(x, y, z));
+	shape->mesh = NULL;
+	return shape;
+}
+
+rbCollisionShape *RB_shape_new_sphere(float radius)
+{
+	rbCollisionShape *shape = new rbCollisionShape;
+	shape->cshape = new btSphereShape(radius);
+	shape->mesh = NULL;
+	return shape;
+}
+
+rbCollisionShape *RB_shape_new_capsule(float radius, float height)
+{
+	rbCollisionShape *shape = new rbCollisionShape;
+	shape->cshape = new btCapsuleShapeZ(radius, height);
+	shape->mesh = NULL;
+	return shape;
+}
+
+rbCollisionShape *RB_shape_new_cone(float radius, float height)
+{
+	rbCollisionShape *shape = new rbCollisionShape;
+	shape->cshape = new btConeShapeZ(radius, height);
+	shape->mesh = NULL;
+	return shape;
+}
+
+rbCollisionShape *RB_shape_new_cylinder(float radius, float height)
+{
+	rbCollisionShape *shape = new rbCollisionShape;
+	shape->cshape = new btCylinderShapeZ(btVector3(radius, radius, height));
+	shape->mesh = NULL;
+	return shape;
+}
+
+/* Setup (Convex Hull) ------------ */
+
+rbCollisionShape *RB_shape_new_convex_hull(float *verts, int stride, int count, float margin, bool *can_embed)
+{
+	btConvexHullComputer hull_computer = btConvexHullComputer();
+	
+	// try to embed the margin, if that fails don't shrink the hull
+	if (hull_computer.compute(verts, stride, count, margin, 0.0f) < 0.0f) {
+		hull_computer.compute(verts, stride, count, 0.0f, 0.0f);
+		*can_embed = false;
+	}
+	
+	rbCollisionShape *shape = new rbCollisionShape;
+	btConvexHullShape *hull_shape = new btConvexHullShape(&(hull_computer.vertices[0].getX()), hull_computer.vertices.size());
+	
+	shape->cshape = hull_shape;
+	shape->mesh = NULL;
+	return shape;
+}
+
+/* Setup (Triangle Mesh) ---------- */
+
+/* Need to call rbTriMeshNewData() followed by rbTriMeshAddTriangle() several times 
+ * to set up the mesh buffer BEFORE calling rbShapeNewTriMesh(). Otherwise,
+ * we get nasty crashes...
+ */
+
+rbMeshData *RB_trimesh_data_new()
+{
+	// XXX: welding threshold?
+	return (rbMeshData *) new btTriangleMesh(true, false);
+}
+ 
+void RB_trimesh_add_triangle(rbMeshData *mesh, const float v1[3], const float v2[3], const float v3[3])
+{
+	btTriangleMesh *meshData = reinterpret_cast<btTriangleMesh*>(mesh);
+	
+	/* cast vertices to usable forms for Bt-API */
+	btVector3 vtx1((btScalar)v1[0], (btScalar)v1[1], (btScalar)v1[2]);
+	btVector3 vtx2((btScalar)v2[0], (btScalar)v2[1], (btScalar)v2[2]);
+	btVector3 vtx3((btScalar)v3[0], (btScalar)v3[1], (btScalar)v3[2]);
+	
+	/* add to the mesh 
+	 *	- remove duplicated verts is enabled
+	 */
+	meshData->addTriangle(vtx1, vtx2, vtx3, false);
+}
+ 
+rbCollisionShape *RB_shape_new_trimesh(rbMeshData *mesh)
+{
+	rbCollisionShape *shape = new rbCollisionShape;
+	btTriangleMesh *tmesh = reinterpret_cast<btTriangleMesh*>(mesh);
+	
+	/* triangle-mesh we create is a BVH wrapper for triangle mesh data (for faster lookups) */
+	// RB_TODO perhaps we need to allow saving out this for performance when rebuilding?
+	btBvhTriangleMeshShape *unscaledShape = new btBvhTriangleMeshShape(tmesh, true, true);
+	
+	shape->cshape = new btScaledBvhTriangleMeshShape(unscaledShape, btVector3(1.0f, 1.0f, 1.0f));
+	shape->mesh = tmesh;
+	return shape;
+}
+
+rbCollisionShape *RB_shape_new_gimpact_mesh(rbMeshData *mesh)
+{
+	rbCollisionShape *shape = new rbCollisionShape;
+	/* interpret mesh buffer as btTriangleIndexVertexArray (i.e. an impl of btStridingMeshInterface) */
+	btTriangleMesh *tmesh = reinterpret_cast<btTriangleMesh*>(mesh);
+	
+	btGImpactMeshShape *gimpactShape = new btGImpactMeshShape(tmesh);
+	gimpactShape->updateBound(); // TODO: add this to the update collision margin call?
+	
+	shape->cshape = gimpactShape;
+	shape->mesh = tmesh;
+	return shape;
+}
+
+/* Cleanup --------------------------- */
+
+void RB_shape_delete(rbCollisionShape *shape)
+{
+	if (shape->cshape->getShapeType() == SCALED_TRIANGLE_MESH_SHAPE_PROXYTYPE) {
+		btBvhTriangleMeshShape *child_shape = ((btScaledBvhTriangleMeshShape *)shape->cshape)->getChildShape();
+		if (child_shape)
+			delete child_shape;
+	}
+	if (shape->mesh)
+		delete shape->mesh;
+	delete shape->cshape;
+	delete shape;
+}
+
+/* Settings --------------------------- */
+
+float RB_shape_get_margin(rbCollisionShape *shape)
+{
+	return shape->cshape->getMargin();
+}
+
+void RB_shape_set_margin(rbCollisionShape *shape, float value)
+{
+	shape->cshape->setMargin(value);
+}
+
+/* ********************************** */
+/* Constraints */
+
+/* Setup ----------------------------- */
+
+void RB_dworld_add_constraint(rbDynamicsWorld *world, rbConstraint *con, int disable_collisions)
+{
+	btTypedConstraint *constraint = reinterpret_cast<btTypedConstraint*>(con);
+	
+	world->dynamicsWorld->addConstraint(constraint, disable_collisions);
+}
+
+void RB_dworld_remove_constraint(rbDynamicsWorld *world, rbConstraint *con)
+{
+	btTypedConstraint *constraint = reinterpret_cast<btTypedConstraint*>(con);
+	
+	world->dynamicsWorld->removeConstraint(constraint);
+}
+
+/* ............ */
+
+static void make_constraint_transforms(btTransform &transform1, btTransform &transform2, btRigidBody *body1, btRigidBody *body2, float pivot[3], float orn[4])
+{
+	btTransform pivot_transform = btTransform();
+	pivot_transform.setOrigin(btVector3(pivot[0], pivot[1], pivot[2]));
+	pivot_transform.setRotation(btQuaternion(orn[1], orn[2], orn[3], orn[0]));
+	
+	transform1 = body1->getWorldTransform().inverse() * pivot_transform;
+	transform2 = body2->getWorldTransform().inverse() * pivot_transform;
+}
+
+rbConstraint *RB_constraint_new_point(float pivot[3], rbRigidBody *rb1, rbRigidBody *rb2)
+{
+	btRigidBody *body1 = rb1->body;
+	btRigidBody *body2 = rb2->body;
+	
+	btVector3 pivot1 = body1->getWorldTransform().inverse() * btVector3(pivot[0], pivot[1], pivot[2]);
+	btVector3 pivot2 = body2->getWorldTransform().inverse() * btVector3(pivot[0], pivot[1], pivot[2]);
+	
+	btTypedConstraint *con = new btPoint2PointConstraint(*body1, *body2, pivot1, pivot2);
+	
+	return (rbConstraint *)con;
+}
+
+rbConstraint *RB_constraint_new_fixed(float pivot[3], float orn[4], rbRigidBody *rb1, rbRigidBody *rb2)
+{
+	btRigidBody *body1 = rb1->body;
+	btRigidBody *body2 = rb2->body;
+	btTransform transform1;
+	btTransform transform2;
+	
+	make_constraint_transforms(transform1, transform2, body1, body2, pivot, orn);
+	
+	btGeneric6DofConstraint *con = new btGeneric6DofConstraint(*body1, *body2, transform1, transform2, true);
+	
+	/* lock all axes */
+	for (int i = 0; i < 6; i++)
+		con->setLimit(i, 0, 0);
+	
+	return (rbConstraint *)con;
+}
+
+rbConstraint *RB_constraint_new_hinge(float pivot[3], float orn[4], rbRigidBody *rb1, rbRigidBody *rb2)
+{
+	btRigidBody *body1 = rb1->body;
+	btRigidBody *body2 = rb2->body;
+	btTransform transform1;
+	btTransform transform2;
+	
+	make_constraint_transforms(transform1, transform2, body1, body2, pivot, orn);
+	
+	btHingeConstraint *con = new btHingeConstraint(*body1, *body2, transform1, transform2);
+	
+	return (rbConstraint *)con;
+}
+
+rbConstraint *RB_constraint_new_slider(float pivot[3], float orn[4], rbRigidBody *rb1, rbRigidBody *rb2)
+{
+	btRigidBody *body1 = rb1->body;
+	btRigidBody *body2 = rb2->body;
+	btTransform transform1;
+	btTransform transform2;
+	
+	make_constraint_transforms(transform1, transform2, body1, body2, pivot, orn);
+	
+	btSliderConstraint *con = new btSliderConstraint(*body1, *body2, transform1, transform2, true);
+	
+	return (rbConstraint *)con;
+}
+
+rbConstraint *RB_constraint_new_piston(float pivot[3], float orn[4], rbRigidBody *rb1, rbRigidBody *rb2)
+{
+	btRigidBody *body1 = rb1->body;
+	btRigidBody *body2 = rb2->body;
+	btTransform transform1;
+	btTransform transform2;
+	
+	make_constraint_transforms(transform1, transform2, body1, body2, pivot, orn);
+	
+	btSliderConstraint *con = new btSliderConstraint(*body1, *body2, transform1, transform2, true);
+	con->setUpperAngLimit(-1.0f); // unlock rotation axis
+	
+	return (rbConstraint *)con;
+}
+
+rbConstraint *RB_constraint_new_6dof(float pivot[3], float orn[4], rbRigidBody *rb1, rbRigidBody *rb2)
+{
+	btRigidBody *body1 = rb1->body;
+	btRigidBody *body2 = rb2->body;
+	btTransform transform1;
+	btTransform transform2;
+	
+	make_constraint_transforms(transform1, transform2, body1, body2, pivot, orn);
+	
+	btTypedConstraint *con = new btGeneric6DofConstraint(*body1, *body2, transform1, transform2, true);
+	
+	return (rbConstraint *)con;
+}
+
+rbConstraint *RB_constraint_new_6dof_spring(float pivot[3], float orn[4], rbRigidBody *rb1, rbRigidBody *rb2)
+{
+	btRigidBody *body1 = rb1->body;
+	btRigidBody *body2 = rb2->body;
+	btTransform transform1;
+	btTransform transform2;
+	
+	make_constraint_transforms(transform1, transform2, body1, body2, pivot, orn);
+	
+	btTypedConstraint *con = new btGeneric6DofSpringConstraint(*body1, *body2, transform1, transform2, true);
+	
+	return (rbConstraint *)con;
+}
+
+/* Cleanup ----------------------------- */
+
+void RB_constraint_delete(rbConstraint *con)
+{
+	btTypedConstraint *constraint = reinterpret_cast<btTypedConstraint*>(con);
+	delete constraint;
+}
+
+/* Settings ------------------------- */
+
+void RB_constraint_set_enabled(rbConstraint *con, int enabled)
+{
+	btTypedConstraint *constraint = reinterpret_cast<btTypedConstraint*>(con);
+	
+	constraint->setEnabled(enabled);
+}
+
+void RB_constraint_set_limits_hinge(rbConstraint *con, float lower, float upper)
+{
+	btHingeConstraint *constraint = reinterpret_cast<btHingeConstraint*>(con);
+	
+	// RB_TODO expose these
+	float softness = 0.9f;
+	float bias_factor = 0.3f;
+	float relaxation_factor = 1.0f;
+	
+	constraint->setLimit(lower, upper, softness, bias_factor, relaxation_factor);
+}
+
+void RB_constraint_set_limits_slider(rbConstraint *con, float lower, float upper)
+{
+	btSliderConstraint *constraint = reinterpret_cast<btSliderConstraint*>(con);
+	
+	constraint->setLowerLinLimit(lower);
+	constraint->setUpperLinLimit(upper);
+}
+
+void RB_constraint_set_limits_piston(rbConstraint *con, float lin_lower, float lin_upper, float ang_lower, float ang_upper)
+{
+	btSliderConstraint *constraint = reinterpret_cast<btSliderConstraint*>(con);
+	
+	constraint->setLowerLinLimit(lin_lower);
+	constraint->setUpperLinLimit(lin_upper);
+	constraint->setLowerAngLimit(ang_lower);
+	constraint->setUpperAngLimit(ang_upper);
+}
+
+void RB_constraint_set_limits_6dof(rbConstraint *con, int axis, float lower, float upper)
+{
+	btGeneric6DofConstraint *constraint = reinterpret_cast<btGeneric6DofConstraint*>(con);
+	
+	constraint->setLimit(axis, lower, upper);
+}
+
+void RB_constraint_set_stiffness_6dof_spring(rbConstraint *con, int axis, float stiffness)
+{
+	btGeneric6DofSpringConstraint *constraint = reinterpret_cast<btGeneric6DofSpringConstraint*>(con);
+	
+	constraint->setStiffness(axis, stiffness);
+}
+
+void RB_constraint_set_damping_6dof_spring(rbConstraint *con, int axis, float damping)
+{
+	btGeneric6DofSpringConstraint *constraint = reinterpret_cast<btGeneric6DofSpringConstraint*>(con);
+	
+	// invert damping range so that 0 = no damping
+	constraint->setDamping(axis, 1.0f - damping);
+}
+
+void RB_constraint_set_spring_6dof_spring(rbConstraint *con, int axis, int enable)
+{
+	btGeneric6DofSpringConstraint *constraint = reinterpret_cast<btGeneric6DofSpringConstraint*>(con);
+	
+	constraint->enableSpring(axis, enable);
+}
+
+void RB_constraint_set_equilibrium_6dof_spring(rbConstraint *con)
+{
+	btGeneric6DofSpringConstraint *constraint = reinterpret_cast<btGeneric6DofSpringConstraint*>(con);
+	
+	constraint->setEquilibriumPoint();
+}
+
+void RB_constraint_set_solver_iterations(rbConstraint *con, int num_solver_iterations)
+{
+	btTypedConstraint *constraint = reinterpret_cast<btTypedConstraint*>(con);
+	
+	constraint->setOverrideNumSolverIterations(num_solver_iterations);
+}
+
+void RB_constraint_set_breaking_threshold(rbConstraint *con, float threshold)
+{
+	btTypedConstraint *constraint = reinterpret_cast<btTypedConstraint*>(con);
+	
+	constraint->setBreakingImpulseThreshold(threshold);
+}
+
+/* ********************************** */
diff --git a/intern/smoke/intern/FLUID_3D.cpp b/intern/smoke/intern/FLUID_3D.cpp
index 4eb11a46f5b..a1d39286a90 100644
--- a/intern/smoke/intern/FLUID_3D.cpp
+++ b/intern/smoke/intern/FLUID_3D.cpp
@@ -1055,7 +1055,7 @@ void FLUID_3D::project()
 			for (x = 1; x < _xRes - 1; x++, index++)
 			{
 				float vMask[3] = {1.0f, 1.0f, 1.0f}, vObst[3] = {0, 0, 0};
-				float vR = 0.0f, vL = 0.0f, vT = 0.0f, vB = 0.0f, vD = 0.0f, vU = 0.0f;
+				// float vR = 0.0f, vL = 0.0f, vT = 0.0f, vB = 0.0f, vD = 0.0f, vU = 0.0f;  // UNUSED
 
 				float pC = _pressure[index]; // center
 				float pR = _pressure[index + 1]; // right