 {
     Lock lock;
     profile("/ launch kernel");
-    //train_model_kernel<<<1024, 256>>>();
-    printf("| -- (cpu) cache_length: %lld\n", cache_length);
-    train_model_kernel<<<1024, 256>>>(gpu_cache, cache_length, layer1_size,
-                                 gpu_vocab, gpu_expTable,
-                                 gpu_syn0, gpu_syn1,
-                                 &lock);
+    
+    FILE *debug_in, *debug_out;
+
+    // debug in
+    debug_in = fopen("debug_gpu_in", "w");
+    for (long long i = 0; i < vocab_size; ++i) {
+        fprintf(debug_in, "syn0[%lld]:", i);
+        for (long long j = 0; j < layer1_size; ++j) {
+            fprintf(debug_in, " %f", syn0[i * layer1_size + j]);
+        }
+        fprintf(debug_in, "\n");
+    }
+    fclose(debug_in);
+    
+    train_model_kernel<<<1024, 512>>>(gpu_cache, cache_length, layer1_size,
+                                       gpu_vocab, gpu_expTable,
+                                       gpu_syn0, gpu_syn1,
+                                       &lock);
+
+    // debug out
+    debug_out = fopen("debug_gpu_out", "w");
+    for (long long i = 0; i < vocab_size; ++i) {
+        fprintf(debug_out, "syn0[%lld]:", i);
+        for (long long j = 0; j < layer1_size; ++j) {
+            fprintf(debug_out, " %f", syn0[i * layer1_size + j]);
+        }
+        fprintf(debug_out, "\n");
+    }
+    fclose(debug_out);
+    
     profile("\\ launch kernel");
 }
 

}

__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;

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;
