 
     baq_hmm.start();
 
-    if (baq_use_lmem)
-    {
+#if ENABLE_LMEM_PATH
+    if (!baq_use_lmem)
+#endif
+    {
+        // slow path: store the matrices in global memory
+        parallel<system>::for_each(thrust::make_zip_iterator(thrust::make_tuple(active_baq_read_list.begin(),
+                                                                      baq.matrix_index.begin(),
+                                                                      baq.scaling_index.begin())),
+                         thrust::make_zip_iterator(thrust::make_tuple(active_baq_read_list.end(),
+                                                                      baq.matrix_index.end(),
+                                                                      baq.scaling_index.end())),
+                         hmm_glocal<system>(context, batch.device, baq_state));
+    }
+#if ENABLE_LMEM_PATH
+    else {
         // fast path: use local memory for the matrices
         parallel<system>::for_each(thrust::make_zip_iterator(thrust::make_tuple(active_baq_read_list.begin(),
                                                                       baq.matrix_index.begin(),

                               context.active_read_list.end(),
                               compute_hmm_windows<system>(context, batch.device));

    if (!baq_use_lmem)
    {
        // compute the index and size of the HMM matrices
        baq.matrix_index.resize(num_active + 1);

    vector<system, uint32>& baq_state = context.temp_u32_2;

    // check if we can use the lmem kernel
#if ENABLE_LMEM_PATH
    const bool baq_use_lmem = (batch.host->max_read_size <= LMEM_MAX_READ_LEN);
    static bool baq_lmem_warning_printed = false;

    if (!baq_use_lmem && !baq_lmem_warning_printed)
    {
        fprintf(stderr, "WARNING: read size exceeds LMEM_MAX_READ_LEN, using fallback path for BAQ\n");
        baq_lmem_warning_printed = true;
    }
#endif

    timer<system> baq_setup, baq_hmm, baq_postprocess;

