[Metal] Backend skeleton compiles

bin/occa.cpp

@@ -6,6 +6,7 @@
 #include <occa/lang/modes/cuda.hpp>
 #include <occa/lang/modes/hip.hpp>
 #include <occa/lang/modes/opencl.hpp>
+#include <occa/lang/modes/metal.hpp>
 
 using namespace occa;
 
@@ -160,6 +161,8 @@ bool runTranslate(const json &args) {
     parser = new lang::okl::hipParser(kernelProps);
   } else if (mode == "OpenCL") {
     parser = new lang::okl::openclParser(kernelProps);
+  } else if (mode == "Metal") {
+    parser = new lang::okl::metalParser(kernelProps);
   }
 
   if (!parser) {
@@ -194,10 +197,10 @@ bool runTranslate(const json &args) {
       << "*/\n";
   }
 
-  if (printLauncher && (
-        (mode == "CUDA")
-        || (mode == "HIP")
-        || (mode == "OpenCL"))) {
+  if (printLauncher && ((mode == "CUDA")
+                        || (mode == "HIP")
+                        || (mode == "OpenCL")
+                        || (mode == "Metal"))) {
     launcherParser = &(((occa::lang::okl::withLauncher*) parser)->launcherParser);
     std::cout << launcherParser->toString();
   } else {
@@ -254,6 +257,7 @@ bool runEnv(const json &args) {
             << "    - OCCA_OPENMP_ENABLED        : " << envEcho("OCCA_OPENMP_ENABLED", OCCA_OPENMP_ENABLED) << "\n"
             << "    - OCCA_CUDA_ENABLED          : " << envEcho("OCCA_CUDA_ENABLED", OCCA_CUDA_ENABLED) << "\n"
             << "    - OCCA_OPENCL_ENABLED        : " << envEcho("OCCA_OPENCL_ENABLED", OCCA_OPENCL_ENABLED) << "\n"
+            << "    - OCCA_METAL_ENABLED         : " << envEcho("OCCA_METAL_ENABLED", OCCA_METAL_ENABLED) << "\n"
 
             << "  Run-Time Options:\n"
             << "    - OCCA_CXX                   : " << envEcho("OCCA_CXX") << "\n"

include/occa/core/launchedDevice.hpp

@@ -1,11 +1,16 @@
 #ifndef OCCA_CORE_LAUNCHEDDEVICE_HEADER
 #define OCCA_CORE_LAUNCHEDDEVICE_HEADER
 
+#include <vector>
+
 #include <occa/core/device.hpp>
+#include <occa/lang/kernelMetadata.hpp>
 #include <occa/lang/modes/withLauncher.hpp>
 #include <occa/tools/properties.hpp>
 
 namespace occa {
+  typedef std::vector<lang::kernelMetadata> orderedKernelMetadata;
+
   class launchedModeDevice_t : public modeDevice_t {
   public:
     launchedModeDevice_t(const occa::properties &properties_);
@@ -33,6 +38,11 @@ namespace occa {
                                       const std::string &kernelName,
                                       lang::kernelMetadata &launcherMetadata);
 
+    orderedKernelMetadata getLaunchedKernelsMetadata(
+      const std::string &kernelName,
+      lang::kernelMetadataMap &deviceMetadata
+    );
+
     //---[ Virtual Methods ]------------
     virtual lang::okl::withLauncher* createParser(const occa::properties &props) const = 0;
 

include/occa/lang/modes/metal.hpp

@@ -1,19 +1,18 @@
-#ifndef OCCA_LANG_MODES_OPENCL_HEADER
-#define OCCA_LANG_MODES_OPENCL_HEADER
+#ifndef OCCA_LANG_MODES_METAL_HEADER
+#define OCCA_LANG_MODES_METAL_HEADER
 
 #include <occa/lang/modes/withLauncher.hpp>
 
 namespace occa {
   namespace lang {
     namespace okl {
-      class openclParser : public withLauncher {
+      class metalParser : public withLauncher {
       public:
-        qualifier_t constant;
-        qualifier_t kernel;
-        qualifier_t global;
-        qualifier_t local;
+        qualifier_t kernel_q;
+        qualifier_t device_q;
+        qualifier_t shared_q;
 
-        openclParser(const occa::properties &settings_ = occa::properties());
+        metalParser(const occa::properties &settings_ = occa::properties());
 
         virtual void onClear();
         virtual void beforePreprocessing();
@@ -26,23 +25,12 @@ namespace occa {
 
         virtual std::string getInnerIterator(const int loopIndex);
 
-        void addExtensions();
-
-        void updateConstToConstant();
-
-        void setLocalQualifiers();
+        void setSharedQualifiers();
 
         static bool sharedVariableMatcher(exprNode &expr);
-        static bool updateScopeStructVariables(statement_t &smnt);
-        static void addStructToVariable(variable_t &var);
-        static void addStructToFunctionArgs(function_t &func);
 
         void addBarriers();
 
-        void addFunctionPrototypes();
-
-        void addStructQualifiers();
-
         void setupKernels();
 
         void migrateLocalDecls(functionDeclStatement &kernelSmnt);

include/occa/modes/metal.hpp

@@ -4,7 +4,7 @@
 #  ifndef OCCA_MODES_METAL_HEADER
 #  define OCCA_MODES_METAL_HEADER
 
-#include <occa/modes/metal/utils.hpp>
+// No utility methods yet
 
 #  endif
 #endif

include/occa/modes/metal/device.hpp

@@ -1,27 +1,26 @@
 #include <occa/defines.hpp>
 
-#if OCCA_OPENCL_ENABLED
-#  ifndef OCCA_MODES_OPENCL_DEVICE_HEADER
-#  define OCCA_MODES_OPENCL_DEVICE_HEADER
+#if OCCA_METAL_ENABLED
+#  ifndef OCCA_MODES_METAL_DEVICE_HEADER
+#  define OCCA_MODES_METAL_DEVICE_HEADER
 
 #include <occa/core/launchedDevice.hpp>
-#include <occa/modes/opencl/headers.hpp>
+#include <occa/modes/metal/headers.hpp>
 
 namespace occa {
-  namespace opencl {
+  namespace metal {
     class info_t;
 
     class device : public occa::launchedModeDevice_t {
-      friend cl_context getContext(occa::device device);
 
     private:
       mutable hash_t hash_;
 
     public:
-      int platformID, deviceID;
+      int deviceID;
 
-      cl_device_id clDevice;
-      cl_context clContext;
+      metalDevice_t metalDevice;
+      metalCommandQueue_t metalCommandQueue;
 
       device(const occa::properties &properties_);
       virtual ~device();
@@ -43,8 +42,6 @@ namespace occa {
       virtual void waitFor(streamTag tag);
       virtual double timeBetween(const streamTag &startTag,
                                  const streamTag &endTag);
-
-      cl_command_queue& getCommandQueue() const;
       //================================
 
       //---[ Kernel ]-------------------

include/occa/modes/metal/headers.hpp

@@ -1,10 +1,14 @@
-#include <occa/defines.hpp>
+#ifndef OCCA_MODES_METAL_HEADER_HEADER
+#  define OCCA_MODES_METAL_HEADER_HEADER
+
+namespace occa {
+  namespace metal {
+    typedef struct _metalBuffer_t*       metalBuffer_t;
+    typedef struct _metalCommandQueue_t* metalCommandQueue_t;
+    typedef struct _metalDevice_t*       metalDevice_t;
+    typedef struct _metalEvent_t*        metalEvent_t;
+    typedef struct _metalKernel_t*       metalKernel_t;
+  }
+}
 
-#if   (OCCA_OS & OCCA_LINUX_OS)
-#  include <CL/cl.h>
-#  include <CL/cl_gl.h>
-#elif (OCCA_OS & OCCA_MACOS_OS)
-#  include <OpenCL/OpenCl.h>
-#else
-#  include "CL/opencl.h"
 #endif

include/occa/modes/metal/kernel.hpp

@@ -1,42 +1,31 @@
 #include <occa/defines.hpp>
 
-#if OCCA_OPENCL_ENABLED
-#  ifndef OCCA_MODES_OPENCL_KERNEL_HEADER
-#  define OCCA_MODES_OPENCL_KERNEL_HEADER
+#if OCCA_METAL_ENABLED
+#  ifndef OCCA_MODES_METAL_KERNEL_HEADER
+#  define OCCA_MODES_METAL_KERNEL_HEADER
 
 #include <occa/core/launchedKernel.hpp>
-#include <occa/modes/opencl/headers.hpp>
-#include <occa/modes/opencl/utils.hpp>
+#include <occa/modes/metal/headers.hpp>
 
 namespace occa {
-  namespace opencl {
+  namespace metal {
     class device;
 
     class kernel : public occa::launchedModeKernel_t {
       friend class device;
-      friend cl_kernel getCLKernel(occa::kernel kernel);
 
     private:
-      cl_device_id clDevice;
-      cl_kernel clKernel;
+      metalDevice_t metalDevice;
+      metalKernel_t metalKernel;
 
     public:
       kernel(modeDevice_t *modeDevice_,
              const std::string &name_,
              const std::string &sourceFilename_,
              const occa::properties &properties_);
 
-      kernel(modeDevice_t *modeDevice_,
-             const std::string &name_,
-             const std::string &sourceFilename_,
-             cl_device_id clDevice_,
-             cl_kernel clKernel_,
-             const occa::properties &properties_);
-
       ~kernel();
 
-      cl_command_queue& getCommandQueue() const;
-
       int maxDims() const;
       dim maxOuterDims() const;
       dim maxInnerDims() const;

include/occa/modes/metal/memory.hpp

@@ -1,40 +1,28 @@
 #include <occa/defines.hpp>
 
-#if OCCA_OPENCL_ENABLED
-#  ifndef OCCA_MODES_OPENCL_MEMORY_HEADER
-#  define OCCA_MODES_OPENCL_MEMORY_HEADER
+#if OCCA_METAL_ENABLED
+#  ifndef OCCA_MODES_METAL_MEMORY_HEADER
+#  define OCCA_MODES_METAL_MEMORY_HEADER
 
 #include <occa/core/memory.hpp>
-#include <occa/modes/opencl/headers.hpp>
+#include <occa/modes/metal/headers.hpp>
 
 namespace occa {
-  namespace opencl {
+  namespace metal {
     class device;
 
     class memory : public occa::modeMemory_t {
-      friend class opencl::device;
-
-      friend cl_mem getCLMemory(occa::memory memory);
-
-      friend void* getMappedPtr(occa::memory memory);
-
-      friend occa::memory wrapMemory(occa::device device,
-                                     cl_mem clMem,
-                                     const udim_t bytes,
-                                     const occa::properties &props);
+      friend class metal::device;
 
     private:
-      cl_mem clMem;
-      void *mappedPtr;
+      metalBuffer_t metalBuffer;
 
     public:
       memory(modeDevice_t *modeDevice_,
              udim_t size_,
              const occa::properties &properties_ = occa::properties());
       ~memory();
 
-      cl_command_queue& getCommandQueue() const;
-
       kernelArg makeKernelArg() const;
 
       modeMemory_t* addOffset(const dim_t offset);

include/occa/modes/metal/registration.hpp

@@ -1,18 +1,18 @@
 #include <occa/defines.hpp>
 
-#if OCCA_OPENCL_ENABLED
-#  ifndef OCCA_MODES_OPENCL_REGISTRATION_HEADER
-#  define OCCA_MODES_OPENCL_REGISTRATION_HEADER
+#if OCCA_METAL_ENABLED
+#  ifndef OCCA_MODES_METAL_REGISTRATION_HEADER
+#  define OCCA_MODES_METAL_REGISTRATION_HEADER
 
 #include <occa/modes.hpp>
-#include <occa/modes/opencl/device.hpp>
-#include <occa/modes/opencl/kernel.hpp>
-#include <occa/modes/opencl/memory.hpp>
+#include <occa/modes/metal/device.hpp>
+#include <occa/modes/metal/kernel.hpp>
+#include <occa/modes/metal/memory.hpp>
 #include <occa/tools/styling.hpp>
 #include <occa/core/base.hpp>
 
 namespace occa {
-  namespace opencl {
+  namespace metal {
     class modeInfo : public modeInfo_v {
     public:
       modeInfo();
@@ -21,8 +21,8 @@ namespace occa {
       styling::section& getDescription();
     };
 
-    extern occa::mode<opencl::modeInfo,
-                      opencl::device> mode;
+    extern occa::mode<metal::modeInfo,
+                      metal::device> mode;
   }
 }
 

include/occa/modes/metal/stream.hpp

@@ -1,21 +1,21 @@
 #include <occa/defines.hpp>
 
-#if OCCA_OPENCL_ENABLED
-#  ifndef OCCA_MODES_OPENCL_STREAM_HEADER
-#  define OCCA_MODES_OPENCL_STREAM_HEADER
+#if OCCA_METAL_ENABLED
+#  ifndef OCCA_MODES_METAL_STREAM_HEADER
+#  define OCCA_MODES_METAL_STREAM_HEADER
 
 #include <occa/core/stream.hpp>
-#include <occa/modes/opencl/headers.hpp>
+#include <occa/modes/metal/headers.hpp>
 
 namespace occa {
-  namespace opencl {
+  namespace metal {
     class stream : public occa::modeStream_t {
     public:
-      cl_command_queue commandQueue;
+      metalCommandQueue_t metalCommandQueue;
 
       stream(modeDevice_t *modeDevice_,
              const occa::properties &properties_,
-             cl_command_queue commandQueue_);
+             metalCommandQueue_t metalCommandQueue_);
 
       virtual ~stream();
     };

include/occa/modes/metal/streamTag.hpp

@@ -1,21 +1,21 @@
 #include <occa/defines.hpp>
 
-#if OCCA_OPENCL_ENABLED
-#  ifndef OCCA_MODES_OPENCL_STREAMTAG_HEADER
-#  define OCCA_MODES_OPENCL_STREAMTAG_HEADER
+#if OCCA_METAL_ENABLED
+#  ifndef OCCA_MODES_METAL_STREAMTAG_HEADER
+#  define OCCA_MODES_METAL_STREAMTAG_HEADER
 
 #include <occa/core/streamTag.hpp>
-#include <occa/modes/opencl/headers.hpp>
+#include <occa/modes/metal/headers.hpp>
 
 namespace occa {
-  namespace opencl {
+  namespace metal {
     class streamTag : public occa::modeStreamTag_t {
     public:
-      cl_event clEvent;
+      metalEvent_t metalEvent;
       double time;
 
       streamTag(modeDevice_t *modeDevice_,
-                cl_event clEvent_);
+                metalEvent_t metalEvent_);
 
       virtual ~streamTag();