 
 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");
 }
 

#define MAX_SENTENCE_LENGTH 1000
#define MAX_CODE_LENGTH 40
//#define CACHE_SIZE_BY_BYTES (long long) cache_size * 1024 * 1024 * 1024
#define CACHE_SIZE_BY_BYTES (long long) cache_size * 10240
#define CACHE_SIZE_BY_UNITS CACHE_SIZE_BY_BYTES / sizeof(long long)
#define CUDA_CALL(x) {\
        const cudaError_t error_code = (x);\

    return -1;
}

__global__ void train_model_kernel(const long long *cache, const long long cache_length, const long long layer1_size,
                                   struct vocab_word *vocab, real *expTable, real *syn0, real *syn1,
                                   Lock *lock)
{
    const unsigned int thread_idx = (blockIdx.x * blockDim.x) + threadIdx.x;
    const long long cache_begin = get_fragment_begin(cache, cache_length, thread_idx);
    const long long cache_end = get_fragment_end(cache, cache_length, thread_idx);
    long long word, last_word, word_count = 0, last_word_count = 0;
    long long word_count_actual = 0; // overwrite __device__ global variable word_count_acutal
    long long cache_idx = cache_begin, local_iter = iter;
    long long sen[MAX_SENTENCE_LENGTH + 1], sentence_length = 0, sentence_position = 0;
    long long a, b, c, d, cw, l2;
    unsigned long long next_random = (long long) thread_idx;
    real f, g, alpha = 0.05; // overwrite __device__ global variable alpha
    //real *neu1 = (real *) malloc(layer1_size * sizeof(*neu1));
    //real *neu1e = (real *) malloc(layer1_size * sizeof(*neu1e));
    real neu1[200];
    real neu1e[200];

    while (1) {
        if (word_count - last_word_count > 1000) {
            //lock->lock();
            word_count_actual += word_count - last_word_count;
            last_word_count = word_count;
            alpha = starting_alpha * (1 - (long long) (word_count_actual * blockDim.x * gridDim.x) / (real) (iter * gpu_train_words + 1));
            if (alpha < starting_alpha * 0.0001) alpha = starting_alpha * 0.0001;
            printf("%dAlpha: %f  word_count_actual: %lld  blockDim.x: %ud  gridDim.x: %ud  iter: %lld  gpu_train_words: %lld", alpha, word_count_actual, blockDim.x, gridDim.x, iter, gpu_train_words);
            //lock->unlock();
        }
        
        if (sentence_length == 0) {
            while (1) {
                word = cache[cache_idx++];
                if (cache_idx >= cache_end) break;
                if (word == -1) continue;
                ++word_count;
                if (word == 0) break;
                if (sample > 0) {
                    real ran = (sqrt(vocab[word].cn / (sample * gpu_train_words)) + 1) * (sample * gpu_train_words) / vocab[word].cn;
                    next_random = next_random * (unsigned long long)25214903917 + 11;
                    if (ran < (next_random & 0xFFFF) / (real)65536) continue;
                }
                sen[sentence_length++] = word;
                if (sentence_length >= MAX_SENTENCE_LENGTH) break;
            }
            sentence_position = 0;
        }
    
        if (cache_idx >= cache_end) {
            //lock->lock();
            word_count_actual += word_count - last_word_count;
            //lock->unlock();
            --local_iter;
            if (local_iter == 0) break;
            word_count = 0;
            last_word_count = 0;
            sentence_length = 0;
            cache_idx = cache_begin;
            continue;
        }

        word = sen[sentence_position];
        if (word == -1) continue;
        for (c = 0; c < layer1_size; ++c) neu1[c] = 0;
        for (c = 0; c < layer1_size; ++c) neu1e[c] = 0;
        next_random = next_random * (unsigned long long) 25214903917 + 11;
        b = next_random % window;
        if (1) { // train cbow: if (cbow) {
            // in -> hidden            
            cw = 0;
            for (a = b; a < window * 2 + 1 - b; ++a) if (a != window) {
                    c = sentence_position - window + a;
                    if (c < 0) continue;
                    if (c >= sentence_length) continue;
                    last_word = sen[c];
                    if (last_word == -1) continue;
                    for (c = 0; c < layer1_size; c++) neu1[c] += syn0[c + last_word * layer1_size];
                    ++cw;
                }
            if (cw) {
                for (c = 0; c < layer1_size; c++) neu1[c] /= cw;
                if (1) for (d = 0; d < vocab[word].codelen; ++d) { // using hs: if (hs) { for ...
                        f = 0;
                        l2 = vocab[word].point[d] * layer1_size;
                        // Propagate hidden -> output
                        for (c = 0; c < layer1_size; c++) f += neu1[c] * syn1[c + l2];
                        if (f <= -MAX_EXP) continue;
                        else if (f >= MAX_EXP) continue;
                        else f = expTable[(int)((f + MAX_EXP) * (EXP_TABLE_SIZE / MAX_EXP / 2))];
                        // 'g' is the gradient multiplied by the learning rate
                        g = (1 - vocab[word].code[d] - f) * alpha;
                        // Propagate errors output -> hidden
                        for (c = 0; c < layer1_size; c++) neu1e[c] += g * syn1[c + l2];
                        // Learn weights hidden -> output
                        for (c = 0; c < layer1_size; c++) syn1[c + l2] += g * neu1[c];
                    }
                // NEGATIVE SAMPLING
                /*
                if (negative > 0) {
                    printf("negative not implemented\n");
                    assert(0);
                }
                */
                // hidden -> in
                for (a = b; a < window * 2 + 1 - b; a++) if (a != window) {
                        c = sentence_position - window + a;
                        if (c < 0) continue;
                        if (c >= sentence_length) continue;
                        last_word = sen[c];
                        if (last_word == -1) continue;
                        for (c = 0; c < layer1_size; c++) syn0[c + last_word * layer1_size] += neu1e[c];
                    }
            }
        } else { // train skip-gram
            printf("skip-gram not implemented\n");
            assert(0);
        }
    
        ++sentence_position;
        if (sentence_position >= sentence_length) {
            sentence_length = 0;
            continue;
        }
    }
}

__device__ long long get_fragment_begin(const long long *cache, const long long cache_length, const unsigned int thread_idx)
