 
 void launch_kernel()
 {
+    Lock lock;
     profile("/ launch kernel");
     //train_model_kernel<<<1024, 256>>>();
     printf("| -- (cpu) cache_length: %lld\n", cache_length);
-    train_model_kernel<<<4, 8>>>(gpu_cache, cache_length);
-    fflush(stdout);
-    //pause();
+    train_model_kernel<<<1024, 256>>>(gpu_cache, cache_length, layer1_size,
+                                 gpu_vocab, gpu_expTable,
+                                 gpu_syn0, gpu_syn1,
+                                 &lock);
     profile("\\ launch kernel");
 }
 

    if ((i = ArgPos((char *)"-debug", argc, argv)) > 0) debug_mode = atoi(argv[i + 1]);
    if ((i = ArgPos((char *)"-binary", argc, argv)) > 0) binary = atoi(argv[i + 1]);
    if ((i = ArgPos((char *)"-cbow", argc, argv)) > 0) cbow = atoi(argv[i + 1]);
    if (cbow) alpha = 0.05;
    if ((i = ArgPos((char *)"-alpha", argc, argv)) > 0) alpha = atof(argv[i + 1]);
    if ((i = ArgPos((char *)"-output", argc, argv)) > 0) strcpy(output_file, argv[i + 1]);
    if ((i = ArgPos((char *)"-window", argc, argv)) > 0) window = atoi(argv[i + 1]);
    if ((i = ArgPos((char *)"-sample", argc, argv)) > 0) sample = atof(argv[i + 1]);
    if ((i = ArgPos((char *)"-hs", argc, argv)) > 0) hs = atoi(argv[i + 1]);
    if ((i = ArgPos((char *)"-negative", argc, argv)) > 0) negative = atoi(argv[i + 1]);
    if ((i = ArgPos((char *)"-threads", argc, argv)) > 0) num_threads = atoi(argv[i + 1]);
    if ((i = ArgPos((char *)"-iter", argc, argv)) > 0) iter = atoi(argv[i + 1]);
    if ((i = ArgPos((char *)"-min-count", argc, argv)) > 0) min_count = atoi(argv[i + 1]);
    if ((i = ArgPos((char *)"-classes", argc, argv)) > 0) classes = atoi(argv[i + 1]);
    if ((i = ArgPos((char *)"-cache", argc, argv)) > 0) cache_size = atoi(argv[i + 1]);
    vocab = (struct vocab_word *)calloc(vocab_max_size, sizeof(struct vocab_word));
    vocab_hash = (int *)calloc(vocab_hash_size, sizeof(int));
    expTable = (real *)malloc((EXP_TABLE_SIZE + 1) * sizeof(real));\

    /* profile: computing exp_table */
    profile("/ computing exp_table");


struct vocab_word {
    long long cn;
    int point[MAX_CODE_LENGTH];
    char *word, code[MAX_CODE_LENGTH], codelen;
};

struct Lock {
    int *mutex;
    Lock(void) {
        int state = 0;
        cudaMalloc((void**) &mutex, sizeof(int));
        cudaMemcpy(mutex, &state, sizeof(int), cudaMemcpyHostToDevice);
    }
    ~Lock(void) {
        cudaFree(mutex);
    }
    __device__ void lock(void) {
        while(atomicCAS(mutex, 0, 1) != 0);
    }
    __device__ void unlock(void) {
        atomicExch(mutex, 0);
    }
};

char train_file[MAX_STRING], output_file[MAX_STRING];
char save_vocab_file[MAX_STRING], read_vocab_file[MAX_STRING];
struct vocab_word *vocab, *gpu_vocab;
real *gpu_expTable, *gpu_syn0, *gpu_syn1;
// int binary = 0, cbow = 1, debug_mode = 2, window = 5, min_count = 5, num_threads = 12, min_reduce = 1, cache_size = 1;
int binary = 0, cbow = 1, debug_mode = 2, min_count = 5, num_threads = 12, min_reduce = 1, cache_size = 1;
int *vocab_hash;
long long vocab_max_size = 1000, vocab_size = 0, layer1_size = 100, cache_length = 0, *cpu_cache, *gpu_cache;
//long long train_words = 0, word_count_actual = 0, iter = 5, file_size = 0, classes = 0;
long long train_words = 0, file_size = 0, classes = 0;
//real alpha = 0.025, starting_alpha, sample = 1e-3;
real *syn0, *syn1, *syn1neg, *expTable;
clock_t start;

int hs = 0, negative = 5;
const int table_size = 1e8;
int *table;

__device__ long long word_count_actual = 0;
//__device__ real alpha = 0.25;
__constant__ long long gpu_train_words = 0, iter = 5;
__constant__ real starting_alpha, sample = 1e-3;
__constant__ int window;

void paulse(void);
void profile(const char *);
