/*
 * cu_syntax.h
 *
 *  Created on: 21/09/2019
 *      Author: pedro
 */

#ifndef SRC_UTILS_CU_SYNTAX_H_
#define SRC_UTILS_CU_SYNTAX_H_

// To avoid cuda syntax errors
#if CUDA_VERSION
#define __global
#define __kernel extern "C" __global__
#define __local
#define __constant
#define atomic_add atomicAdd
#define atomic_inc(a) atomicAdd(a,1)
#define get_local_id(a) threadIdx.x
#define get_local_size(a) blockDim.x
#define get_global_id(a) blockIdx.x * blockDim.x + threadIdx.x
#define get_global_size(a) gridDim.x * blockDim.x
#define atomic_cmpxchg atomicCAS
#define atomic_xchg atomicExch
#define atomic_dec(a) atomicSub(a,1)
#define atomic_and atomicAnd
#define CUDA_FUNC __device__
#define convert_ushort (ushort)
#define convert_int (int)
#define convert_uint (uint)
#define convert_ulong (ulong)
#define barrier(a) __syncthreads()
#define CLK_GLOBAL_MEM_FENCE
#define CLK_GLOBAL_MEM_FENCE

#if CL_D_MAX < 32
#define popcount __popc
#define clz __clz
#else
#define popcount __popcll
#define clz __clzll
#endif

#endif

#endif /* SRC_UTILS_CU_SYNTAX_H_ */