diff --git a/pyopencl/__init__.py b/pyopencl/__init__.py
index d0904c29f54e66cdb403233fd0c984c685d99c5f..bc8cf16c6da2a6115320e1792c9977e0d93ba554 100644
--- a/pyopencl/__init__.py
+++ b/pyopencl/__init__.py
@@ -555,7 +555,8 @@ class Program(object):
     def compile(self, options=[], devices=None, headers=[]):
         options_bytes, _ = self._process_build_options(self._context, options)
 
-        return self._prg.compile(options_bytes, devices, headers)
+        self._get_prg().compile(options_bytes, devices, headers)
+        return self
 
     def __eq__(self, other):
         return self._get_prg() == other._get_prg()
@@ -577,7 +578,9 @@ def create_program_with_built_in_kernels(context, devices, kernel_names):
 
 def link_program(context, programs, options=[], devices=None):
     options_bytes, _ = Program._process_build_options(context, options)
-    return Program(_Program.link(context, programs, options_bytes, devices))
+    programs = [prg._get_prg() for prg in programs]
+    raw_prg = _Program.link(context, programs, options_bytes, devices)
+    return Program(raw_prg)
 
 # }}}
 
diff --git a/src/wrap_cl.hpp b/src/wrap_cl.hpp
index ace4bd2556a37541dcb5cfdafdc5c9df963bd31e..601c3dd1aebacfbbfcf56f49a6b349f5b534e303 100644
--- a/src/wrap_cl.hpp
+++ b/src/wrap_cl.hpp
@@ -394,6 +394,8 @@
 
 namespace pyopencl
 {
+  class program;
+
   // {{{ error
   class error : public std::runtime_error
   {
@@ -401,11 +403,30 @@ namespace pyopencl
       std::string m_routine;
       cl_int m_code;
 
+      // This is here because clLinkProgram returns a program
+      // object *just* so that there is somewhere for it to
+      // stuff the linker logs. :/
+      bool m_program_initialized;
+      cl_program m_program;
+
     public:
       error(const char *routine, cl_int c, const char *msg="")
-        : std::runtime_error(msg), m_routine(routine), m_code(c)
+        : std::runtime_error(msg), m_routine(routine), m_code(c),
+        m_program_initialized(false), m_program(nullptr)
       { }
 
+      error(const char *routine, cl_program prg, cl_int c,
+          const char *msg="")
+        : std::runtime_error(msg), m_routine(routine), m_code(c),
+        m_program_initialized(true), m_program(prg)
+      { }
+
+      virtual ~error()
+      {
+        if (m_program_initialized)
+          clReleaseProgram(m_program);
+      }
+
       const std::string &routine() const
       {
         return m_routine;
@@ -423,6 +444,8 @@ namespace pyopencl
             || code() == CL_OUT_OF_HOST_MEMORY);
       }
 
+      program *get_program() const;
+
   };
 
   // }}}
@@ -4075,7 +4098,7 @@ namespace pyopencl
         &status_code);
 
     if (status_code != CL_SUCCESS)
-      throw pyopencl::error("clLinkPorgram", status_code);
+      throw pyopencl::error("clLinkProgram", result, status_code);
 
     try
     {
@@ -4736,6 +4759,11 @@ namespace pyopencl
 
   // {{{ deferred implementation bits
 
+  inline program *error::get_program() const
+  {
+    return new program(m_program, /* retain */ true);
+  }
+
   inline py::object create_mem_object_wrapper(cl_mem mem, bool retain=true)
   {
     cl_mem_object_type mem_obj_type;
diff --git a/src/wrap_constants.cpp b/src/wrap_constants.cpp
index 7b6a97f16fbf2083534d8175b0395981461c0b70..597e49c21219a7ef7220f2d3e7947bef16fb9cc2 100644
--- a/src/wrap_constants.cpp
+++ b/src/wrap_constants.cpp
@@ -108,6 +108,7 @@ void pyopencl_expose_constants(py::module &m)
       .DEF_SIMPLE_METHOD(code)
       .DEF_SIMPLE_METHOD(what)
       .DEF_SIMPLE_METHOD(is_out_of_memory)
+      .def("_program", &cls::get_program)
       ;
   }
 
diff --git a/test/test_wrapper.py b/test/test_wrapper.py
index 4d729642163b0bebc63f2bd356a6e7ff00868ab7..5652d754b635bc96a384ae8675105e4d7f184742 100644
--- a/test/test_wrapper.py
+++ b/test/test_wrapper.py
@@ -1042,6 +1042,35 @@ def test_map_dtype(ctx_factory, dtype):
         assert array.dtype == dt
 
 
+def test_compile_link(ctx_factory):
+    ctx = ctx_factory()
+
+    if ctx._get_cl_version() < (1, 2) or cl.get_cl_header_version() < (1, 2):
+        pytest.skip("Context and ICD loader must understand CL1.2 for compile/link")
+
+    platform = ctx.devices[0].platform
+    if platform.name == "Apple":
+        pytest.skip("Apple doesn't like our compile/link test")
+
+    queue = cl.CommandQueue(ctx)
+    vsink_prg = cl.Program(ctx, """//CL//
+        void value_sink(float x)
+        {
+        }
+        """).compile()
+    main_prg = cl.Program(ctx, """//CL//
+        void value_sink(float x);
+
+        __kernel void experiment()
+        {
+            value_sink(3.1415f + get_global_id(0));
+        }
+        """).compile()
+    z = cl.link_program(ctx, [vsink_prg, main_prg], devices=ctx.devices)
+    z.experiment(queue, (128**2,), (128,))
+    queue.finish()
+
+
 if __name__ == "__main__":
     # make sure that import failures get reported, instead of skipping the tests.
     import pyopencl  # noqa