Removed unused returnStruct field from function Included stdlib.h for all files

diff --git a/backend/kernels/add.cl b/backend/kernels/add.cl
index 7786e87..e03781d 100644 (file)
--- a/backend/kernels/add.cl
+++ b/backend/kernels/add.cl
@@ -1,3 +1,4 @@
+#include <stdlib.h>
 __kernel unsigned int add(unsigned int x, unsigned int y)
 {
   return x + y;

diff --git a/backend/kernels/add2.cl b/backend/kernels/add2.cl
index 8070576..58a8d2f 100644 (file)
--- a/backend/kernels/add2.cl
+++ b/backend/kernels/add2.cl
@@ -1,3 +1,4 @@
+#include <stdlib.h>
 struct big{
   unsigned int a, b;
 };

diff --git a/backend/kernels/cmp.cl b/backend/kernels/cmp.cl
index aab1615..789f852 100644 (file)
--- a/backend/kernels/cmp.cl
+++ b/backend/kernels/cmp.cl
@@ -1,3 +1,4 @@
+#include <stdlib.h>
 __kernel void test_cmp(__global bool *dst, int x, int y, float z, float w)
 {
   dst[0] = (x < y) + (z > w);

diff --git a/backend/kernels/cmp_cvt.cl b/backend/kernels/cmp_cvt.cl
index 266eae7..ab39fba 100644 (file)
--- a/backend/kernels/cmp_cvt.cl
+++ b/backend/kernels/cmp_cvt.cl
@@ -1,4 +1,4 @@
-#include "stdlib.h"
+#include <stdlib.h>
 
 __kernel void cmp_cvt(__global int *dst, int x, int y)
 {

diff --git a/backend/kernels/cycle.cl b/backend/kernels/cycle.cl
index fe9135c..8b0be54 100644 (file)
--- a/backend/kernels/cycle.cl
+++ b/backend/kernels/cycle.cl
@@ -1,3 +1,4 @@
+#include <stdlib.h>
 __kernel void cycle(global int *dst)
 {
   int x, y;

diff --git a/backend/kernels/extract.cl b/backend/kernels/extract.cl
index ca2ef19..fb8a5c5 100644 (file)
--- a/backend/kernels/extract.cl
+++ b/backend/kernels/extract.cl
@@ -1,4 +1,4 @@
-#include "stdlib.h"
+#include <stdlib.h>
 __kernel void extract(__global int4 *dst, __global int4 *src, int c)
 {
   const int4 from = src[0];

diff --git a/backend/kernels/function_param.cl b/backend/kernels/function_param.cl
index 4c29d21..46a7fd0 100644 (file)
--- a/backend/kernels/function_param.cl
+++ b/backend/kernels/function_param.cl
@@ -1,10 +1,10 @@
+#include <stdlib.h>
 struct struct0
 {
   int hop[5];
   int x, y, z;
 };
 
-#define __global __attribute__((address_space(1)))
 __kernel void param(__global struct struct0 *dst, struct struct0 s, __local int *h, int x, int y)
 {
   s.hop[4] += x + h[4];

diff --git a/backend/kernels/get_global_id.cl b/backend/kernels/get_global_id.cl
index 010beed..6c2b554 100644 (file)
--- a/backend/kernels/get_global_id.cl
+++ b/backend/kernels/get_global_id.cl
@@ -1,4 +1,4 @@
-#include "stdlib.h"
+#include <stdlib.h>
 
 __kernel void test_global_id(__global int *dst, __global int *p)
 {

diff --git a/backend/kernels/insert.cl b/backend/kernels/insert.cl
index 1711fea..6497c8c 100644 (file)
--- a/backend/kernels/insert.cl
+++ b/backend/kernels/insert.cl
@@ -1,4 +1,5 @@
-#include "stdlib.h"
+#include <stdlib.h>
+
 __kernel void insert(__global int4 *dst, __global int4 *src, int c)
 {
   int4 x = src[0];

diff --git a/backend/kernels/loop.cl b/backend/kernels/loop.cl
index e8ab630..af64abd 100644 (file)
--- a/backend/kernels/loop.cl
+++ b/backend/kernels/loop.cl
@@ -1,3 +1,4 @@
+#include <stdlib.h>
 __kernel void add(__global int *dst, unsigned int x)
 {
   for (int i = 0; i < x; ++i) dst[i]++;

diff --git a/backend/kernels/mad.cl b/backend/kernels/mad.cl
index 9589ff6..14c5987 100644 (file)
--- a/backend/kernels/mad.cl
+++ b/backend/kernels/mad.cl
@@ -1,4 +1,4 @@
-typedef float float4 __attribute__((ext_vector_type(4)));
+#include <stdlib.h>
 __attribute__((pure, overloadable)) int mad(int,int,int);
 __attribute__((pure, overloadable)) float mad(float,float,float);
 __attribute__((pure, overloadable)) float4 mad(float4,float4,float4);

diff --git a/backend/kernels/select.cl b/backend/kernels/select.cl
index 86a0b08..8f35915 100644 (file)
--- a/backend/kernels/select.cl
+++ b/backend/kernels/select.cl
@@ -1,4 +1,4 @@
-#include "stdlib.h"
+#include <stdlib.h>
 
 __kernel void test_select(__global int4 *dst,
                           __global int4 *src0,

diff --git a/backend/kernels/short.cl b/backend/kernels/short.cl
index 83e38a0..e2d5b19 100644 (file)
--- a/backend/kernels/short.cl
+++ b/backend/kernels/short.cl
@@ -1,3 +1,4 @@
+#include <stdlib.h>
 __kernel void short_write(__global short *dst, short x, short y)
 {
   dst[0] = x + y;

diff --git a/backend/kernels/shuffle.cl b/backend/kernels/shuffle.cl
index 6d49621..b39ae85 100644 (file)
--- a/backend/kernels/shuffle.cl
+++ b/backend/kernels/shuffle.cl
@@ -1,4 +1,4 @@
-#include "stdlib.h"
+#include <stdlib.h>
 __kernel void shuffle(__global int4 *dst, __global int4 *src, int c)
 {
   const int4 from = src[0];

diff --git a/backend/kernels/simple_float4.cl b/backend/kernels/simple_float4.cl
index 7b47a18..d9dbe51 100644 (file)
--- a/backend/kernels/simple_float4.cl
+++ b/backend/kernels/simple_float4.cl
@@ -1,4 +1,4 @@
-#include "stdlib.h"
+#include <stdlib.h>
 
 __kernel void simple_float4(__global float4 *dst, __global float4 *src)
 {

diff --git a/backend/kernels/simple_float4_2.cl b/backend/kernels/simple_float4_2.cl
index 1776d7c..6788511 100644 (file)
--- a/backend/kernels/simple_float4_2.cl
+++ b/backend/kernels/simple_float4_2.cl
@@ -1,4 +1,4 @@
-#include "stdlib.h"
+#include <stdlib.h>
 
 __kernel void simple_float4(__global float4 *dst, __global float4 *src)
 {

diff --git a/backend/kernels/simple_float4_3.cl b/backend/kernels/simple_float4_3.cl
index 57b84f2..3d20f42 100644 (file)
--- a/backend/kernels/simple_float4_3.cl
+++ b/backend/kernels/simple_float4_3.cl
@@ -1,4 +1,4 @@
-#include "stdlib.h"
+#include <stdlib.h>
 
 __kernel void simple_float4(__global float4 *dst, __global float4 *src, bool b)
 {

diff --git a/backend/kernels/stdlib.h b/backend/kernels/stdlib.h
index 49ed4f0..8ee4b83 100644 (file)
--- a/backend/kernels/stdlib.h
+++ b/backend/kernels/stdlib.h
@@ -77,3 +77,6 @@ __attribute__((overloadable)) inline int4 select(int4 src0, int4 src1, int4 cond
   return dst;
 }
 
+#define __global __attribute__((address_space(1)))
+#define global __global
+

diff --git a/backend/kernels/store.cl b/backend/kernels/store.cl
index b23b13c..337ba02 100644 (file)
--- a/backend/kernels/store.cl
+++ b/backend/kernels/store.cl
@@ -1,3 +1,4 @@
+#include <stdlib.h>
 __kernel void store(__global int *dst, __local int *dst0, int x)
 {
   dst[0] = 1;

diff --git a/backend/kernels/struct.cl b/backend/kernels/struct.cl
index a4fa09c..d72de6e 100644 (file)
--- a/backend/kernels/struct.cl
+++ b/backend/kernels/struct.cl
@@ -1,3 +1,4 @@
+#include <stdlib.h>
 struct my_struct {
   int a;
   int b[2];

diff --git a/backend/kernels/struct2.cl b/backend/kernels/struct2.cl
index 5c3aa41..adaace3 100644 (file)
--- a/backend/kernels/struct2.cl
+++ b/backend/kernels/struct2.cl
@@ -1,3 +1,4 @@
+#include <stdlib.h>
 struct my_struct {
   int a;
   int b[2];

diff --git a/backend/kernels/test_select.cl b/backend/kernels/test_select.cl
index 7973d23..8676c0d 100644 (file)
--- a/backend/kernels/test_select.cl
+++ b/backend/kernels/test_select.cl
@@ -1,4 +1,4 @@
-#include "stdlib.h"
+#include <stdlib.h>
 __kernel void test_select(__global int *dst, __global int *src)
 {
 

diff --git a/backend/kernels/undefined.cl b/backend/kernels/undefined.cl
index f9153ff..a1df672 100644 (file)
--- a/backend/kernels/undefined.cl
+++ b/backend/kernels/undefined.cl
@@ -1,3 +1,4 @@
+#include <stdlib.h>
 __kernel void undefined(__global int *dst)
 {
   int x;

diff --git a/backend/kernels/void.cl b/backend/kernels/void.cl
index ef59500..a5a5331 100644 (file)
--- a/backend/kernels/void.cl
+++ b/backend/kernels/void.cl
@@ -1,2 +1,3 @@
+#include <stdlib.h>
 __kernel void hop() {}
 

diff --git a/backend/src/ir/context.hpp b/backend/src/ir/context.hpp
index cc8545c..23bc048 100644 (file)
--- a/backend/src/ir/context.hpp
+++ b/backend/src/ir/context.hpp
@@ -159,7 +159,7 @@ namespace ir {
       BasicBlock *bb;              //!< Basic block currently processed
       vector<uint8_t> *usedLabels; //!< Store all labels that are defined
     };
-    vector<StackElem> fnStack;  //!< Stack of functions still to finish
+    vector<StackElem> fnStack;     //!< Stack of functions still to finish
     GBE_CLASS(Context);
   };
 

diff --git a/backend/src/ir/function.cpp b/backend/src/ir/function.cpp
index 06c2697..e2f06da 100644 (file)
--- a/backend/src/ir/function.cpp
+++ b/backend/src/ir/function.cpp
@@ -28,7 +28,7 @@ namespace gbe {
 namespace ir {
 
   Function::Function(const std::string &name, Profile profile) :
-    name(name), structReturned(false), profile(profile)
+    name(name), profile(profile)
   {
       initProfile(*this);
   }
@@ -68,7 +68,6 @@ namespace ir {
   std::ostream &operator<< (std::ostream &out, const Function &fn)
   {
     out << ".decl_function " << fn.getName() << std::endl;
-    out << ".return_struct " << (fn.isStructReturned() ? "true" : "false") << std::endl;
     out << fn.getRegisterFile();
     out << "## " << fn.inputNum() << " input register"
         << plural(fn.inputNum())  << " ##" << std::endl;

diff --git a/backend/src/ir/function.hpp b/backend/src/ir/function.hpp
index ae1fc43..c4aaf04 100644 (file)
--- a/backend/src/ir/function.hpp
+++ b/backend/src/ir/function.hpp
@@ -143,10 +143,6 @@ namespace ir {
       GBE_ASSERT(blocks[ID] != NULL);
       return *blocks[ID];
     }
-    /*! Function returns a structure by pointer (see ptx32 ABI) */
-    INLINE void setStructReturned(bool isReturned) { structReturned = isReturned; }
-    /*! Indicate if a structure is returned from the function */
-    INLINE bool isStructReturned(void) const { return structReturned; }
     /*! Create a new label (still not bound to a basic block) */
     LabelIndex newLabel(void);
     /*! Number of registers in the register file */
@@ -175,7 +171,6 @@ namespace ir {
     vector<BasicBlock*> blocks;   //!< All chained basic blocks
     RegisterFile file;            //!< RegisterDatas used by the instructions
     GrowingPool<Instruction> insnPool; //!< For fast instruction allocation
-    bool structReturned;               //!< First argument is pointer to struct
     Profile profile;                   //!< Current function profile
     GBE_CLASS(Function);
   };

diff --git a/backend/src/llvm/llvm_gen_backend.cpp b/backend/src/llvm/llvm_gen_backend.cpp
index 956722f..74948dc 100644 (file)
--- a/backend/src/llvm/llvm_gen_backend.cpp
+++ b/backend/src/llvm/llvm_gen_backend.cpp
@@ -595,16 +595,12 @@ namespace gbe
 
   void GenWriter::emitFunctionPrototype(Function &F)
   {
-    const bool returnStruct = F.hasStructRetAttr();
-
+    GBE_ASSERTM(F.hasStructRetAttr() == false,
+                "Returned value for kernel functions");
     // Loop over the arguments and output registers for them
     if (!F.arg_empty()) {
       Function::arg_iterator I = F.arg_begin(), E = F.arg_end();
 
-      // When a struct is returned, first argument is pointer to the structure
-      if (returnStruct)
-        ctx.getFunction().setStructReturned(true);
-
       // Insert a new register for each function argument
       for (; I != E; ++I) {
         const Type *type = I->getType();
@@ -616,14 +612,9 @@ namespace gbe
 
     // When returning a structure, first input register is the pointer to the
     // structure
-    if (!returnStruct) {
-      const Type *type = F.getReturnType();
-      if (type->isVoidTy() == false) {
-        const ir::RegisterFamily family = getFamily(ctx, type);
-        const ir::Register reg = ctx.reg(family);
-        ctx.output(reg);
-      }
-    }
+    const Type *type = F.getReturnType();
+    GBE_ASSERTM(type->isVoidTy() == true,
+                "Returned value for kernel functions");
 
 #if GBE_DEBUG
     // Variable number of arguments is not supported
@@ -643,6 +634,14 @@ namespace gbe
 
   void GenWriter::emitFunction(Function &F)
   {
+    switch (F.getCallingConv()) {
+      case CallingConv::PTX_Device: // we do not emit device function
+        return;
+      case CallingConv::PTX_Kernel:
+        break;
+      default: GBE_ASSERTM(false, "Unsupported calling convention");
+    }
+
     ctx.startFunction(F.getName());
     this->regTranslator.clear();
     this->labelMap.clear();
@@ -722,7 +721,7 @@ namespace gbe
         case Instruction::LShr: ctx.SHR(type, dst, src0, src1); break;
         case Instruction::AShr: ctx.ASR(type, dst, src0, src1); break;
         default: NOT_SUPPORTED;
-      };
+      }
     }
   }
 
@@ -776,7 +775,7 @@ namespace gbe
           case ICmpInst::ICMP_UGT: ctx.LE(unsignedType, dst, src0, src1); break;
           case ICmpInst::ICMP_SGT: ctx.LE(signedType, dst, src0, src1); break;
           default: NOT_SUPPORTED;
-        };
+        }
       }
       // Nothing special to do
       else {
@@ -792,7 +791,7 @@ namespace gbe
           case ICmpInst::ICMP_UGT: ctx.GT(unsignedType, dst, src0, src1); break;
           case ICmpInst::ICMP_SGT: ctx.GT(signedType, dst, src0, src1); break;
           default: NOT_SUPPORTED;
-        };
+        }
       }
     }
   }
@@ -829,7 +828,7 @@ namespace gbe
         case ICmpInst::FCMP_OGT:
         case ICmpInst::FCMP_UGT: ctx.GT(type, dst, src0, src1); break;
         default: NOT_SUPPORTED;
-      };
+      }
     }
   }
 
@@ -920,7 +919,7 @@ namespace gbe
       }
       break;
       default: NOT_SUPPORTED;
-    };
+    }
   }
 
   /*! Once again, it is a templated functor. No lambda */
@@ -1146,7 +1145,8 @@ namespace gbe
 
   static INLINE ir::MemorySpace addressSpaceLLVMToGen(unsigned llvmMemSpace) {
     switch (llvmMemSpace) {
-      case 0: return ir::MEM_GLOBAL;
+      case 0: return ir::MEM_PRIVATE;
+      case 1: return ir::MEM_GLOBAL;
       case 4: return ir::MEM_LOCAL;
     }
     GBE_ASSERT(false);

diff --git a/backend/src/utest/utest_llvm.cpp b/backend/src/utest/utest_llvm.cpp
index 9ecc145..76ed389 100644 (file)
--- a/backend/src/utest/utest_llvm.cpp
+++ b/backend/src/utest/utest_llvm.cpp
@@ -77,7 +77,8 @@ runTests:
   fclose(dummyKernel);
 
   //UTEST_EXPECT_SUCCESS(utestLLVM2Gen("function_param.ll"));
-  UTEST_EXPECT_SUCCESS(utestLLVM2Gen("loop.ll"));
+  //UTEST_EXPECT_SUCCESS(utestLLVM2Gen("loop.ll"));
+  UTEST_EXPECT_SUCCESS(utestLLVM2Gen("function.ll"));
   //UTEST_EXPECT_SUCCESS(utestLLVM2Gen("mad.ll"));
 #if 0
   UTEST_EXPECT_SUCCESS(utestLLVM2Gen("select.ll"));