public static void ApplyKernel(GThread thread, int[] outputData)
        {
            //int[,] cache = thread.AllocateShared<int>("cache", X_SIZE, Y_SIZE);
            int targetX = thread.blockIdx.x;
            int targetY = 0;

            float value = 0;

            while(targetY < Y_SIZE)
            {
                for (int kernelX = KERNEL_SIZE / -2; kernelX <= KERNEL_SIZE / 2; kernelX++)
                    for (int kernelY = KERNEL_SIZE / -2; kernelY <= KERNEL_SIZE / 2; kernelY++)
                    {
                        int realX = targetX + kernelX;
                        int realY = targetY + kernelY;

                        if (realX >= 0 && realX < X_SIZE &&
                            realY >= 0 && realY < Y_SIZE)
                            value += MemoryKernel[kernelX + KERNEL_SIZE / 2, kernelY + KERNEL_SIZE / 2] * MemoryMain2D[realX, realY];

                        //Debug.WriteLine(String.Format("hoge: {0}",kernelX));
                    }

                //cache[targetX, targetY] = (int)value;
                //outputData[targetX + targetY * X_SIZE] = cache[targetX, targetY];
                outputData[targetX + targetY * X_SIZE] = (int)value;
                targetY++;
                value = 0;
            }
        }

 public static void calc_e(GThread thread, int n, int[] dx, int[] dy, int[] e)
 {
     for (int i = 0; i < n; i++)
     {
         e[i] = 2 * dy[i] - dx[i];
     }
 }

        public static void MuxArray(GThread thread, int[] a, int[] b, int[] c)
        {
            int tid = thread.blockIdx.x;

            if (tid < N)
                c[tid] = a[tid] * b[tid];
        }

        public static void Product(GThread thread, int[] a, int[] b, int[] c)
        {
            int tid = thread.threadIdx.x + thread.blockIdx.x * thread.blockDim.x;
            int[] cache = thread.AllocateShared<int>("cache", 4);
            int temp = 0;
            int cacheIndex=thread.threadIdx.x;
            while (tid < N)
            {
                temp = temp + a[tid] * b[tid];
                tid += thread.blockDim.x * thread.gridDim.x;
            }
            cache[thread.threadIdx.x] = temp;

            thread.SyncThreads();

            int i = thread.blockDim.x / 2;
            while (i != 0)
            {
                if (cacheIndex < i)
                {
                    cache[cacheIndex] += cache[cacheIndex + i];
                }
                thread.SyncThreads();

                i /= 2;
            }
            if (cacheIndex == 0)
            {
                c[thread.blockIdx.x] = cache[0];
            }
        }

        public static void thekernel(GThread thread, SphereOpenCL[] s, byte[] ptr)
        {
            int x = thread.threadIdx.x + thread.blockIdx.x * thread.blockDim.x;
            int y = thread.threadIdx.y + thread.blockIdx.y * thread.blockDim.y;
            int offset = x + y * thread.blockDim.x * thread.gridDim.x;
            float ox = (x - ray_gui.DIM / 2);
            float oy = (y - ray_gui.DIM / 2);

            float r = 0, g = 0, b = 0;
            float maxz = -INF;
            for (int i = 0; i < SPHERES; i++)
            {
                float n = 0;                
                float t = hit(s[i], ox, oy, ref n);
                if (t > maxz)
                {
                    float fscale = n;
                    r = s[i].r * fscale;
                    g = s[i].g * fscale;
                    b = s[i].b * fscale;
                    maxz = t;
                }
            }

            ptr[offset * 4 + 0] = (byte)(r * 255);
            ptr[offset * 4 + 1] = (byte)(g * 255);
            ptr[offset * 4 + 2] = (byte)(b * 255);
            ptr[offset * 4 + 3] = 255;
        }

        public static void Dot(GThread thread, float[] a, float[] b, float[] c)
        {
            float[] cache = thread.AllocateShared<float>("cache", threadsPerBlock);

            int tid = thread.threadIdx.x + thread.blockIdx.x * thread.blockDim.x;
            int cacheIndex = thread.threadIdx.x;

            float temp = 0;
            while (tid < N)
            {
                temp += a[tid] * b[tid];
                tid += thread.blockDim.x * thread.gridDim.x;
            }

            // set the cache values
            cache[cacheIndex] = temp;

            // synchronize threads in this block
            thread.SyncThreads();

            // for reductions, threadsPerBlock must be a power of 2
            // because of the following code
            int i = thread.blockDim.x / 2;
            while (i != 0)
            {
                if (cacheIndex < i)
                    cache[cacheIndex] += cache[cacheIndex + i];
                thread.SyncThreads();
                i /= 2;
            }

            if (cacheIndex == 0)
                c[thread.blockIdx.x] = cache[0];
        }

        public static void histo_kernel(GThread thread, byte[] buffer, int size, uint[] histo) 
        {
            // clear out the accumulation buffer called temp
            // since we are launched with 256 threads, it is easy
            // to clear that memory with one write per thread
            uint[] temp = thread.AllocateShared<uint>("temp", 256);
            temp[thread.threadIdx.x] = 0;
            thread.SyncThreads();

            // calculate the starting index and the offset to the next
            // block that each thread will be processing
            int i = thread.threadIdx.x + thread.blockIdx.x * thread.blockDim.x;
            int stride = thread.blockDim.x * thread.gridDim.x;
            while (i < size) 
            {
                thread.atomicAdd(ref temp[buffer[i]], 1 );
                i += stride;
            }
            // sync the data from the above writes to shared memory
            // then add the shared memory values to the values from
            // the other thread blocks using global memory
            // atomic adds
            // same as before, since we have 256 threads, updating the
            // global histogram is just one write per thread!
            thread.SyncThreads();

            thread.atomicAdd(ref (histo[thread.threadIdx.x]), temp[thread.threadIdx.x]);
        }

 public static void calc_e_v2(GThread thread, int n, int[] dx, int[] dy, int[] e)
 {
     int i = thread.blockDim.x * thread.blockIdx.x + thread.threadIdx.x;
     while(i < n)
     {
         e[i] = 2 * dy[i] - dx[i];
         i += (thread.blockDim.x * thread.gridDim.x);
     }
 }

        public static void SetValueGPUDouble(GThread thread, int n, double[] vector, double value)
        {
            int tid = thread.blockIdx.x;

            if (tid < n)
            {
                vector[tid] = value;
            }
        }

        public static void SetValueGPUSingle(GThread thread, int n, float[] vector, float value)
        {
            int tid = thread.blockIdx.x;

            if (tid < n)
            {
                vector[tid] = value;
            }
        }

 public static void add(GThread thread, int[] a, int[] b, int[] sum)
 {
     int index = thread.threadIdx.x + thread.blockIdx.x * thread.blockDim.x;
     while (index < N)
     {
         sum[index] = sum[index] + a[index] + b[index];
         index = index + thread.blockDim.x * thread.gridDim.x;
     }
 }

 public static void add_0(GThread thread, int[] a, int[] b, int[] c)
 {
     int tid = thread.blockIdx.x;
     while (tid < N)
     {
         c[tid] = a[tid] + b[tid];
         tid += thread.gridDim.x;
     }
 }

 public static void parentKernel(GThread thread, int[] a, int[] c, short coeff)
 {
     //childKernel(thread, a, c, coeff);
     int rc;
     //BROKEN thread.Launch(N / 2, numberYouFirstThoughtOf() * coeff, "childKernel", a, c, numberYouFirstThoughtOf() * coeff + 23 * a[0]);
     thread.Launch(N, 1, "childKernel", a, c, coeff * numberYouFirstThoughtOf());//a[0]);//numberYouFirstThoughtOf() * coeff + 23 * 
     rc = thread.SynchronizeDevice();
     int count = 0;
     rc = thread.GetDeviceCount(ref count);
 }

        public static void SyncThreadCountKernel(GThread thread, int[] input, int[] output)
        {
            var tid = thread.threadIdx.x;

            int value = input[tid];
            bool predicate = value == 1;
            var count = thread.SyncThreadsCount(predicate);

            if (tid == 0)
                output[0] = count;
        }

        public static void add(GThread thread, int[] a, int[] b, int[] c, int[] d, int[] e, int[] sum)
        {
            //To get Array Index for each Thread
            int index = thread.threadIdx.x + thread.blockIdx.x * thread.blockDim.x;

            while (index < N)
            {
                sum[index] = a[index] + b[index] + c[index] + d[index] + e[index];
                index = index + thread.blockDim.x * thread.gridDim.x;
            }
        }

        public static void thekernel(GThread thread, byte[] ptr)
        {
            int x = thread.blockIdx.x;
            int y = thread.blockIdx.y;
            int offset = x + y * thread.gridDim.x;

            int juliaValue = julia(x, y);
            ptr[offset * 4 + 0] = (byte)(255.0F * juliaValue);
            ptr[offset * 4 + 1] = 0;
            ptr[offset * 4 + 2] = 0;
            ptr[offset * 4 + 3] = 255;                         
        }

 public static void thekernel(GThread thread, int[] a, int[] b, int[] c)
 {
     int idx = thread.threadIdx.x + thread.blockIdx.x * thread.blockDim.x;
     if (idx < N) 
     {
         int idx1 = (idx + 1) % 256;
         int idx2 = (idx + 2) % 256;
         float aS = (a[idx] + a[idx1] + a[idx2]) / 3.0f;
         float bS = (b[idx] + b[idx1] + b[idx2]) / 3.0f;
         c[idx] = (int)(aS + bS) / 2;
     }
 }

        public static void BallotKernel(GThread thread, int[] input, int[] output)
        {
            var tid = thread.threadIdx.x;
            var wid = thread.threadIdx.x / 32;
            var twid = thread.threadIdx.x % 32;

            int value = input[tid];
            bool predicate = value == 1;
            var ballot = thread.Ballot(predicate);

            if (twid == 0)
                output[wid] = ballot;
        }

        public static void GPU_MA(GThread thread, int[,] GPU_A, int[,] GPU_B, int[,] GPU_C, int Size)
        {
            int x = thread.threadIdx.x + thread.blockDim.x * thread.blockIdx.x;
            int y = thread.threadIdx.y + thread.blockDim.y * thread.blockIdx.y;

            if (x < Size && y < Size)
            {
                //GPU_C[y, x] = y;
                GPU_C[y, x] = 0;
                for (int z = 0; z < Size; z++)
                {
                    GPU_C[y, x] += GPU_A[y, z] * GPU_B[z, x];
                }
            }
        }

        public static void DoSomeMath(GThread thread, int[] start, int[] end, double[] result)
        {
            int tid = thread.blockIdx.x;
            int i = 0;
            int tot = end[0] - start[0];
            while (tid < N)
            {
              while (i < tot)
              {
            result[i] = Math.Sin((i + 3.14));
            i++;
              }

              tid += thread.gridDim.x;
            }
        }