#ifndef OPENCL_H
#define OPENCL_H

#define __OPENCL__

#ifdef __CUDA__
#error Cannot include both opencl.h and cuda.h
#endif

typedef unsigned char uchar;
typedef unsigned short ushort;
typedef unsigned int uint;
typedef unsigned int size_t;

/* Vectors */
#include "opencl_vectors.h"


/* Images */
typedef struct image2d_s
{
} image2d_t;

typedef struct image3d_s
{
} image3d_t;

#define CLK_NORMALIZED_COORDS_TRUE  0
#define CLK_NORMALIZED_COORDS_FALSE 1
#define CLK_ADDRESS_MIRRORED_REPEAT 0
#define CLK_ADDRESS_REPEAT          2
#define CLK_ADDRESS_CLAMP_TO_EDGE   4
#define CLK_ADDRESS_CLAMP           6
#define CLK_ADDRESS_NONE            8
#define CLK_FILTER_NEAREST          0
#define CLK_FILTER_LINEAR           16

typedef unsigned int sampler_t;

float4 read_imagef(image2d_t image, sampler_t sampler, uint2 coord);
int4 read_imagei(image2d_t image, sampler_t sampler, uint2 coord);
uint4 read_imageui(image2d_t image, sampler_t sampler, uint2 coord);

void write_imagef(image2d_t image, uint2 coords, float4 color);
void write_imagei(image2d_t image, uint2 coords, int4 color);
void write_imageui(image2d_t image, uint2 coords, uint4 color);

int get_image_height (image2d_t image);
int get_image_width (image2d_t image);

#include "annotations.h"

/* Predefined constants */
#define M_PI_F 3.14f /* Standard says this should be value of PI.  Does not specify to what precision */


int mul24(int, int);

int get_group_id(int);
int get_local_id(int);
int get_local_size(int);
int get_global_id(int);
int get_num_groups(int);
int get_global_size(int);

float3 vload3(int, const __global float *);

/* Deprecated functions with explicit vector width */
float4 cbrt_4(float4);
int4 clamp_4(int4, int4, int4);
float4 clamp_4_f(float4, float4, float4);
float4 convert_float4_i(int4);
float4 convert_float4_ui(uint4);
float4 convert_float4_rtp_i(int4);
float4 convert_float4_rtp_ui(uint4);
float3 cross_3(float3, float3);
float4 cross_4(float4, float4);
float dot_2(float2, float2);
float dot_3(float3, float3);
float dot_4(float4, float4);
float fast_length_2(float2);
float fast_length_3(float3);
float fast_length_4(float4);
float3 fast_normalize_3(float3);
float4 fast_normalize_4(float4);
float4 half_exp_4(float4);
float4 half_powr_4(float4, float4);
float length_4(float4);
float2 max_2(float2, float2);
float3 max_3(float3, float3);
float4 max_4(float4, float4);
float2 min_2(float2, float2);
float3 min_3(float3, float3);
float4 min_4(float4, float4);

/* Overloaded functions */
#include "opencl_common_functions.h"
#include "opencl_geometric_functions.h"
#include "opencl_math_functions.h"
// TODO: Table 6.10 and 6.11 (Integer functions)
//       Table 6.14 (Scalar and Vector Relational functions)
//       Table 6.15 (Vector Data Load and Store Functions)
float4 __attribute__((overloadable)) convert_float4(int4);
float4 __attribute__((overloadable)) convert_float4(uint4);
float4 __attribute__((overloadable)) convert_float4(uchar4);

double4 __attribute__((overloadable)) convert_double4(int4);
double4 __attribute__((overloadable)) convert_double4(uint4);

int4 __attribute__((overloadable)) convert_int4(float4);
int4 __attribute__((overloadable)) convert_int4(uint4);
int4 __attribute__((overloadable)) convert_int4(uchar4);
int8 __attribute__((overloadable)) convert_int8(double8);

uchar4 __attribute__((overloadable)) convert_uchar4(int4);
uchar4 __attribute__((overloadable)) convert_uchar4(float4);
uchar4 __attribute__((overloadable)) convert_uchar4_sat(float4);

uint4 __attribute__((overloadable)) convert_uint4(float4);
uint4 __attribute__((overloadable)) convert_uint4(uchar4);
int __attribute__((overloadable)) clamp(int, int, int);
int2 __attribute__((overloadable)) clamp(int2, int2, int2);
int3 __attribute__((overloadable)) clamp(int3, int3, int3);
int4 __attribute__((overloadable)) clamp(int4, int4, int4);
float3 __attribute__((overloadable)) select(float3, float3, int3);
float4 __attribute__((overloadable)) select(float4, float4, int4);
double4 __attribute__((overloadable)) select(double4, double4, long4);

/*
typedef enum { CLK_LOCAL_MEM_FENCE } barrier_t;

void barrier(barrier_t);
*/

#define CLK_LOCAL_MEM_FENCE
#define CLK_GLOBAL_MEM_FENCE

void barrier();



/* Work group dimensions */

// Must define a dimension

#ifndef __1D_WORK_GROUP
#ifndef __2D_WORK_GROUP
#ifndef __3D_WORK_GROUP

#error You must specify the dimension of a work group by defining one of __1D_WORK_GROUP, __2D_WORK_GROUP or __3D_WORK_GROUP

#endif
#endif
#endif

// Must define only one dimension

#ifdef __1D_WORK_GROUP
#ifdef __2D_WORK_GROUP
#error Cannot define __1D_WORK_GROUP and __2D_WORK_GROUP
#endif
#ifdef __3D_WORK_GROUP
#error Cannot define __1D_WORK_GROUP and __3D_WORK_GROUP
#endif
#endif

#ifdef __2D_WORK_GROUP
#ifdef __1D_WORK_GROUP
#error Cannot define __2D_WORK_GROUP and __1D_WORK_GROUP
#endif
#ifdef __3D_WORK_GROUP
#error Cannot define __2D_WORK_GROUP and __3D_WORK_GROUP
#endif
#endif

#ifdef __3D_WORK_GROUP
#ifdef __1D_WORK_GROUP
#error Cannot define __3D_WORK_GROUP and __1D_WORK_GROUP
#endif
#ifdef __2D_WORK_GROUP
#error Cannot define __3D_WORK_GROUP and __2D_WORK_GROUP
#endif
#endif

// Generate axioms for different work group sizes

#ifdef __1D_WORK_GROUP
__axiom(get_local_size(1) == 1);
__axiom(get_local_size(2) == 1);
#endif

#ifdef __2D_WORK_GROUP
__axiom(get_local_size(2) == 1);
#endif


/* Work group grid dimensions */

// Must define a dimension

#ifndef __1D_GRID
#ifndef __2D_GRID
#ifndef __3D_GRID

#error You must specify the dimension of the grid of work groups by defining one of __1D_GRID, __2D_GRID or __3D_GRID

#endif
#endif
#endif

// Must define only one dimension

#ifdef __1D_GRID
#ifdef __2D_GRID
#error Cannot define __1D_GRID and __2D_GRID
#endif
#ifdef __3D_GRID
#error Cannot define __1D_GRID and __3D_GRID
#endif
#endif

#ifdef __2D_GRID
#ifdef __1D_GRID
#error Cannot define __2D_GRID and __1D_GRID
#endif
#ifdef __3D_GRID
#error Cannot define __2D_GRID and __3D_GRID
#endif
#endif

#ifdef __3D_GRID
#ifdef __1D_GRID
#error Cannot define __3D_GRID and __1D_GRID
#endif
#ifdef __2D_GRID
#error Cannot define __3D_GRID and __2D_GRID
#endif
#endif

// Generate axioms for different grid sizes

#ifdef __1D_GRID
__axiom(get_num_groups(1) == 1);
__axiom(get_num_groups(2) == 1);
#endif

#ifdef __2D_GRID
__axiom(get_num_groups(2) == 1);
#endif



#endif
