OpenCL backend

include/oklt/core/target_backends.h

@@ -14,6 +14,7 @@ enum struct TargetBackend : unsigned char {
     CUDA,    ///< CUDA backend.
     HIP,     ///< HIP backend.
     DPCPP,   ///< DPCPP backend.
+    OPENCL,  ///< OPENCL backend.
 
     _LAUNCHER,  ///< Launcher backend.
 };

lib/CMakeLists.txt

@@ -105,6 +105,22 @@ set (OCCA_TRANSPILER_SOURCES
     attributes/backend/dpcpp/common.cpp
     attributes/backend/dpcpp/common.h
 
+    # OPENCL
+    attributes/backend/opencl/kernel.cpp
+    attributes/backend/opencl/translation_unit.cpp
+    attributes/backend/opencl/global_constant.cpp
+    attributes/backend/opencl/global_function.cpp
+    attributes/backend/opencl/outer.cpp
+    attributes/backend/opencl/inner.cpp
+    attributes/backend/opencl/tile.cpp
+    attributes/backend/opencl/shared.cpp
+    attributes/backend/opencl/restrict.cpp
+    attributes/backend/opencl/atomic.cpp
+    attributes/backend/opencl/barrier.cpp
+    attributes/backend/opencl/exclusive.cpp
+    attributes/backend/opencl/common.cpp
+    attributes/backend/opencl/common.h
+
     # Serial subset
     attributes/utils/serial_subset/empty.cpp
     attributes/utils/serial_subset/kernel.cpp

lib/attributes/backend/cuda/translation_unit.cpp

@@ -17,7 +17,7 @@ std::vector<std::string_view> getBackendHeader(SessionStage& s) {
 }
 
 HandleResult handleTranslationUnit(SessionStage& s, const TranslationUnitDecl& d) {
-    return handleTranslationUnit(s, d, getBackendHeader(s));
+    return handleTranslationUnit(s, d, getBackendHeader(s), {});
 }
 
 __attribute__((constructor)) void registerAttrBackend() {

lib/attributes/backend/dpcpp/translation_unit.cpp

@@ -11,7 +11,7 @@ const std::string_view SYCL_INCLUDE = "<CL/sycl.hpp>";
 const std::string_view SYCL_NS = "sycl";
 
 HandleResult handleTranslationUnitDpcpp(SessionStage& s, const clang::TranslationUnitDecl& decl) {
-    return oklt::handleTranslationUnit(s, decl, {SYCL_INCLUDE}, {SYCL_NS});
+    return oklt::handleTranslationUnit(s, decl, {SYCL_INCLUDE}, {}, {SYCL_NS});
 }
 
 __attribute__((constructor)) void registerTranslationUnitAttrBackend() {

lib/attributes/backend/hip/translation_unit.cpp

@@ -13,7 +13,7 @@ using namespace clang;
 
 const std::string HIP_RT_INC = "<hip/hip_runtime.h>";
 HandleResult handleTU(SessionStage& s, const TranslationUnitDecl& d) {
-    return handleTranslationUnit(s, d, {HIP_RT_INC});
+    return handleTranslationUnit(s, d, {HIP_RT_INC}, {});
 }
 
 __attribute__((constructor)) void registerAttrBackend() {

lib/attributes/backend/opencl/atomic.cpp

@@ -0,0 +1,31 @@
+#include "attributes/attribute_names.h"
+#include "attributes/utils/cuda_subset/handle.h"
+#include "core/handler_manager/backend_handler.h"
+#include "core/transpiler_session/session_stage.h"
+#include "core/utils/attributes.h"
+#include "core/utils/range_to_string.h"
+#include "pipeline/core/error_codes.h"
+
+#include <spdlog/spdlog.h>
+
+namespace {
+using namespace oklt;
+using namespace clang;
+
+HandleResult handleAtomicAttribute(SessionStage& stage, const Stmt& stmt, const Attr& attr) {
+    SPDLOG_DEBUG("Handle [@atomic] attribute (stmt)");
+
+    removeAttribute(stage, attr);
+    return {};
+}
+
+
+__attribute__((constructor)) void registerAttrBackend() {
+    auto ok = registerBackendHandler(
+        TargetBackend::OPENCL, ATOMIC_ATTR_NAME, handleAtomicAttribute);
+
+    if (!ok) {
+        SPDLOG_ERROR("[OPENCL] Failed to register {} attribute handler", ATOMIC_ATTR_NAME);
+    }
+}
+}  // namespace

lib/attributes/backend/opencl/barrier.cpp

@@ -0,0 +1,35 @@
+#include "attributes/attribute_names.h"
+#include "attributes/backend/opencl/common.h"
+#include "attributes/utils/cuda_subset/handle.h"
+#include "core/handler_manager/backend_handler.h"
+#include "core/utils/attributes.h"
+
+#include <clang/AST/Attr.h>
+#include <clang/AST/Stmt.h>
+#include <spdlog/spdlog.h>
+
+namespace {
+using namespace oklt;
+using namespace clang;
+
+const std::string BARRIER_STR = "barrier(CLK_LOCAL_MEM_FENCE);\n";
+
+HandleResult handleBarrierAttribute(SessionStage& s,
+                                    const clang::Stmt& stmt,
+                                    const clang::Attr& a) {
+    SPDLOG_DEBUG("Handle [@barrier] attribute");
+
+    SourceRange range(getAttrFullSourceRange(a).getBegin(), stmt.getEndLoc());
+    s.getRewriter().ReplaceText(range, BARRIER_STR);
+    return {};
+}
+
+__attribute__((constructor)) void registerAttrBackend() {
+    auto ok =
+        registerBackendHandler(TargetBackend::OPENCL, BARRIER_ATTR_NAME, handleBarrierAttribute);
+
+    if (!ok) {
+        SPDLOG_ERROR("[OPENCL] Failed to register {} attribute handler", BARRIER_ATTR_NAME);
+    }
+}
+}  // namespace

lib/attributes/backend/opencl/common.cpp

@@ -0,0 +1,59 @@
+#include "attributes/backend/dpcpp/common.h"
+#include "util/string_utils.hpp"
+#include "core/sema/okl_sema_ctx.h"
+#include "core/utils/range_to_string.h"
+
+#include <clang/Rewrite/Core/Rewriter.h>
+
+#include <spdlog/spdlog.h>
+
+namespace oklt::opencl {
+using namespace clang;
+
+std::string axisToStr(const Axis& axis) {
+    static std::map<Axis, std::string> mapping{{Axis::X, "0"}, {Axis::Y, "1"}, {Axis::Z, "2"}};
+    return mapping[axis];
+}
+
+std::string getIdxVariable(const AttributedLoop& loop) {
+    auto strAxis = axisToStr(loop.axis);
+    switch (loop.type) {
+        case (LoopType::Inner):
+            return util::fmt("get_local_id({})", strAxis).value();
+        case (LoopType::Outer):
+            return util::fmt("get_group_id({})", strAxis).value();
+        default:  // Incorrect case
+            return "";
+    }
+}
+std::string buildInnerOuterLoopIdxLine(const OklLoopInfo& forLoop,
+                                       const AttributedLoop& loop,
+                                       int& openedScopeCounter,
+                                       oklt::Rewriter& rewriter) {
+    static_cast<void>(openedScopeCounter);
+    auto idx = getIdxVariable(loop);
+    auto op = forLoop.IsInc() ? "+" : "-";
+
+    std::string res;
+    if (forLoop.isUnary()) {
+        res = std::move(util::fmt("{} {} = ({}) {} {};\n",
+                                  forLoop.var.typeName,
+                                  forLoop.var.name,
+                                  getLatestSourceText(forLoop.range.start, rewriter),
+                                  op,
+                                  idx)
+                            .value());
+    } else {
+        res = std::move(util::fmt("{} {} = ({}) {} (({}) * {});\n",
+                                  forLoop.var.typeName,
+                                  forLoop.var.name,
+                                  getLatestSourceText(forLoop.range.start, rewriter),
+                                  op,
+                                  getLatestSourceText(forLoop.inc.val, rewriter),
+                                  idx)
+                            .value());
+    }
+    return res;
+}
+
+}

lib/attributes/backend/opencl/common.h

@@ -0,0 +1,24 @@
+#include "core/rewriter/rewriter_proxy.h"
+#include "attributes/frontend/params/loop.h"
+
+#include <string>
+
+
+namespace clang {
+class Rewriter;
+}
+
+namespace oklt {
+struct OklLoopInfo;
+}
+
+namespace oklt::opencl {
+std::string axisToStr(const Axis& axis);
+std::string getIdxVariable(const AttributedLoop& loop);
+std::string buildInnerOuterLoopIdxLine(const OklLoopInfo& forLoop,
+                                       const AttributedLoop& loop,
+                                       int& openedScopeCounter,
+                                       oklt::Rewriter& rewriter);
+
+const std::string SYNC_THREADS_BARRIER = "barrier(CLK_LOCAL_MEM_FENCE)";
+}

lib/attributes/backend/opencl/exclusive.cpp

@@ -0,0 +1,76 @@
+#include "attributes/attribute_names.h"
+#include "attributes/utils/cuda_subset/handle.h"
+#include "attributes/utils/default_handlers.h"
+#include "core/handler_manager/backend_handler.h"
+#include "core/sema/okl_sema_ctx.h"
+#include "core/transpiler_session/session_stage.h"
+#include "core/utils/attributes.h"
+
+#include <clang/AST/Attr.h>
+#include <clang/AST/DeclBase.h>
+#include <spdlog/spdlog.h>
+
+namespace {
+using namespace oklt;
+using namespace clang;
+
+HandleResult handleExclusiveDeclAttribute(SessionStage& s , const Decl& decl, const Attr& a) {
+    SPDLOG_DEBUG("Handle [@exclusive] attribute (decl)");
+
+    removeAttribute(s, a);
+    return {};
+}
+
+HandleResult handleExclusiveVarDeclAttribute(SessionStage& s , const VarDecl& decl, const Attr& a) {
+    SPDLOG_DEBUG("Handle [@exclusive] attribute (decl)");
+
+    auto& sema = s.tryEmplaceUserCtx<OklSemaCtx>();
+    auto loopInfo = sema.getLoopInfo();
+    if (!loopInfo) {
+        return tl::make_unexpected(
+            Error{{}, "@exclusive: failed to fetch loop meta data from sema"});
+    }
+
+    auto compStmt = dyn_cast_or_null<CompoundStmt>(loopInfo->stmt.getBody());
+    if (!compStmt || !loopInfo->is(LoopType::Outer)) {
+        return tl::make_unexpected(
+            Error{{}, "Must define [@exclusive] variables between [@outer] and [@inner] loops"});
+    }
+
+    auto child = loopInfo->getFirstAttributedChild();
+    if (!child || !child->is(LoopType::Inner)) {
+        return tl::make_unexpected(
+            Error{{}, "Must define [@exclusive] variables between [@outer] and [@inner] loops"});
+    }
+
+    removeAttribute(s, a);
+    return {};
+}
+
+HandleResult handleExclusiveExprAttribute(SessionStage& s , const DeclRefExpr& expr,const Attr& a) {
+    SPDLOG_DEBUG("Handle [@exclusive] attribute (stmt)");
+
+    auto& sema = s.tryEmplaceUserCtx<OklSemaCtx>();
+    auto loopInfo = sema.getLoopInfo();
+    if (!loopInfo) {
+        return tl::make_unexpected(
+            Error{{}, "@exclusive: failed to fetch loop meta data from sema"});
+    }
+
+    removeAttribute(s, a);
+    return {};
+}
+
+__attribute__((constructor)) void registerAttrBackend() {
+    auto ok = registerBackendHandler(
+        TargetBackend::OPENCL, EXCLUSIVE_ATTR_NAME, handleExclusiveExprAttribute);
+    ok &= registerBackendHandler(
+        TargetBackend::OPENCL, EXCLUSIVE_ATTR_NAME, handleExclusiveDeclAttribute);
+    ok &= registerBackendHandler(
+        TargetBackend::OPENCL, EXCLUSIVE_ATTR_NAME, handleExclusiveVarDeclAttribute);
+
+    if (!ok) {
+        SPDLOG_ERROR("[OPENCL] Failed to register {} attribute handler", EXCLUSIVE_ATTR_NAME);
+    }
+}
+}  // namespace

lib/attributes/backend/opencl/global_constant.cpp

@@ -0,0 +1,22 @@
+#include "attributes/utils/replace_attribute.h"
+#include "core/handler_manager/implicid_handler.h"
+
+#include <spdlog/spdlog.h>
+
+namespace {
+using namespace oklt;
+using namespace clang;
+
+HandleResult handleGlobalConstant(oklt::SessionStage& s, const clang::VarDecl& decl) {
+    const std::string OPENCL_CONST_QUALIFIER = "__constant";
+    return oklt::handleGlobalConstant(s, decl, OPENCL_CONST_QUALIFIER);
+}
+
+__attribute__((constructor)) void registeCUDAGlobalConstantHandler() {
+    auto ok = registerImplicitHandler(TargetBackend::OPENCL, handleGlobalConstant);
+
+    if (!ok) {
+        SPDLOG_ERROR("[OPENCL] Failed to register implicit handler for global constant");
+    }
+}
+}  // namespace

lib/attributes/backend/opencl/global_function.cpp

@@ -0,0 +1,54 @@
+#include "attributes/utils/replace_attribute.h"
+#include "attributes/attribute_names.h"
+#include "core/handler_manager/implicid_handler.h"
+#include "core/transpiler_session/header_info.h"
+#include "core/transpiler_session/session_stage.h"
+#include "core/utils/var_decl.h"
+
+#include <clang/AST/AST.h>
+#include <spdlog/spdlog.h>
+
+namespace {
+using namespace oklt;
+using namespace clang;
+
+HandleResult handleGlobalFunction(oklt::SessionStage& s, const clang::FunctionDecl& decl) {
+    if (decl.getLocation().isInvalid() || decl.isInlineBuiltinDeclaration() || !decl.hasBody()) {
+        return {};
+    }
+
+    if (decl.hasAttrs()) {
+        for (auto* attr : decl.getAttrs()) {
+            if (attr->getNormalizedFullName() == KERNEL_ATTR_NAME) {
+                SPDLOG_DEBUG(
+                    "Global function handler skipped function {}, since it has @kernel attribute",
+                    decl.getNameAsString());
+                return {};
+            }
+        }
+    }
+
+    auto& r = s.getRewriter();  
+    auto loc = decl.getFunctionTypeLoc();
+    auto funcr = SourceRange(decl.getBeginLoc(), loc.getRParenLoc());
+    auto str = r.getRewrittenText(funcr);
+
+    str += ";\n";
+
+    SPDLOG_DEBUG("Handle global function '{}' at {}",
+                 decl.getNameAsString(),
+                 decl.getLocation().printToString(s.getCompiler().getSourceManager()));
+
+    r.InsertTextBefore(decl.getSourceRange().getBegin(), str);
+
+    return {};
+}
+
+__attribute__((constructor)) void registerTranslationUnitAttrBackend() {
+    auto ok = registerImplicitHandler(TargetBackend::OPENCL, handleGlobalFunction);
+
+    if (!ok) {
+        SPDLOG_ERROR("[OPENCL] Failed to register implicit handler for global function");
+    }
+}
+}  // namespace