Hat mark buffers 1

Author: Gary Frost

hat/backends/ffi/mock/cpp/mock_backend.cpp

@@ -30,6 +30,12 @@ class MockBackend : public Backend {
     public :
     };
 
+      class MockQueue : public Backend::Queue {
+        public :
+         MockQueue():Backend::Queue(){}
+         virtual ~MockQueue(){}
+        };
+
     class MockProgram : public Backend::Program {
         class MockKernel : public Backend::Program::Kernel {
         public:
@@ -65,8 +71,8 @@ class MockBackend : public Backend {
 
 public:
 
-    MockBackend(MockConfig *mockConfig, int mockConfigSchemeLen, char *mockBackendSchema)
-            : Backend(mockConfig, mockConfigSchemeLen, mockBackendSchema) {
+    MockBackend(MockConfig *mockConfig, int mockConfigSchemeLen, char *mockBackendSchema, MockQueue *mockQueue)
+            : Backend(mockConfig, mockConfigSchemeLen, mockBackendSchema, mockQueue) {
         if (mockConfig == nullptr) {
             std::cout << "mockConfig == null" << std::endl;
         } else {
@@ -99,5 +105,6 @@ class MockBackend : public Backend {
 
 long getBackend(void *config, int configSchemaLen, char *configSchema) {
     MockBackend::MockConfig *mockConfig = (MockBackend::MockConfig *) config;
-    return (long) new MockBackend(mockConfig, configSchemaLen, configSchema);
+    MockBackend::MockQueue *mockQueue = (MockBackend::MockQueue *) new MockBackend::MockQueue();
+    return (long) new MockBackend(mockConfig, configSchemaLen, configSchema, mockQueue);
 }

hat/backends/ffi/opencl/cpp/opencl_backend.cpp

@@ -43,7 +43,18 @@ OpenCLBackend::OpenCLProgram::OpenCLKernel::OpenCLBuffer::OpenCLBuffer(Backend::
         std::cerr << OpenCLBackend::errorMsg(status) << std::endl;
         exit(1);
     }
-    arg->value.buffer.vendorPtr = static_cast<void *>(this);
+
+    BufferState_s * bufferState = BufferState_s::of(
+      arg->value.buffer.memorySegment,
+      arg->value.buffer.sizeInBytes
+      );
+    if (INFO){
+       bufferState->dump("on allocation before assign");
+    }
+    bufferState->vendorPtr =  static_cast<void *>(this);
+    if (INFO){
+        bufferState->dump("after assign ");
+    }
     if (INFO){
          std::cout << "created buffer " << std::endl;
     }
@@ -58,15 +69,15 @@ void OpenCLBackend::OpenCLProgram::OpenCLKernel::OpenCLBuffer::copyToDevice() {
      */
     auto openclKernel = dynamic_cast<OpenCLKernel *>(kernel);
     auto openclBackend = dynamic_cast<OpenCLBackend *>(openclKernel->program->backend);
-    cl_int status = clEnqueueWriteBuffer(openclBackend->command_queue,
+    cl_int status = clEnqueueWriteBuffer( dynamic_cast<OpenCLQueue *>(openclKernel->program->backend->queue)->command_queue,
                                          clMem,
                                          CL_FALSE,
                                          0,
                                          arg->value.buffer.sizeInBytes,
                                          arg->value.buffer.memorySegment,
-                                         openclKernel->eventc,
-                                         ((openclKernel->eventc == 0) ? NULL : openclKernel->events),
-                                         &(openclKernel->events[openclKernel->eventc]));
+                                         dynamic_cast<OpenCLQueue *>(openclKernel->program->backend->queue)->eventc,
+                                         ((dynamic_cast<OpenCLQueue *>(openclKernel->program->backend->queue)->eventc == 0) ? NULL : dynamic_cast<OpenCLQueue *>(openclKernel->program->backend->queue)->events),
+                                         &(dynamic_cast<OpenCLQueue *>(openclKernel->program->backend->queue)->events[dynamic_cast<OpenCLQueue *>(openclKernel->program->backend->queue)->eventc]));
 
 
 
@@ -76,7 +87,7 @@ void OpenCLBackend::OpenCLProgram::OpenCLKernel::OpenCLBuffer::copyToDevice() {
         std::cerr << OpenCLBackend::errorMsg(status) << std::endl;
         exit(1);
     }
-    openclKernel->eventc++;
+    dynamic_cast<OpenCLQueue *>(openclKernel->program->backend->queue)->eventc++;
     if (INFO){
         std::cout << "enqueued buffer copyToDevice " << std::endl;
     }
@@ -85,21 +96,21 @@ void OpenCLBackend::OpenCLProgram::OpenCLKernel::OpenCLBuffer::copyToDevice() {
 void OpenCLBackend::OpenCLProgram::OpenCLKernel::OpenCLBuffer::copyFromDevice() {
     auto openclKernel = dynamic_cast<OpenCLKernel *>(kernel);
     auto openclBackend = dynamic_cast<OpenCLBackend *>(openclKernel->program->backend);
-    cl_int status = clEnqueueReadBuffer(openclBackend->command_queue,
+    cl_int status = clEnqueueReadBuffer( dynamic_cast<OpenCLQueue *>(openclKernel->program->backend->queue)->command_queue,
                                         clMem,
                                         CL_FALSE,
                                         0,
                                         arg->value.buffer.sizeInBytes,
                                         arg->value.buffer.memorySegment,
-                                        openclKernel->eventc,
-                                        ((openclKernel->eventc == 0) ? NULL : openclKernel->events),
-                                        &(openclKernel->events[openclKernel->eventc]));
+                                        dynamic_cast<OpenCLQueue *>(openclKernel->program->backend->queue)->eventc,
+                                        ((dynamic_cast<OpenCLQueue *>(openclKernel->program->backend->queue)->eventc == 0) ? NULL : dynamic_cast<OpenCLQueue *>(openclKernel->program->backend->queue)->events),
+                                        &(dynamic_cast<OpenCLQueue *>(openclKernel->program->backend->queue)->events[dynamic_cast<OpenCLQueue *>(openclKernel->program->backend->queue)->eventc]));
 
     if (status != CL_SUCCESS) {
         std::cerr << OpenCLBackend::errorMsg(status) << std::endl;
         exit(1);
     }
-    openclKernel->eventc++;
+    dynamic_cast<OpenCLQueue *>(openclKernel->program->backend->queue)->eventc++;
     if (INFO){
        std::cout << "enqueued buffer copyFromDevice " << std::endl;
     }
@@ -110,8 +121,7 @@ OpenCLBackend::OpenCLProgram::OpenCLKernel::OpenCLBuffer::~OpenCLBuffer() {
 }
 
 OpenCLBackend::OpenCLProgram::OpenCLKernel::OpenCLKernel(Backend::Program *program, char* name, cl_kernel kernel)
-        : Backend::Program::Kernel(program, name), kernel(kernel), eventMax(0), events(nullptr),
-          eventc(0) {
+        : Backend::Program::Kernel(program, name), kernel(kernel){
 }
 
 OpenCLBackend::OpenCLProgram::OpenCLKernel::~OpenCLKernel() {
@@ -124,12 +134,13 @@ long OpenCLBackend::OpenCLProgram::OpenCLKernel::ndrange(void *argArray) {
     if (INFO){
        Sled::show(std::cout, argArray);
     }
-    if (events != nullptr || eventc != 0) {
-        std::cerr << "opencl issue, we might have leaked events!" << std::endl;
-    }
-    eventMax = argSled.argc() * 4 + 1;
-    eventc = 0;
-    events = new cl_event[eventMax];
+   // if (events != nullptr || eventc != 0) {
+     //   std::cerr << "opencl issue, we might have leaked events!" << std::endl;
+    //}
+   // eventMax = argSled.argc() * 4 + 1;
+    //eventc = 0;
+   // events = new cl_event[eventMax];
+    OpenCLQueue *openclQueue = dynamic_cast<OpenCLQueue *>(program->backend->queue);
     NDRange *ndrange = nullptr;
     for (int i = 0; i < argSled.argc(); i++) {
         Arg_s *arg = argSled.arg(i);
@@ -138,33 +149,6 @@ long OpenCLBackend::OpenCLProgram::OpenCLKernel::ndrange(void *argArray) {
                auto openclBuffer = new OpenCLBuffer(this, arg);
                 if (arg->idx == 0){
                     ndrange = static_cast<NDRange *>(arg->value.buffer.memorySegment);
-                }else{
-                   IfaceBufferBits_s *ifacebufferbitz = IfaceBufferBits_s::of(
-                      arg->value.buffer.memorySegment,
-                      arg->value.buffer.sizeInBytes
-                   );
-                   if (ifacebufferbitz->ok()){
-                      if (INFO){
-                         if (ifacebufferbitz->isJavaDirty()){
-                             printf(" java dirty (javaDirty:%08x)\n", ifacebufferbitz->payload.javaDirty);
-                         }else{
-                            printf(" NOT java dirty (javaDirty:%08x)\n", ifacebufferbitz->payload.javaDirty);
-                         }
-                         if (ifacebufferbitz->isGpuDirty()){
-                            printf(" gpu dirty (gpuDirty:%08x)\n", ifacebufferbitz->payload.gpuDirty);
-                         }else{
-                            printf(" NOT gpu dirty (gpuDirty:%08x)\n", ifacebufferbitz->payload.gpuDirty);
-                         }
-                      }
-                   }else{
-                      printf("bad magic \n");
-                      printf("(magic1:%016lx,", ifacebufferbitz->magic1);
-                      printf("javaDirty:%08x,", ifacebufferbitz->payload.javaDirty);
-                      printf("gpuDirty:%08x,", ifacebufferbitz->payload.gpuDirty);
-                      printf("unused[0]:%08x,", ifacebufferbitz->payload.unused[0]);
-                      printf("unused[1]:%08x,", ifacebufferbitz->payload.unused[1]);
-                      printf("magic2:%016lx)\n", ifacebufferbitz->magic2);
-                   }
                 }
                 openclBuffer->copyToDevice();
                 cl_int status = clSetKernelArg(kernel, arg->idx, sizeof(cl_mem), &openclBuffer->clMem);
@@ -214,7 +198,7 @@ long OpenCLBackend::OpenCLProgram::OpenCLKernel::ndrange(void *argArray) {
                 break;
             }
             default: {
-                std::cerr << "unexpected variant " << (char) arg->variant << std::endl;
+                std::cerr << "unexpected variant (ndrange) " << (char) arg->variant << std::endl;
                 exit(1);
             }
         }
@@ -226,15 +210,15 @@ long OpenCLBackend::OpenCLProgram::OpenCLKernel::ndrange(void *argArray) {
     }
     size_t dims = 1;
     cl_int status = clEnqueueNDRangeKernel(
-            dynamic_cast<OpenCLBackend *>(program->backend)->command_queue,
+            openclQueue->command_queue,
             kernel,
             dims,
             nullptr,
             &globalSize,
             nullptr,
-            eventc,
-            ((eventc == 0) ? nullptr : events),
-            &(events[eventc]));
+            openclQueue->eventc,
+            (openclQueue->eventc == 0) ? nullptr : openclQueue->events,
+            &(openclQueue->events[openclQueue->eventc]));
     if (status != CL_SUCCESS) {
         std::cerr << OpenCLBackend::errorMsg(status) << std::endl;
         exit(1);
@@ -244,34 +228,49 @@ long OpenCLBackend::OpenCLProgram::OpenCLKernel::ndrange(void *argArray) {
        std::cout <<  " globalSize=" << globalSize << " " << std::endl;
     }
 
-    eventc++;
+    openclQueue->eventc++;
     for (int i = 0; i < argSled.argc(); i++) {
         Arg_s *arg = argSled.arg(i);
         if (arg->variant == '&') {
-            static_cast<OpenCLBuffer *>(arg->value.buffer.vendorPtr)->copyFromDevice();
+            BufferState_s * bufferState = BufferState_s::of(
+              arg->value.buffer.memorySegment,
+              arg->value.buffer.sizeInBytes
+              );
+            static_cast<OpenCLBuffer *>(bufferState->vendorPtr)->copyFromDevice();
+            if (INFO){
+               bufferState->dump("After copy from device");
+            }
+
         }
     }
-    status = clWaitForEvents(eventc, events);
+    status = clWaitForEvents(openclQueue->eventc, openclQueue->events);
     if (status != CL_SUCCESS) {
         std::cerr << OpenCLBackend::errorMsg(status) << std::endl;
         exit(1);
     }
-    for (int i = 0; i < eventc; i++) {
-        status = clReleaseEvent(events[i]);
+    for (int i = 0; i < openclQueue->eventc; i++) {
+        status = clReleaseEvent(openclQueue->events[i]);
         if (status != CL_SUCCESS) {
             std::cerr << OpenCLBackend::errorMsg(status) << std::endl;
             exit(1);
         }
     }
-    delete[] events;
-    eventMax = 0;
-    eventc = 0;
-    events = nullptr;
+   // delete[] events;
+    //eventMax = 0;
+   openclQueue->eventc = 0;
+    //events = nullptr;
     for (int i = 0; i < argSled.argc(); i++) {
         Arg_s *arg = argSled.arg(i);
         if (arg->variant == '&') {
-            delete static_cast<OpenCLBuffer *>(arg->value.buffer.vendorPtr);
-            arg->value.buffer.vendorPtr = nullptr;
+            BufferState_s * bufferState = BufferState_s::of(
+                      arg->value.buffer.memorySegment,
+                      arg->value.buffer.sizeInBytes
+                      );
+            delete static_cast<OpenCLBuffer *>(bufferState->vendorPtr);
+            bufferState->vendorPtr = nullptr;
+            if (INFO){
+               bufferState->dump("After deleting buffer ");
+            }
         }
     }
     return 0;
@@ -297,7 +296,7 @@ bool OpenCLBackend::OpenCLProgram::programOK() {
 }
 
 OpenCLBackend::OpenCLBackend(OpenCLBackend::OpenCLConfig *openclConfig, int configSchemaLen, char *configSchema)
-        : Backend((Backend::Config *) openclConfig, configSchemaLen, configSchema) {
+        : Backend((Backend::Config *) openclConfig,  configSchemaLen, configSchema, (Backend::Queue *) new OpenCLQueue()) {
 
     if (INFO){
        if (openclConfig == nullptr) {
@@ -346,13 +345,14 @@ OpenCLBackend::OpenCLBackend(OpenCLBackend::OpenCLConfig *openclConfig, int conf
 
     cl_command_queue_properties queue_props = CL_QUEUE_PROFILING_ENABLE;
 
-    if ((command_queue = clCreateCommandQueue(context, device_ids[0], queue_props, &status)) == NULL ||
+    if ((dynamic_cast<OpenCLQueue *>(queue)->command_queue = clCreateCommandQueue(context, device_ids[0], queue_props, &status)) == NULL ||
         status != CL_SUCCESS) {
         clReleaseContext(context);
         delete[] platforms;
         delete[] device_ids;
         return;
     }
+
     device_id = device_ids[0];
     delete[] device_ids;
     delete[] platforms;
@@ -365,15 +365,16 @@ OpenCLBackend::OpenCLBackend()
 
 OpenCLBackend::~OpenCLBackend() {
     clReleaseContext(context);
-    clReleaseCommandQueue(command_queue);
+    clReleaseCommandQueue(dynamic_cast<OpenCLQueue *>(queue)->command_queue);
 }
 
 void OpenCLBackend::OpenCLProgram::OpenCLKernel::showEvents(int width) {
-    cl_ulong *samples = new cl_ulong[4 * eventc]; // queued, submit, start, end
+    OpenCLQueue * openclQueue =  dynamic_cast<OpenCLQueue *>(program->backend->queue);
+    cl_ulong *samples = new cl_ulong[4 * openclQueue->eventc]; // queued, submit, start, end
     int sample = 0;
     cl_ulong min;
     cl_ulong max;
-    for (int event = 0; event < eventc; event++) {
+    for (int event = 0; event < openclQueue->eventc; event++) {
         for (int type = 0; type < 4; type++) {
             cl_profiling_info info;
             switch (type) {
@@ -391,7 +392,7 @@ void OpenCLBackend::OpenCLProgram::OpenCLKernel::showEvents(int width) {
                     break;
             }
 
-            if ((clGetEventProfilingInfo(events[event], info, sizeof(samples[sample]), &samples[sample], NULL)) !=
+            if ((clGetEventProfilingInfo(openclQueue->events[event], info, sizeof(samples[sample]), &samples[sample], NULL)) !=
                 CL_SUCCESS) {
                 std::cerr << "failed to get profile info " << info << std::endl;
             }
@@ -414,7 +415,7 @@ void OpenCLBackend::OpenCLProgram::OpenCLKernel::showEvents(int width) {
     std::cout << "Range: " << range << "(ns)" << std::endl;
     std::cout << "Scale: " << scale << " range (ns) per char" << std::endl;
 
-    for (int event = 0; event < eventc; event++) {
+    for (int event = 0; event < openclQueue->eventc; event++) {
         cl_ulong queue = (samples[sample++] - min) / scale;
         cl_ulong submit = (samples[sample++] - min) / scale;
         cl_ulong start = (samples[sample++] - min) / scale;

hat/backends/ffi/opencl/include/opencl_backend.h

@@ -62,28 +62,31 @@ class OpenCLBackend : public Backend {
     public:
         boolean gpu;
     };
+    class OpenCLQueue : public Backend::Queue {
+    public:
+       size_t eventMax;
+       cl_event *events;
+       size_t eventc;
+       cl_command_queue command_queue;
+       OpenCLQueue():Backend::Queue(), eventMax(256), events(new cl_event[eventMax]), eventc(0){
+       }
+       virtual ~OpenCLQueue(){}
+    };
 
     class OpenCLProgram : public Backend::Program {
         class OpenCLKernel : public Backend::Program::Kernel {
 
             class OpenCLBuffer : public Backend::Program::Kernel::Buffer {
             public:
                 cl_mem clMem;
-
                 void copyToDevice();
-
                 void copyFromDevice();
-
                 OpenCLBuffer(Backend::Program::Kernel *kernel, Arg_s *arg);
-
                 virtual ~OpenCLBuffer();
             };
 
         private:
             cl_kernel kernel;
-            size_t eventMax;
-            cl_event *events;
-            size_t eventc;
         protected:
             void showEvents(int width);
         public:
@@ -110,18 +113,11 @@ class OpenCLBackend : public Backend {
 public:
     cl_platform_id platform_id;
     cl_context context;
-    cl_command_queue command_queue;
     cl_device_id device_id;
-
-
     OpenCLBackend();
-
     OpenCLBackend(OpenCLConfig *config, int configSchemaLen, char *configSchema);
-
     ~OpenCLBackend();
-
     int getMaxComputeUnits();
-
     void info();
     void dumpSled(std::ostream &out,void *argArray);
     char *dumpSchema(std::ostream &out,int depth, char *ptr, void *data);

hat/backends/ffi/shared/cpp/shared.cpp

@@ -109,7 +109,7 @@ void Sled::show(std::ostream &out, void *argArray) {
                 break;
             }
             default: {
-                std::cerr << "unexpected variant '" << (char) arg->variant << "'" << std::endl;
+                std::cerr << "unexpected variant (shared.cpp) '" << (char) arg->variant << "'" << std::endl;
                 exit(1);
             }
         }