diff --git a/amd/m01500_a3.cl b/amd/m01500_a3.cl
index cc907de0e..91e2ba6ca 100644
--- a/amd/m01500_a3.cl
+++ b/amd/m01500_a3.cl
@@ -837,7 +837,11 @@ static void m01500m (__local u32 *s_S, __global pw_t *pws, __global gpu_rule_t *
    * salt
    */
 
+  #ifdef DESCRYPT_SALT
+  const u32 salt = DESCRYPT_SALT;
+  #else
   const u32 salt = salt_bufs[salt_pos].salt_buf[0];
+  #endif
 
   /**
    * keys
@@ -1199,7 +1203,11 @@ static void m01500s (__local u32 *s_S, __global pw_t *pws, __global gpu_rule_t *
    * salt
    */
 
+  #ifdef DESCRYPT_SALT
+  const u32 salt = DESCRYPT_SALT;
+  #else
   const u32 salt = salt_bufs[salt_pos].salt_buf[0];
+  #endif
 
   /**
    * digest
diff --git a/include/ext_cuda.h b/include/ext_cuda.h
index c034e469c..18692deab 100644
--- a/include/ext_cuda.h
+++ b/include/ext_cuda.h
@@ -56,5 +56,9 @@ void hc_cuCtxSynchronize (void);
 void hc_cuCtxSetCacheConfig (CUfunc_cache config);
 void hc_cuDriverGetVersion (int *driverVersion);
 void hc_cuModuleLoadDataEx (CUmodule *module, const void *image, unsigned int numOptions, CUjit_option *options, void **optionValues);
+void hc_cuLinkAddFile (CUlinkState state, CUjitInputType type, const char *path, unsigned int numOptions, CUjit_option *options, void **optionValues);
+void hc_cuLinkComplete (CUlinkState state, void **cubinOut, size_t *sizeOut);
+void hc_cuLinkCreate (unsigned int numOptions, CUjit_option *options, void **optionValues, CUlinkState *stateOut);
+void hc_cuLinkDestroy (CUlinkState state);
 
 #endif
diff --git a/nv/m01500_a3.cu b/nv/m01500_a3.cu
index f28e09a59..86fe186f9 100644
--- a/nv/m01500_a3.cu
+++ b/nv/m01500_a3.cu
@@ -1204,7 +1204,11 @@ __device__ static void m01500m (const pw_t *pws, const gpu_rule_t *rules_buf, co
    * salt
    */
 
+  #ifdef DESCRYPT_SALT
+  const u32 salt = DESCRYPT_SALT;
+  #else
   const u32 salt = salt_bufs[salt_pos].salt_buf[0];
+  #endif
 
   /**
    * keys
@@ -1568,7 +1572,11 @@ __device__ static void m01500s (const pw_t *pws, const gpu_rule_t *rules_buf, co
    * salt
    */
 
+  #ifdef DESCRYPT_SALT
+  const u32 salt = DESCRYPT_SALT;
+  #else
   const u32 salt = salt_bufs[salt_pos].salt_buf[0];
+  #endif
 
   /**
    * digest
diff --git a/src/ext_cuda.c b/src/ext_cuda.c
index 77bab0c9c..74b4b1d5a 100644
--- a/src/ext_cuda.c
+++ b/src/ext_cuda.c
@@ -556,3 +556,51 @@ void hc_cuModuleLoadDataEx (CUmodule *module, const void *image, unsigned int nu
     exit (-1);
   }
 }
+
+void hc_cuLinkAddFile (CUlinkState state, CUjitInputType type, const char *path, unsigned int numOptions, CUjit_option *options, void **optionValues)
+{
+  CUresult CU_err = cuLinkAddFile (state, type, path, numOptions, options, optionValues);
+
+  if (CU_err != CUDA_SUCCESS)
+  {
+    log_error ("ERROR: %s %d\n", "cuLinkAddFile()", CU_err);
+
+    exit (-1);
+  }
+}
+
+void hc_cuLinkComplete (CUlinkState state, void **cubinOut, size_t *sizeOut)
+{
+  CUresult CU_err = cuLinkComplete (state, cubinOut, sizeOut);
+
+  if (CU_err != CUDA_SUCCESS)
+  {
+    log_error ("ERROR: %s %d\n", "cuLinkComplete()", CU_err);
+
+    exit (-1);
+  }
+}
+
+void hc_cuLinkCreate (unsigned int numOptions, CUjit_option *options, void **optionValues, CUlinkState *stateOut)
+{
+  CUresult CU_err = cuLinkCreate (numOptions, options, optionValues, stateOut);
+
+  if (CU_err != CUDA_SUCCESS)
+  {
+    log_error ("ERROR: %s %d\n", "cuLinkCreate()", CU_err);
+
+    exit (-1);
+  }
+}
+
+void hc_cuLinkDestroy (CUlinkState state)
+{
+  CUresult CU_err = cuLinkDestroy (state);
+
+  if (CU_err != CUDA_SUCCESS)
+  {
+    log_error ("ERROR: %s %d\n", "cuLinkDestroy()", CU_err);
+
+    exit (-1);
+  }
+}
diff --git a/src/oclHashcat.c b/src/oclHashcat.c
index 245b29aec..1856ebc36 100644
--- a/src/oclHashcat.c
+++ b/src/oclHashcat.c
@@ -315,7 +315,7 @@ hc_thread_mutex_t mux_display;
 
 hc_global_data_t data;
 
-const char *PROMPT = "[s]tatus [p]ause [r]esume [b]ypass [c]heckpoint stop [q]uit => ";
+const char *PROMPT = "[s]tatus [p]ause [r]esume [b]ypass [c]heckpoint [q]uit => ";
 
 const char *USAGE_MINI[] =
 {
@@ -1607,6 +1607,43 @@ static void status_benchmark ()
  * oclHashcat -only- functions
  */
 
+#ifdef _CUDA
+
+static void generate_source_kernel_filename (const uint attack_exec, const uint attack_kern, const uint kern_type, char *install_dir, char *kernel_file)
+{
+  if (attack_exec == ATTACK_EXEC_ON_GPU)
+  {
+    if (attack_kern == ATTACK_KERN_STRAIGHT)
+      snprintf (kernel_file, 255, "%s/nv/m%05d_a0.cu", install_dir, (int) kern_type);
+    else if (attack_kern == ATTACK_KERN_COMBI)
+      snprintf (kernel_file, 255, "%s/nv/m%05d_a1.cu", install_dir, (int) kern_type);
+    else if (attack_kern == ATTACK_KERN_BF)
+      snprintf (kernel_file, 255, "%s/nv/m%05d_a3.cu", install_dir, (int) kern_type);
+  }
+  else
+    snprintf (kernel_file, 255, "%s/nv/m%05d.cu", install_dir, (int) kern_type);
+}
+
+#elif _OCL
+
+static void generate_source_kernel_filename (const uint attack_exec, const uint attack_kern, const uint kern_type, char *install_dir, char *kernel_file)
+{
+  if (attack_exec == ATTACK_EXEC_ON_GPU)
+  {
+    if (attack_kern == ATTACK_KERN_STRAIGHT)
+      snprintf (kernel_file, 255, "%s/amd/m%05d_a0.cl", install_dir, (int) kern_type);
+    else if (attack_kern == ATTACK_KERN_COMBI)
+      snprintf (kernel_file, 255, "%s/amd/m%05d_a1.cl", install_dir, (int) kern_type);
+    else if (attack_kern == ATTACK_KERN_BF)
+      snprintf (kernel_file, 255, "%s/amd/m%05d_a3.cl", install_dir, (int) kern_type);
+  }
+  else
+    snprintf (kernel_file, 255, "%s/amd/m%05d.cl", install_dir, (int) kern_type);
+}
+
+#endif
+
+
 static uint convert_from_hex (char *line_buf, const uint line_len)
 {
   if (line_len & 1) return (line_len); // not in hex
@@ -12780,6 +12817,25 @@ int main (int argc, char **argv)
       }
     }
 
+    /**
+     * Some algorithm, like descrypt, can benefit from JIT compilation
+     */
+
+    uint force_jit_compilation = 0;
+
+    if (hash_mode == 8900)
+    {
+      force_jit_compilation = 8900;
+    }
+    else if (hash_mode == 9300)
+    {
+      force_jit_compilation = 8900;
+    }
+    else if (hash_mode == 1500 && attack_mode == ATTACK_MODE_BF && data.salts_cnt == 1)
+    {
+      force_jit_compilation = 1500;
+    }
+
     /**
      * generate bitmap tables
      */
@@ -13847,52 +13903,91 @@ int main (int argc, char **argv)
 
       hc_cuStreamCreate (&device_param->stream, 0);
 
+      /**
+       * In theory we'd need a real JIT solution as we have it with OpenCL, but CUDA does not provide such a feature, what a shame!
+       * There's NVRTC library which is able to compile sourcecode to PTX which we could use, but for some unknown reason this works only for 64 bit
+       * There's also the problem that the user needs to install the CUDA SDK to get this to work.
+       */
+
+      force_jit_compilation = 0;
+
       /**
        * module find
        */
 
+      struct stat st;
+
       char module_file[256];
 
       memset (module_file, 0, sizeof (module_file));
 
-      #ifdef __x86_64__
-      if (attack_exec == ATTACK_EXEC_ON_GPU)
+      #ifdef BINARY_KERNEL
+
+      if (force_jit_compilation == 0)
       {
-        if (attack_kern == ATTACK_KERN_STRAIGHT)
-          snprintf (module_file, sizeof (module_file) - 1, "%s/kernels/4318/m%05d_a0.sm_%d%d.64.cubin", install_dir, (int) kern_type, sm_major, sm_minor);
-        else if (attack_kern == ATTACK_KERN_COMBI)
-          snprintf (module_file, sizeof (module_file) - 1, "%s/kernels/4318/m%05d_a1.sm_%d%d.64.cubin", install_dir, (int) kern_type, sm_major, sm_minor);
-        else if (attack_kern == ATTACK_KERN_BF)
-          snprintf (module_file, sizeof (module_file) - 1, "%s/kernels/4318/m%05d_a3.sm_%d%d.64.cubin", install_dir, (int) kern_type, sm_major, sm_minor);
+        #ifdef __x86_64__
+        if (attack_exec == ATTACK_EXEC_ON_GPU)
+        {
+          if (attack_kern == ATTACK_KERN_STRAIGHT)
+            snprintf (module_file, sizeof (module_file) - 1, "%s/kernels/4318/m%05d_a0.sm_%d%d.64.cubin", install_dir, (int) kern_type, sm_major, sm_minor);
+          else if (attack_kern == ATTACK_KERN_COMBI)
+            snprintf (module_file, sizeof (module_file) - 1, "%s/kernels/4318/m%05d_a1.sm_%d%d.64.cubin", install_dir, (int) kern_type, sm_major, sm_minor);
+          else if (attack_kern == ATTACK_KERN_BF)
+            snprintf (module_file, sizeof (module_file) - 1, "%s/kernels/4318/m%05d_a3.sm_%d%d.64.cubin", install_dir, (int) kern_type, sm_major, sm_minor);
+        }
+        else
+        {
+          snprintf (module_file, sizeof (module_file) - 1, "%s/kernels/4318/m%05d.sm_%d%d.64.cubin", install_dir, (int) kern_type, sm_major, sm_minor);
+
+          if ((hash_mode == 8900) || (hash_mode == 9300))
+          {
+            snprintf (module_file, sizeof (module_file) - 1, "%s/kernels/4318/m%05d_%d_%d_%d_%d.sm_%d%d.64.cubin", install_dir, (int) kern_type, data.salts_buf[0].scrypt_N, data.salts_buf[0].scrypt_r, data.salts_buf[0].scrypt_p, data.salts_buf[0].scrypt_tmto, sm_major, sm_minor);
+          }
+        }
+
+        #else
+        if (attack_exec == ATTACK_EXEC_ON_GPU)
+        {
+          if (attack_kern == ATTACK_KERN_STRAIGHT)
+            snprintf (module_file, sizeof (module_file) - 1, "%s/kernels/4318/m%05d_a0.sm_%d%d.32.cubin", install_dir, (int) kern_type, sm_major, sm_minor);
+          else if (attack_kern == ATTACK_KERN_COMBI)
+            snprintf (module_file, sizeof (module_file) - 1, "%s/kernels/4318/m%05d_a1.sm_%d%d.32.cubin", install_dir, (int) kern_type, sm_major, sm_minor);
+          else if (attack_kern == ATTACK_KERN_BF)
+            snprintf (module_file, sizeof (module_file) - 1, "%s/kernels/4318/m%05d_a3.sm_%d%d.32.cubin", install_dir, (int) kern_type, sm_major, sm_minor);
+        }
+        else
+        {
+          snprintf (module_file, sizeof (module_file) - 1, "%s/kernels/4318/m%05d.sm_%d%d.32.cubin", install_dir, (int) kern_type, sm_major, sm_minor);
+
+          if ((hash_mode == 8900) || (hash_mode == 9300))
+          {
+            snprintf (module_file, sizeof (module_file) - 1, "%s/kernels/4318/m%05d_%d_%d_%d_%d.sm_%d%d.32.cubin", install_dir, (int) kern_type, data.salts_buf[0].scrypt_N, data.salts_buf[0].scrypt_r, data.salts_buf[0].scrypt_p, data.salts_buf[0].scrypt_tmto, sm_major, sm_minor);
+          }
+        }
+
+        #endif
       }
       else
       {
-        snprintf (module_file, sizeof (module_file) - 1, "%s/kernels/4318/m%05d.sm_%d%d.64.cubin", install_dir, (int) kern_type, sm_major, sm_minor);
+        generate_source_kernel_filename (attack_exec, attack_kern, kern_type, install_dir, module_file);
 
-        if ((hash_mode == 8900) || (hash_mode == 9300))
+        if (stat (module_file, &st) == -1)
         {
-          snprintf (module_file, sizeof (module_file) - 1, "%s/kernels/4318/m%05d_%d_%d_%d_%d.sm_%d%d.64.cubin", install_dir, (int) kern_type, data.salts_buf[0].scrypt_N, data.salts_buf[0].scrypt_r, data.salts_buf[0].scrypt_p, data.salts_buf[0].scrypt_tmto, sm_major, sm_minor);
+          log_error ("ERROR: %s: %s", module_file, strerror (errno));
+
+          return -1;
         }
       }
 
       #else
-      if (attack_exec == ATTACK_EXEC_ON_GPU)
-      {
-        if (attack_kern == ATTACK_KERN_STRAIGHT)
-          snprintf (module_file, sizeof (module_file) - 1, "%s/kernels/4318/m%05d_a0.sm_%d%d.32.cubin", install_dir, (int) kern_type, sm_major, sm_minor);
-        else if (attack_kern == ATTACK_KERN_COMBI)
-          snprintf (module_file, sizeof (module_file) - 1, "%s/kernels/4318/m%05d_a1.sm_%d%d.32.cubin", install_dir, (int) kern_type, sm_major, sm_minor);
-        else if (attack_kern == ATTACK_KERN_BF)
-          snprintf (module_file, sizeof (module_file) - 1, "%s/kernels/4318/m%05d_a3.sm_%d%d.32.cubin", install_dir, (int) kern_type, sm_major, sm_minor);
-      }
-      else
+
+      generate_source_kernel_filename (attack_exec, attack_kern, kern_type, install_dir, module_file);
+
+      if (stat (module_file, &st) == -1)
       {
-        snprintf (module_file, sizeof (module_file) - 1, "%s/kernels/4318/m%05d.sm_%d%d.32.cubin", install_dir, (int) kern_type, sm_major, sm_minor);
+        log_error ("ERROR: %s: %s", module_file, strerror (errno));
 
-        if ((hash_mode == 8900) || (hash_mode == 9300))
-        {
-          snprintf (module_file, sizeof (module_file) - 1, "%s/kernels/4318/m%05d_%d_%d_%d_%d.sm_%d%d.32.cubin", install_dir, (int) kern_type, data.salts_buf[0].scrypt_N, data.salts_buf[0].scrypt_r, data.salts_buf[0].scrypt_p, data.salts_buf[0].scrypt_tmto, sm_major, sm_minor);
-        }
+        return -1;
       }
 
       #endif
@@ -14598,87 +14693,92 @@ int main (int argc, char **argv)
       const unsigned char **kernel_sources = (const unsigned char **) mymalloc (sizeof (unsigned char *));
 
       #ifdef BINARY_KERNEL
-      if (attack_exec == ATTACK_EXEC_ON_GPU)
-      {
-        if (attack_kern == ATTACK_KERN_STRAIGHT)
-          snprintf (kernel_file, sizeof (kernel_file) - 1, "%s/kernels/4098/m%05d_a0.%s_%s_%s_%d.kernel", install_dir, (int) kern_type, device_name, device_version, driver_version, COMPTIME);
-        else if (attack_kern == ATTACK_KERN_COMBI)
-          snprintf (kernel_file, sizeof (kernel_file) - 1, "%s/kernels/4098/m%05d_a1.%s_%s_%s_%d.kernel", install_dir, (int) kern_type, device_name, device_version, driver_version, COMPTIME);
-        else if (attack_kern == ATTACK_KERN_BF)
-          snprintf (kernel_file, sizeof (kernel_file) - 1, "%s/kernels/4098/m%05d_a3.%s_%s_%s_%d.kernel", install_dir, (int) kern_type, device_name, device_version, driver_version, COMPTIME);
-      }
-      else
-      {
-        snprintf (kernel_file, sizeof (kernel_file) - 1, "%s/kernels/4098/m%05d.%s_%s_%s_%d.kernel", install_dir, (int) kern_type, device_name, device_version, driver_version, COMPTIME);
-
-        if ((hash_mode == 8900) || (hash_mode == 9300))
-        {
-          snprintf (kernel_file, sizeof (kernel_file) - 1, "%s/kernels/4098/m%05d_%d_%d_%d_%d.%s_%s_%s_%d.kernel", install_dir, (int) kern_type, data.salts_buf[0].scrypt_N, data.salts_buf[0].scrypt_r, data.salts_buf[0].scrypt_p, data.salts_buf[0].scrypt_tmto, device_name, device_version, driver_version, COMPTIME);
-        }
-      }
-
-      if (stat (kernel_file, &st) == -1)
+      if (force_jit_compilation == 0)
       {
-        if (quiet == 0) log_info ("Device #%u: Kernel %s not found in cache! Building may take a while...", device_id + 1, kernel_file);
-
-        char module_file[256];
-
-        memset (module_file, 0, sizeof (module_file));
-
         if (attack_exec == ATTACK_EXEC_ON_GPU)
         {
           if (attack_kern == ATTACK_KERN_STRAIGHT)
-            snprintf (module_file, sizeof (module_file) - 1, "%s/kernels/4098/m%05d_a0.VLIW%d.llvmir", install_dir, (int) kern_type, vliw);
+            snprintf (kernel_file, sizeof (kernel_file) - 1, "%s/kernels/4098/m%05d_a0.%s_%s_%s_%d.kernel", install_dir, (int) kern_type, device_name, device_version, driver_version, COMPTIME);
           else if (attack_kern == ATTACK_KERN_COMBI)
-            snprintf (module_file, sizeof (module_file) - 1, "%s/kernels/4098/m%05d_a1.VLIW%d.llvmir", install_dir, (int) kern_type, vliw);
+            snprintf (kernel_file, sizeof (kernel_file) - 1, "%s/kernels/4098/m%05d_a1.%s_%s_%s_%d.kernel", install_dir, (int) kern_type, device_name, device_version, driver_version, COMPTIME);
           else if (attack_kern == ATTACK_KERN_BF)
-            snprintf (module_file, sizeof (module_file) - 1, "%s/kernels/4098/m%05d_a3.VLIW%d.llvmir", install_dir, (int) kern_type, vliw);
+            snprintf (kernel_file, sizeof (kernel_file) - 1, "%s/kernels/4098/m%05d_a3.%s_%s_%s_%d.kernel", install_dir, (int) kern_type, device_name, device_version, driver_version, COMPTIME);
         }
         else
         {
-          snprintf (module_file, sizeof (module_file) - 1, "%s/kernels/4098/m%05d.VLIW%d.llvmir", install_dir, (int) kern_type, vliw);
+          snprintf (kernel_file, sizeof (kernel_file) - 1, "%s/kernels/4098/m%05d.%s_%s_%s_%d.kernel", install_dir, (int) kern_type, device_name, device_version, driver_version, COMPTIME);
 
           if ((hash_mode == 8900) || (hash_mode == 9300))
           {
-            snprintf (module_file, sizeof (module_file) - 1, "%s/kernels/4098/m%05d_%d_%d_%d_%d.VLIW%d.llvmir", install_dir, (int) kern_type, data.salts_buf[0].scrypt_N, data.salts_buf[0].scrypt_r, data.salts_buf[0].scrypt_p, data.salts_buf[0].scrypt_tmto, vliw);
+            snprintf (kernel_file, sizeof (kernel_file) - 1, "%s/kernels/4098/m%05d_%d_%d_%d_%d.%s_%s_%s_%d.kernel", install_dir, (int) kern_type, data.salts_buf[0].scrypt_N, data.salts_buf[0].scrypt_r, data.salts_buf[0].scrypt_p, data.salts_buf[0].scrypt_tmto, device_name, device_version, driver_version, COMPTIME);
           }
         }
 
-        load_kernel (module_file, 1, kernel_lengths, kernel_sources);
+        if (stat (kernel_file, &st) == -1)
+        {
+          if (quiet == 0) log_info ("Device #%u: Kernel %s not found in cache! Building may take a while...", device_id + 1, kernel_file);
 
-        cl_program program = hc_clCreateProgramWithBinary (device_param->context, 1, &device_param->device, kernel_lengths, (const unsigned char **) kernel_sources, NULL);
+          char module_file[256];
 
-        local_free (kernel_sources[0]);
+          memset (module_file, 0, sizeof (module_file));
 
-        hc_clBuildProgram (program, 1, &device_param->device, "-cl-std=CL1.2", NULL, NULL);
+          if (attack_exec == ATTACK_EXEC_ON_GPU)
+          {
+            if (attack_kern == ATTACK_KERN_STRAIGHT)
+              snprintf (module_file, sizeof (module_file) - 1, "%s/kernels/4098/m%05d_a0.VLIW%d.llvmir", install_dir, (int) kern_type, vliw);
+            else if (attack_kern == ATTACK_KERN_COMBI)
+              snprintf (module_file, sizeof (module_file) - 1, "%s/kernels/4098/m%05d_a1.VLIW%d.llvmir", install_dir, (int) kern_type, vliw);
+            else if (attack_kern == ATTACK_KERN_BF)
+              snprintf (module_file, sizeof (module_file) - 1, "%s/kernels/4098/m%05d_a3.VLIW%d.llvmir", install_dir, (int) kern_type, vliw);
+          }
+          else
+          {
+            snprintf (module_file, sizeof (module_file) - 1, "%s/kernels/4098/m%05d.VLIW%d.llvmir", install_dir, (int) kern_type, vliw);
 
-        size_t binary_size;
+            if ((hash_mode == 8900) || (hash_mode == 9300))
+            {
+              snprintf (module_file, sizeof (module_file) - 1, "%s/kernels/4098/m%05d_%d_%d_%d_%d.VLIW%d.llvmir", install_dir, (int) kern_type, data.salts_buf[0].scrypt_N, data.salts_buf[0].scrypt_r, data.salts_buf[0].scrypt_p, data.salts_buf[0].scrypt_tmto, vliw);
+            }
+          }
 
-        clGetProgramInfo (program, CL_PROGRAM_BINARY_SIZES, sizeof (size_t), &binary_size, NULL);
+          load_kernel (module_file, 1, kernel_lengths, kernel_sources);
 
-        unsigned char *binary = (unsigned char *) mymalloc (binary_size);
+          cl_program program = hc_clCreateProgramWithBinary (device_param->context, 1, &device_param->device, kernel_lengths, (const unsigned char **) kernel_sources, NULL);
 
-        clGetProgramInfo (program, CL_PROGRAM_BINARIES, sizeof (binary), &binary, NULL);
+          local_free (kernel_sources[0]);
 
-        writeProgramBin (kernel_file, binary, binary_size);
+          hc_clBuildProgram (program, 1, &device_param->device, "-cl-std=CL1.2", NULL, NULL);
 
-        local_free (binary);
+          size_t binary_size;
 
-        stat (kernel_file, &st); // to reload filesize
-      }
+          clGetProgramInfo (program, CL_PROGRAM_BINARY_SIZES, sizeof (size_t), &binary_size, NULL);
 
-      #else
-      if (attack_exec == ATTACK_EXEC_ON_GPU)
-      {
-        if (attack_kern == ATTACK_KERN_STRAIGHT)
-          snprintf (kernel_file, sizeof (kernel_file) - 1, "%s/amd/m%05d_a0.cl", install_dir, (int) kern_type);
-        else if (attack_kern == ATTACK_KERN_COMBI)
-          snprintf (kernel_file, sizeof (kernel_file) - 1, "%s/amd/m%05d_a1.cl", install_dir, (int) kern_type);
-        else if (attack_kern == ATTACK_KERN_BF)
-          snprintf (kernel_file, sizeof (kernel_file) - 1, "%s/amd/m%05d_a3.cl", install_dir, (int) kern_type);
+          unsigned char *binary = (unsigned char *) mymalloc (binary_size);
+
+          clGetProgramInfo (program, CL_PROGRAM_BINARIES, sizeof (binary), &binary, NULL);
+
+          writeProgramBin (kernel_file, binary, binary_size);
+
+          local_free (binary);
+
+          stat (kernel_file, &st); // to reload filesize
+        }
       }
       else
-        snprintf (kernel_file, sizeof (kernel_file) - 1, "%s/amd/m%05d.cl", install_dir, (int) kern_type);
+      {
+        generate_source_kernel_filename (attack_exec, attack_kern, kern_type, install_dir, kernel_file);
+
+        if (stat (kernel_file, &st) == -1)
+        {
+          log_error ("ERROR: %s: %s", kernel_file, strerror (errno));
+
+          return -1;
+        }
+      }
+
+      #else
+
+      generate_source_kernel_filename (attack_exec, attack_kern, kern_type, install_dir, kernel_file);
 
       if (stat (kernel_file, &st) == -1)
       {
@@ -14694,7 +14794,14 @@ int main (int argc, char **argv)
       if (quiet == 0) log_info ("Device #%u: Kernel %s (%ld bytes)", device_id + 1, kernel_file, st.st_size);
 
       #ifdef BINARY_KERNEL
-      device_param->program = hc_clCreateProgramWithBinary (device_param->context, 1, &device_param->device, kernel_lengths, (const unsigned char **) kernel_sources, NULL);
+      if (force_jit_compilation == 0)
+      {
+        device_param->program = hc_clCreateProgramWithBinary (device_param->context, 1, &device_param->device, kernel_lengths, (const unsigned char **) kernel_sources, NULL);
+      }
+      else
+      {
+        device_param->program = hc_clCreateProgramWithSource (device_param->context, 1, (const char **) kernel_sources, NULL);
+      }
       #else
       device_param->program = hc_clCreateProgramWithSource (device_param->context, 1, (const char **) kernel_sources, NULL);
       #endif
@@ -14892,15 +14999,31 @@ int main (int argc, char **argv)
 
       char *build_opts = NULL;
 
-      #ifndef BINARY_KERNEL
-      build_opts = (char *) mymalloc (256);
+      #ifdef BINARY_KERNEL
 
-      sprintf (build_opts, "-I . -I amd/ -D VLIW%d -x clc++ -cl-std=CL1.2", vliw);
+      if (force_jit_compilation == 0)
+      {
+        // nothing to do
+      }
+      else if (force_jit_compilation == 1500)
+      {
+        build_opts = (char *) mymalloc (256);
 
-      if ((hash_mode == 8900) || (hash_mode == 9300))
+        sprintf (build_opts, "-I . -I amd/ -D VLIW%d -x clc++ -cl-std=CL1.2 -DDESCRYPT_SALT=%d", vliw, data.salts_buf[0].salt_buf[0]);
+      }
+      else if (force_jit_compilation == 8900)
       {
+        build_opts = (char *) mymalloc (256);
+
         sprintf (build_opts, "-I . -I amd/ -D VLIW%d -x clc++ -cl-std=CL1.2 -DSCRYPT_N=%d -DSCRYPT_R=%d -DSCRYPT_P=%d -DSCRYPT_TMTO=%d", vliw, data.salts_buf[0].scrypt_N, data.salts_buf[0].scrypt_r, data.salts_buf[0].scrypt_p, data.salts_buf[0].scrypt_tmto);
       }
+
+      #else
+
+      build_opts = (char *) mymalloc (256);
+
+      sprintf (build_opts, "-I . -I amd/ -D VLIW%d -x clc++ -cl-std=CL1.2", vliw);
+
       #endif
 
       clBuildProgram (device_param->program, 1, &device_param->device, build_opts, NULL, NULL);
diff --git a/src/shared.c b/src/shared.c
index 2e7259cc6..e8eb2b25d 100644
--- a/src/shared.c
+++ b/src/shared.c
@@ -8376,7 +8376,7 @@ void stop_at_checkpoint ()
 
     data.checkpoint_cur_words = get_lowest_words_done ();
 
-    log_info ("Stop at next checkpoint");
+    log_info ("Checkpoint enabled, will quit when Restore Point updates next time");
   }
   else
   {