Added alignment attribute to address space structures and modified pa…

…rallel localmemory initializer to work with apple driver (original version did crash the i5 cpu compiler) #20

TESTING.txt

@@ -28,9 +28,14 @@ T3: POCL driver and Intel processor
     Device: Intel(R) Core(TM) i7-2640M CPU @ 2.80GHz.
 
 T4: Apple driver and Intel processor
+    OS: Mac OSX 10.8.4 
+    Driver: Apple
+    Device: Apple/Intel(R) Core(TM) i5-3210M CPU @ 2.50GHz.
+
+T5: POCL driver and Intel processor
     OS: Mac OSX
-    Driver:
-    Device:
+    Driver: pocl-trunk rev.753 with optimized llvm 3.3
+    Device: Apple/Intel(R) Core(TM) i5-3210M CPU @ 2.50GHz.
 
 See the ISSUES.txt document about problems in different drivers.
 

src/WebCLConsumer.cpp

@@ -433,7 +433,12 @@ class MemoryAccessHandler {
       analyser.getPointerAceesses();
 
     std::map< unsigned, unsigned > maxAccess;
-
+    // add default 8 bit align
+    maxAccess[clang::LangAS::opencl_constant] = 8;
+    maxAccess[clang::LangAS::opencl_global] = 8;
+    maxAccess[clang::LangAS::opencl_local] = 8;
+    maxAccess[0] = 8;
+
     for (WebCLAnalyser::MemoryAccessMap::iterator i = pointerAccesses.begin();
          i != pointerAccesses.end(); ++i) {
 

src/WebCLTransformer.cpp

@@ -226,28 +226,28 @@ void WebCLTransformer::createAddressSpaceLimitsNullInitializer(
 }
 
 void WebCLTransformer::createAddressSpaceTypedef(
-    AddressSpaceInfo &as, const std::string &name)
+    AddressSpaceInfo &as, const std::string &name, const std::string &alignment)
 {
     if (!as.empty()) {
         modulePrologue_ << "typedef struct "
                         << addressSpaceInfoAsStruct(as)
-                        << " " << name << ";\n\n";
+                        << " __attribute__ ((aligned (" << alignment << "))) " << name << ";\n\n";
     }
 }
 
 void WebCLTransformer::createPrivateAddressSpaceTypedef(AddressSpaceInfo &as)
 {
-    createAddressSpaceTypedef(as, cfg_.privateRecordType_);
+    createAddressSpaceTypedef(as, cfg_.privateRecordType_, cfg_.getNameOfAlignMacro("private"));
 }
 
 void WebCLTransformer::createLocalAddressSpaceTypedef(AddressSpaceInfo &as)
 {
-    createAddressSpaceTypedef(as, cfg_.localRecordType_);
+    createAddressSpaceTypedef(as, cfg_.localRecordType_, cfg_.getNameOfAlignMacro("local"));
 }
 
 void WebCLTransformer::createConstantAddressSpaceTypedef(AddressSpaceInfo &as)
 {
-    createAddressSpaceTypedef(as, cfg_.constantRecordType_);
+    createAddressSpaceTypedef(as, cfg_.constantRecordType_, cfg_.getNameOfAlignMacro("constant"));
 }
 
 void WebCLTransformer::createAddressSpaceLimitsTypedef(
@@ -636,9 +636,20 @@ void WebCLTransformer::flushQueuedTransformations() {
 void WebCLTransformer::addMinimumRequiredContinuousAreaLimit(unsigned addressSpace,
                                                              unsigned minWidthInBits)
 {
-    modulePrologue_ << "#define " << cfg_.getNameOfSizeMacro(addressSpace) << " ("
-                    << "((" << minWidthInBits << " + (CHAR_BIT - 1)) / CHAR_BIT)"
-                    << ")\n";
+    preModulePrologue_ << "#define " << cfg_.getNameOfSizeMacro(addressSpace) << " ("
+                       << "((" << minWidthInBits << " + (CHAR_BIT - 1)) / CHAR_BIT)"
+                       << ")\n";
+
+    // get aligment rounded to next power of two
+    unsigned minAlignment = 1;
+    while (minWidthInBits != 0) {
+      minWidthInBits = minWidthInBits>>1;
+      minAlignment = minAlignment<<1;
+    }
+    minAlignment = minAlignment>>1;
+
+    preModulePrologue_ << "#define " << cfg_.getNameOfAlignMacro(addressSpace) << " "
+                       << "(" << minAlignment << "/CHAR_BIT)\n";
 }
 
 void WebCLTransformer::addAddressSpaceNull(std::ostream &out, unsigned addressSpace)
@@ -1091,6 +1102,7 @@ void WebCLTransformer::emitLimitMacros(std::ostream &out)
 
 void WebCLTransformer::emitPrologue(std::ostream &out)
 {
+    out << preModulePrologue_.str();
     out << modulePrologue_.str();
     emitGeneralCode(out);
     emitLimitMacros(out);

src/WebCLTransformer.hpp

@@ -50,7 +50,7 @@ class WebCLTransformer : public WebCLReporter
     /// Applies all AST transformations.
     bool rewrite();
 
-    void createAddressSpaceTypedef(AddressSpaceInfo &as, const std::string &name);
+    void createAddressSpaceTypedef(AddressSpaceInfo &as, const std::string &name, const std::string &alignment);
     void createPrivateAddressSpaceTypedef(AddressSpaceInfo &as);
     void createLocalAddressSpaceTypedef(AddressSpaceInfo &as);
     void createConstantAddressSpaceTypedef(AddressSpaceInfo &as);
@@ -141,6 +141,10 @@ class WebCLTransformer : public WebCLReporter
     RequiredMacroSet usedClampMacros_;
     FunctionPrologueMap kernelPrologues_;
     FunctionPrologueMap functionPrologues_;
+
+    // this is stream for things that needs to be absolutely at first in program even before typedefs
+    std::stringstream preModulePrologue_;
+    // stream for other initializations at start of module like typedefs and address space structures
     std::stringstream modulePrologue_;
 
     // set to keep track that we are not doing paramRelocationInitializations multiple times

src/WebCLTransformerConfiguration.cpp

@@ -107,11 +107,28 @@ const std::string WebCLTransformerConfiguration::getNameOfLimitCheckMacro(
     return result.str();
 }
 
+const std::string WebCLTransformerConfiguration::getNameOfSizeMacro(const std::string &asName) const
+{
+  const std::string name =
+  macroPrefix_ + "_ADDRESS_SPACE_" + asName + "_MIN";
+  return name;
+}
+
 const std::string WebCLTransformerConfiguration::getNameOfSizeMacro(unsigned addressSpaceNum) const
 {
-    const std::string name =
-        macroPrefix_ + "_ADDRESS_SPACE_" + getNameOfAddressSpace(addressSpaceNum) + "_MIN";
-    return name;
+  return getNameOfSizeMacro(getNameOfAddressSpace(addressSpaceNum));
+}
+
+const std::string WebCLTransformerConfiguration::getNameOfAlignMacro(const std::string &asName) const
+{
+  const std::string name =
+  macroPrefix_ + "_ADDRESS_SPACE_" + asName + "_ALIGNMENT";
+  return name;
+}
+
+const std::string WebCLTransformerConfiguration::getNameOfAlignMacro(unsigned addressSpaceNum) const
+{
+  return getNameOfAlignMacro(getNameOfAddressSpace(addressSpaceNum));
 }
 
 const std::string WebCLTransformerConfiguration::getNameOfLimitMacro() const

src/WebCLTransformerConfiguration.hpp

@@ -24,7 +24,10 @@ class WebCLTransformerConfiguration
     const std::string getNameOfAddressSpaceNullPtrRef(unsigned addressSpaceNum) const;
     const std::string getNameOfLimitCheckMacro(
         unsigned addressSpaceNum, int limitCount) const;
+    const std::string getNameOfSizeMacro(const std::string &asName) const;
     const std::string getNameOfSizeMacro(unsigned addressSpaceNum) const;
+    const std::string getNameOfAlignMacro(const std::string &asName) const;
+    const std::string getNameOfAlignMacro(unsigned addressSpaceNum) const;
     const std::string getNameOfLimitMacro() const;
     const std::string getNameOfAddressSpace(clang::QualType type) const;
     const std::string getNameOfType(clang::QualType type) const;

src/general.cl

@@ -36,7 +36,7 @@ void _wcl_local_range_init(__local _WclInitType *begin, __local _WclInitType *en
         get_local_id(2);
     const size_t num_elements = stop - start;
     const size_t item_elements = num_elements / xyz_items;
-
+/*
     __local uchar *item_begin = start + (item_index * item_elements);
     __local const uchar *item_end = item_begin + item_elements;
     for (__local uchar *item_i = item_begin; item_i < item_end; ++item_i)
@@ -46,6 +46,7 @@ void _wcl_local_range_init(__local _WclInitType *begin, __local _WclInitType *en
     __local uchar *item_final = start + (loop_elements + item_index);
     if (item_final < stop)
         *item_final = _WCL_FILLCHAR;
+*/
 }
 
 #endif // cl_khr_initialize_memory