Fix a crash by deleting CPUProgram before the LLVM module.

This commit also fixes memory leaks.

10 files changed, 53 insertions, 1 deletions

diff --git a/src/core/commandqueue.cpp b/src/core/commandqueue.cpp
index 1651457..1a2ddda 100644
--- a/src/core/commandqueue.cpp
+++ b/src/core/commandqueue.cpp
@@ -389,6 +389,14 @@ Event::Event(CommandQueue *parent,
 
 Event::~Event()
 {
+    if (p_parent && p_device_data)
+    {
+        DeviceInterface *device = 0;
+        p_parent->info(CL_QUEUE_DEVICE, sizeof(DeviceInterface *), &device, 0);
+
+        device->freeEventDeviceData(this);
+    }
+
     for (int i=0; i<p_num_events_in_wait_list; ++i)
         clReleaseEvent((cl_event)p_event_wait_list[i]);
 
diff --git a/src/core/cpu/builtins.cpp b/src/core/cpu/builtins.cpp
index f3ea769..69ba404 100644
--- a/src/core/cpu/builtins.cpp
+++ b/src/core/cpu/builtins.cpp
@@ -2,6 +2,11 @@
 
 __thread Coal::CPUKernelWorkGroup *g_work_group;
 
+size_t get_global_id(cl_uint dimindx)
+{
+    return g_work_group->getGlobalId(dimindx);
+}
+
 void setThreadLocalWorkGroup(Coal::CPUKernelWorkGroup *current)
 {
     g_work_group = current;
@@ -9,5 +14,8 @@ void setThreadLocalWorkGroup(Coal::CPUKernelWorkGroup *current)
 
 void *getBuiltin(const std::string &name)
 {
+    if (name == "get_global_id")
+        return (void *)&get_global_id;
+
     return 0;
 }
diff --git a/src/core/cpu/device.cpp b/src/core/cpu/device.cpp
index e9b68dc..4dcd024 100644
--- a/src/core/cpu/device.cpp
+++ b/src/core/cpu/device.cpp
@@ -119,6 +119,21 @@ cl_int CPUDevice::initEventDeviceData(Event *event)
     return CL_SUCCESS;
 }
 
+void CPUDevice::freeEventDeviceData(Event *event)
+{
+    switch (event->type())
+    {
+        case Event::NDRangeKernel:
+        case Event::TaskKernel:
+        {
+            CPUKernelEvent *cpu_e = (CPUKernelEvent *)event->deviceData();
+
+            if (cpu_e)
+                delete cpu_e;
+        }
+    }
+}
+
 void CPUDevice::pushEvent(Event *event)
 {
     // Add an event in the list
diff --git a/src/core/cpu/device.h b/src/core/cpu/device.h
index f19fcbc..c8807a1 100644
--- a/src/core/cpu/device.h
+++ b/src/core/cpu/device.h
@@ -31,6 +31,7 @@ class CPUDevice : public DeviceInterface
                                          llvm::Function *function);
 
         cl_int initEventDeviceData(Event *event);
+        void freeEventDeviceData(Event *event);
 
         void pushEvent(Event *event);
         Event *getEvent(bool &stop);
diff --git a/src/core/cpu/kernel.cpp b/src/core/cpu/kernel.cpp
index db6fcb3..c1180eb 100644
--- a/src/core/cpu/kernel.cpp
+++ b/src/core/cpu/kernel.cpp
@@ -456,3 +456,12 @@ bool CPUKernelWorkGroup::run()
 
     return true;
 }
+
+size_t CPUKernelWorkGroup::getGlobalId(cl_uint dimindx) const
+{
+    if (dimindx > p_event->work_dim())
+        return 0;
+
+    return (p_index[dimindx] * p_event->local_work_size(dimindx))
+           + p_event->global_work_offset(dimindx) + p_current[dimindx];
+}
diff --git a/src/core/cpu/kernel.h b/src/core/cpu/kernel.h
index 787bef6..fbe3b6c 100644
--- a/src/core/cpu/kernel.h
+++ b/src/core/cpu/kernel.h
@@ -54,6 +54,9 @@ class CPUKernelWorkGroup
 
         bool run();
 
+        // Native functions
+        size_t getGlobalId(cl_uint dimindx) const;
+
     private:
         CPUKernel *p_kernel;
         KernelEvent *p_event;
diff --git a/src/core/cpu/program.cpp b/src/core/cpu/program.cpp
index fb1212e..653fb1e 100644
--- a/src/core/cpu/program.cpp
+++ b/src/core/cpu/program.cpp
@@ -27,7 +27,12 @@ CPUProgram::CPUProgram(CPUDevice *device, Program *program)
 CPUProgram::~CPUProgram()
 {
     if (p_jit)
+    {
+        // Dont delete the module
+        p_jit->removeModule(p_module);
+
         delete p_jit;
+    }
 }
 
 bool CPUProgram::linkStdLib() const
diff --git a/src/core/cpu/worker.cpp b/src/core/cpu/worker.cpp
index b2d3c05..90e2668 100644
--- a/src/core/cpu/worker.cpp
+++ b/src/core/cpu/worker.cpp
@@ -90,6 +90,8 @@ void *worker(void *data)
                 if (!instance->run())
                     errcode = CL_INVALID_PROGRAM_EXECUTABLE;
 
+                delete instance;
+
                 break;
             }
             default:
diff --git a/src/core/deviceinterface.h b/src/core/deviceinterface.h
index bca8f13..215642d 100644
--- a/src/core/deviceinterface.h
+++ b/src/core/deviceinterface.h
@@ -42,6 +42,7 @@ class DeviceInterface
 
         /** @note must set mapping address of MapBuffer events */
         virtual cl_int initEventDeviceData(Event *event) = 0;
+	virtual void freeEventDeviceData(Event *event) = 0;
 };
 
 class DeviceBuffer
diff --git a/src/core/program.cpp b/src/core/program.cpp
index 91c8d17..828cd7f 100644
--- a/src/core/program.cpp
+++ b/src/core/program.cpp
@@ -50,8 +50,8 @@ Program::~Program()
         DeviceDependent &dep = p_device_dependent.back();
 
         delete dep.compiler;
-        delete dep.linked_module;
         delete dep.program;
+        delete dep.linked_module;
 
         p_device_dependent.pop_back();
     }