#include <stdio.h>
#include "../context.h"

__code twice(struct Context* context, struct LoopCounter* loopCounter, int index, int prefix, int* array, struct Context* workerContext) {
    int i = loopCounter->i;
    if (i < prefix) {
        array[i+index*prefix] = array[i+index*prefix]*2;
        loopCounter->i++;

        goto meta(context, C_twice);
    }

    loopCounter->i = 0;
    goto meta(workerContext, workerContext->next);
}

__code twice_stub(struct Context* context) {
    struct LoopCounter* loopCounter = &context->data[context->dataNum]->LoopCounter;
    struct Array* array = &context->data[context->dataNum+1]->Array;

    Worker *worker = context->worker;
    CUDAWorker* cudaWorker = (CUDAWorker*)worker->worker;
  // memory allocate
    CUdeviceptr devA;
    CUdeviceptr devOut[num_exec];

    checkCudaErrors(cuMemAlloc(&devA, array->size));

    //twiceカーネルが定義されてなければそれをロードする
    checkCudaErrors(cuModuleLoad(&context->module, "CUDAtwice.ptx"));
    checkCudaErrors(cuModuleGetFunction(context->&function, module, "twice"));

    //入力のDataGearをGPUにbuffer経由で送る
    // Synchronous data transfer(host to device)
    checkCudaErrors(cuMemcpyHtoD(devA, array->array, array->size));

  // Asynchronous launch kernel
     context->num_exec = 1;
     void* args[] = {&devA};
     checkCudaErrors(cuLaunchKernel(function,
                       array->prefix, 1, 1,
                       context->num_exec, 1, 1,
                                 0, NULL , args, NULL));

    //結果を取ってくるコマンドを入力する
    //コマンドの終了待ちを行う   
    checkCudaErrors(cuMemcpyDtoH(array->array, devA, array->size));
    // wait for stream

    //continuationにそってGPUworkerに戻る
    goto meta(context, context->next);
}