Added a new define for the debug mode (GBE_DEBUG)

diff --git a/backend/CMakeLists.txt b/backend/CMakeLists.txt
index d235af9..cea4e58 100644 (file)
--- a/backend/CMakeLists.txt
+++ b/backend/CMakeLists.txt
@@ -60,23 +60,26 @@ if (CMAKE_SYSTEM_NAME STREQUAL "Linux")
     if (NOT GBE_COMPILE_UTESTS)
       set (CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -fno-exceptions")
     endif (NOT GBE_COMPILE_UTESTS)
-    set (CMAKE_CXX_FLAGS_RELEASE "-O2 -DNDEBUG -ftree-vectorize")
+    set (CMAKE_CXX_FLAGS_DEBUG          "-g -DGBE_DEBUG=1")
+    set (CMAKE_CXX_FLAGS_RELWITHDEBINFO "-O2 -g -DGBE_DEBUG=1")
+    set (CMAKE_CXX_FLAGS_MINSIZEREL     "-Os -DNDEBUG -DGBE_DEBUG=0")
+    set (CMAKE_CXX_FLAGS_RELEASE        "-O3 -DNDEBUG -DGBE_DEBUG=0")
   elseif (COMPILER STREQUAL "CLANG")
     set (CMAKE_C_COMPILER             "clang")
     set (CMAKE_C_FLAGS                "-Wall -std=c99")
-    set (CMAKE_C_FLAGS_DEBUG          "-g")
-    set (CMAKE_C_FLAGS_MINSIZEREL     "-Os -DNDEBUG")
-    set (CMAKE_C_FLAGS_RELEASE        "-O3 -DNDEBUG")
-    set (CMAKE_C_FLAGS_RELWITHDEBINFO "-O2 -g")
+    set (CMAKE_C_FLAGS_DEBUG          "-g -DGBE_DEBUG=1")
+    set (CMAKE_C_FLAGS_RELWITHDEBINFO "-O2 -g -DGBE_DEBUG=1")
+    set (CMAKE_C_FLAGS_MINSIZEREL     "-Os -DNDEBUG -DGBE_DEBUG=0")
+    set (CMAKE_C_FLAGS_RELEASE        "-O3 -DNDEBUG -DGBE_DEBUG=0")
     set (CMAKE_CXX_COMPILER             "clang++")
     set (CMAKE_CXX_FLAGS "-fstrict-aliasing -msse2 -ffast-math -fPIC -Wall -Wno-format-security -Wno-invalid-offsetof -std=c++0x")
     set (CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${GBE_DEBUG_MEMORY_FLAG}")
     set (CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${GBE_COMPILE_UTESTS_FLAG}")
     set (CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${VISIBILITY_FLAG}")
-    set (CMAKE_CXX_FLAGS_DEBUG          "-g")
-    set (CMAKE_CXX_FLAGS_MINSIZEREL     "-Os -DNDEBUG")
-    set (CMAKE_CXX_FLAGS_RELEASE        "-O3 -DNDEBUG")
-    set (CMAKE_CXX_FLAGS_RELWITHDEBINFO "-O2 -g")
+    set (CMAKE_CXX_FLAGS_DEBUG          "-g -DGBE_DEBUG=1")
+    set (CMAKE_CXX_FLAGS_RELWITHDEBINFO "-O2 -g -DGBE_DEBUG=1")
+    set (CMAKE_CXX_FLAGS_MINSIZEREL     "-Os -DNDEBUG -DGBE_DEBUG=0")
+    set (CMAKE_CXX_FLAGS_RELEASE        "-O3 -DNDEBUG -DGBE_DEBUG=0")
     set (CMAKE_AR      "/usr/bin/llvm-ar")
     set (CMAKE_LINKER  "/usr/bin/llvm-ld")
     set (CMAKE_NM      "/usr/bin/llvm-nm")
@@ -89,10 +92,11 @@ if (CMAKE_SYSTEM_NAME STREQUAL "Linux")
     set (CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${GBE_DEBUG_MEMORY_FLAG}")
     set (CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${GBE_COMPILE_UTESTS_FLAG}")
     set (CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${VISIBILITY_FLAG} -Wl,-E")
-    set (CMAKE_CXX_FLAGS_DEBUG "-g -O0")
-    set (CMAKE_CXX_FLAGS_RELEASE "-DNDEBUG -O2")
-    set (CCMAKE_CXX_FLAGS_RELWITHDEBINFO "-g -O2")
-    set (CCMAKE_CXX_FLAGS_MINSIZEREL "-Os")
+    set (CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${GBE_DEBUG_MODE_FLAG}")
+    set (CMAKE_CXX_FLAGS_DEBUG "-g -O0 -DGBE_DEBUG=1")
+    set (CCMAKE_CXX_FLAGS_RELWITHDEBINFO "-g -O2 -DGBE_DEBUG=1")
+    set (CMAKE_CXX_FLAGS_RELEASE "-DNDEBUG -O2 -DGBE_DEBUG=0")
+    set (CCMAKE_CXX_FLAGS_MINSIZEREL "-Os -DGBE_DEBUG=0")
     set (CMAKE_EXE_LINKER_FLAGS "")
   endif ()
 
@@ -102,7 +106,8 @@ elseif (CMAKE_SYSTEM_NAME STREQUAL "Windows")
     set (CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -Wno-invalid-offsetof -fstrict-aliasing -msse2 -ffast-math -Wall -fno-rtti -std=c++0x")
     set (CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${GBE_DEBUG_MEMORY_FLAG}")
     set (CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${GBE_COMPILE_UTESTS_FLAG}")
-    set (CMAKE_CXX_FLAGS_RELEASE "-O2 -DNDEBUG")
+    set (CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${GBE_DEBUG_MODE_FLAG}")
+    set (CMAKE_CXX_FLAGS_RELEASE "-O2 -DNDEBUG -DGBE_DEBUG=0")
   else (MINGW)
     set (COMMON_FLAGS "${GBE_DEBUG_MEMORY_FLAG} /arch:SSE2 /D_CRT_SECURE_NO_WARNINGS /DNOMINMAX /GR- /GS- /W3 /wd4275")
     set (CMAKE_CXX_FLAGS ${COMMON_FLAGS})

diff --git a/backend/kernels/struct.cl b/backend/kernels/struct.cl
index b9ec661..5915d31 100644 (file)
--- a/backend/kernels/struct.cl
+++ b/backend/kernels/struct.cl
@@ -3,8 +3,13 @@ struct my_struct {
   int b[2];
 };
 
-void struct_cl (struct my_struct s)
-{
+__constant int g[4] = {0,1,2,3};
 
+__kernel void struct_cl (struct my_struct s, int x, __global int *mem)
+{
+  __local int array[256];
+  for (int i = 0; i < 256; ++i)
+    array[i] = i;
+  mem[0] = array[x] + g[x];
 }
 

diff --git a/backend/src/ir/context.cpp b/backend/src/ir/context.cpp
index ffc7c8c..dcf57e4 100644 (file)
--- a/backend/src/ir/context.cpp
+++ b/backend/src/ir/context.cpp
@@ -125,10 +125,10 @@ namespace ir {
     // Append the instruction in the stream
     Instruction *insnPtr = fn->newInstruction();
     *insnPtr = insn;
-#ifndef NDEBUG
+#if GBE_DEBUG
     std::string whyNot;
     GBE_ASSERTM(insn.wellFormed(*fn, whyNot), whyNot.c_str());
-#endif /* NDEBUG */
+#endif /* GBE_DEBUG */
     bb->append(*insnPtr);
 
     // Close the current block if this is a branch

diff --git a/backend/src/ir/register.hpp b/backend/src/ir/register.hpp
index 81f769d..b2acfdd 100644 (file)
--- a/backend/src/ir/register.hpp
+++ b/backend/src/ir/register.hpp
@@ -25,6 +25,7 @@
 #define __GBE_IR_REGISTER_HPP__
 
 #include "sys/vector.hpp"
+#include "sys/platform.hpp"
 
 namespace gbe {
 namespace ir {

diff --git a/backend/src/llvm/CMakeLists.txt b/backend/src/llvm/CMakeLists.txt
index 10745c5..e33e5ee 100644 (file)
--- a/backend/src/llvm/CMakeLists.txt
+++ b/backend/src/llvm/CMakeLists.txt
@@ -3,6 +3,7 @@ include (${LLVM_DIR}/AddLLVMDefinitions.cmake)
 include (${LLVM_DIR}/HandleLLVMOptions.cmake)
 include (${LLVM_DIR}/LLVMProcessSources.cmake)
 include_directories(${LLVM_INCLUDE_DIRS})
+include_directories(../)
 add_llvm_target(GenBackend GenBackend.cpp)
 
 add_llvm_library_dependencies(LLVMGenBackend

diff --git a/backend/src/llvm/GenBackend.cpp b/backend/src/llvm/GenBackend.cpp
index df47b5b..daebb47 100644 (file)
--- a/backend/src/llvm/GenBackend.cpp
+++ b/backend/src/llvm/GenBackend.cpp
@@ -71,11 +71,10 @@
 #include "llvm/Support/TargetRegistry.h"
 #include "llvm/Support/Host.h"
 #include "llvm/Config/config.h"
+
+#include "ir/context.hpp"
 #include <algorithm>
-// Some ms header decided to define setjmp as _setjmp, undo this for this file.
-#ifdef _MSC_VER
-#undef setjmp
-#endif
+
 using namespace llvm;
 
 extern "C" void LLVMInitializeGenBackendTarget() {
@@ -126,7 +125,6 @@ namespace {
         OpaqueCounter(0), NextAnonValueNumber(0) {
       initializeLoopInfoPass(*PassRegistry::getPassRegistry());
       FPCounter = 0;
-      printf("DDDn\n\n");
     }
 
     virtual const char *getPassName() const { return "Gen backend"; }
@@ -3627,7 +3625,6 @@ bool GenTargetMachine::addPassesToEmitFile(PassManagerBase &PM,
                                            CodeGenOpt::Level OptLevel,
                                            bool DisableVerify) {
   if (FileType != TargetMachine::CGFT_AssemblyFile) return true;
-  
   PM.add(createGCLoweringPass());
   PM.add(createLowerInvokePass());
   PM.add(createCFGSimplificationPass());   // clean up after lower invoke.

diff --git a/backend/src/sys/platform.hpp b/backend/src/sys/platform.hpp
index f74f7fc..b887aca 100644 (file)
--- a/backend/src/sys/platform.hpp
+++ b/backend/src/sys/platform.hpp
@@ -187,11 +187,11 @@
 #endif
 
 /*! Debug syntactic sugar */
-#ifdef NDEBUG
+#if GBE_DEBUG
 #define IF_DEBUG(EXPR)
 #else
 #define IF_DEBUG(EXPR) EXPR
-#endif /* NDEBUG */
+#endif /* GBE_DEBUG */
 
 /*! Debug printing macros */
 #define STRING(x) #x
@@ -208,7 +208,7 @@
 #define _DO_JOIN2(X, Y) X##Y
 
 /*! Run-time assertion */
-#ifndef NDEBUG
+#if GBE_DEBUG
 #define GBE_ASSERT(EXPR) do {                                       \
   if (UNLIKELY(!(EXPR)))                                            \
     gbe::onFailedAssertion(#EXPR, __FILE__, __FUNCTION__, __LINE__);\
@@ -219,7 +219,7 @@
 } while (0)
 #else
 #define GBE_ASSERT(EXPR) do { } while (0)
-#define GBE_ASSERT_M(EXPR) do { } while (0)
+#define GBE_ASSERTM(EXPR, MSG) do { } while (0)
 #endif
 
 /*! Fatal error macros */