Initial version of the builtins

2 files changed, 76 insertions, 30 deletions

diff --git a/src/builtins/coal-internal.h b/src/builtins/coal-internal.h
new file mode 100644
index 0000000..e7f0276
--- /dev/null
+++ b/src/builtins/coal-internal.h
@@ -0,0 +1,44 @@
+#ifndef COAL_INTERNAL_H
+#define COAL_INTERNAL_H
+
+#define __kernel
+
+typedef unsigned char uchar;
+typedef unsigned short ushort;
+typedef unsigned int uint;
+typedef unsigned long ulong;
+
+//typedef float float4 __attribute__((ext_vector_type(4)));
+#define COAL_VECTOR(type, len)                                  \
+   typedef type type##len __attribute__((ext_vector_type(len)))
+#define COAL_VECTOR_SET(type) \
+   COAL_VECTOR(type, 2);      \
+   COAL_VECTOR(type, 3);      \
+   COAL_VECTOR(type, 4)
+
+COAL_VECTOR_SET(char);
+COAL_VECTOR_SET(uchar);
+
+COAL_VECTOR_SET(short);
+COAL_VECTOR_SET(ushort);
+
+COAL_VECTOR_SET(int);
+COAL_VECTOR_SET(uint);
+
+COAL_VECTOR_SET(long);
+COAL_VECTOR_SET(ulong);
+
+COAL_VECTOR_SET(float);
+
+#undef COAL_VECTOR_SET
+#undef COAL_VECTOR
+
+#define __global __attribute__((address_space(1)))
+#define __local __attribute__((address_space(2)))
+#define __constant __attribute__((address_space(3)))
+#define __private __attribute__((address_space(4)))
+
+
+extern int get_global_id(int);
+
+#endif
diff --git a/src/compiler/compiler.cpp b/src/compiler/compiler.cpp
index 87fbc3f..813c4b3 100644
--- a/src/compiler/compiler.cpp
+++ b/src/compiler/compiler.cpp
@@ -39,26 +39,32 @@ Compiler::~Compiler()
 
 bool Compiler::init()
 {
-   llvm::sys::Path Path = llvm::sys::Path::GetCurrentDirectory();
+}
+
+llvm::Module * Compiler::compile(const std::string &text)
+{
    clang::TextDiagnosticPrinter *DiagClient =
-      new TextDiagnosticPrinter(llvm::errs(), DiagnosticOptions());
+      new clang::TextDiagnosticPrinter(llvm::errs(), DiagnosticOptions());
 
    clang::Diagnostic Diags(DiagClient);
-   clang::driver::Driver TheDriver(Path.str(), llvm::sys::getHostTriple(),
-                    "a.out", /*IsProduction=*/false, /*CXXIsProduction=*/false,
-                    Diags);
-   TheDriver.setTitle("clang interpreter");
+   clang::driver::Driver TheDriver(
+      "clc", llvm::sys::getHostTriple(),
+      "a.out", /*IsProduction=*/false, /*CXXIsProduction=*/false,
+      Diags);
+   TheDriver.setTitle("OpenCL Interpreter");
 
    // FIXME: This is a hack to try to force the driver to do something we can
    // recognize. We need to extend the driver library to support this use model
    // (basically, exactly one input, and the operation mode is hard wired).
-
-   llvm::SmallVector<const char *, 16> Args;
-   Args.push_back("-fsyntax-only");
-   llvm::OwningPtr<clang::driver::Compilation> C(TheDriver.BuildCompilation(Args.size(),
-                                                                    Args.data()));
+   llvm::SmallVector<const char *, 16> Args(4);
+   Args[0] = "clc";
+   Args[1] = "-fsyntax-only";
+   Args[2] = "-v";
+   Args[3] = "/home/zack/main.c";
+   llvm::OwningPtr<clang::driver::Compilation> C(
+      TheDriver.BuildCompilation(Args.size(), Args.data()));
    if (!C)
-      return false;
+      return 0;
 
    // FIXME: This is copied from ASTUnit.cpp; simplify and eliminate.
 
@@ -70,13 +76,13 @@ bool Compiler::init()
       llvm::raw_svector_ostream OS(Msg);
       C->PrintJob(OS, C->getJobs(), "; ", true);
       Diags.Report(diag::err_fe_expected_compiler_job) << OS.str();
-      return false;
+      return 0;
    }
 
    const driver::Command *Cmd = cast<driver::Command>(*Jobs.begin());
    if (llvm::StringRef(Cmd->getCreator().getName()) != "clang") {
       Diags.Report(diag::err_fe_expected_clang_command);
-      return false;
+      return 0;
    }
 
    // Initialize a compiler invocation object from the clang (-cc1) arguments.
@@ -98,31 +104,27 @@ bool Compiler::init()
    // FIXME: This is copied from cc1_main.cpp; simplify and eliminate.
 
    // Create a compiler instance to handle the actual work.
-   m_clang.setLLVMContext(new llvm::LLVMContext);
-   m_clang.setInvocation(CI.take());
+   CompilerInstance Clang;
+   Clang.setLLVMContext(new llvm::LLVMContext);
+   Clang.setInvocation(CI.take());
 
    // Create the compilers actual diagnostics engine.
-   m_clang.createDiagnostics(int(CCArgs.size()),const_cast<char**>(
-                                CCArgs.data()));
-   if (!m_clang.hasDiagnostics())
-      return false;
+   Clang.createDiagnostics(int(CCArgs.size()),const_cast<char**>(CCArgs.data()));
+   if (!Clang.hasDiagnostics())
+      return 0;
 
    // Infer the builtin include path if unspecified.
-   if (m_clang.getHeaderSearchOpts().UseBuiltinIncludes &&
-       m_clang.getHeaderSearchOpts().ResourceDir.empty()) {
+   if (Clang.getHeaderSearchOpts().UseBuiltinIncludes &&
+       Clang.getHeaderSearchOpts().ResourceDir.empty()) {
       assert(0);
-      //m_clang.getHeaderSearchOpts().ResourceDir =
+      //Clang.getHeaderSearchOpts().ResourceDir =
       //   CompilerInvocation::GetResourcesPath(argv[0], MainAddr);
    }
-   return true;
-}
 
-llvm::Module * Compiler::compile(const std::string &text)
-{
    // Create and execute the frontend to generate an LLVM bitcode module.
-   llvm::OwningPtr<CodeGenAction> act(new EmitLLVMOnlyAction());
-   if (!m_clang.ExecuteAction(*act))
+   llvm::OwningPtr<CodeGenAction> Act(new EmitLLVMOnlyAction());
+   if (!Clang.ExecuteAction(*Act))
       return 0;
 
-   return act->takeModule();
+   return Act->takeModule();
 }