// Copyright (C) 2013-2018 Altera Corporation, San Jose, California, USA. All rights reserved.
// Permission is hereby granted, free of charge, to any person obtaining a copy of this
// software and associated documentation files (the "Software"), to deal in the Software
// without restriction, including without limitation the rights to use, copy, modify, merge,
// publish, distribute, sublicense, and/or sell copies of the Software, and to permit persons to
// whom the Software is furnished to do so, subject to the following conditions:
// The above copyright notice and this permission notice shall be included in all copies or
// substantial portions of the Software.
// 
// THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND,
// EXPRESS OR IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES
// OF MERCHANTABILITY, FITNESS FOR A PARTICULAR PURPOSE AND
// NONINFRINGEMENT. IN NO EVENT SHALL THE AUTHORS OR COPYRIGHT
// HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER LIABILITY,
// WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
// FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR
// OTHER DEALINGS IN THE SOFTWARE.
// 
// This agreement shall be governed in all respects by the laws of the State of California and
// by the laws of the United States of America.

 // ACL kernel for adding two input vectors





#define size 1000*1000/4

#define DATA_SIZE 4

__kernel void
vector_add(
    __global const float * restrict x,
    __global const float * restrict y,
    __global       float * restrict z,
    const int pad
) {
    float xx[DATA_SIZE],
          yy[DATA_SIZE],
          zz[DATA_SIZE];


	float bufferUmid[30 * 1024 +  DATA_SIZE ]; 


    for(int l=0, i=0; l<size; l++, i=DATA_SIZE+i) {
        


	#pragma unroll
        for (int b = 0; b <  5 * 1024 ; b = b + DATA_SIZE) {

        #pragma unroll DATA_SIZE 
        for (int parallel = 0; parallel < DATA_SIZE; parallel=parallel+1){



	bufferUmid[ b + parallel] = bufferUmid[ b + parallel + DATA_SIZE];
	  
       		}
	}





	//barrier(CLK_LOCAL_MEM_FENCE);

	#pragma unroll DATA_SIZE
        for (int parallel = 0; parallel < DATA_SIZE; parallel = parallel + 1){

        
	bufferUmid[ 5 * 1024 + parallel] = x[0+i+parallel];

         }





	

	/*#pragma unroll
        for(int j=0; j<DATA_SIZE; j++) {
            xx[j] = x[i+j];
        }*/

        








	/*#pragma unroll
        for(int j=0; j<DATA_SIZE; j++) {
            zz[j] = xx[j];
        }*/
       
        #pragma unroll
        for(int j=0; j<DATA_SIZE; j++) {
            z[0+i+j] = bufferUmid[j];
        }
    }
}