     thrust::fill(baq.qualities.begin(), baq.qualities.end(), uint8(-1));
 
     // initialize matrices and scaling factors
+#if ENABLE_LMEM_PATH
     if (!baq_use_lmem)
+#endif
     {
         thrust::fill_n(baq.forward.begin(), baq.forward.size(), 0.0);
         thrust::fill_n(baq.backward.begin(), baq.backward.size(), 0.0);

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

    // check if we can use the lmem kernel
#if 0
    // note: the lmem path is effectively broken due to the varying band width
    const bool baq_use_lmem = (batch.host->max_read_size <= LMEM_MAX_READ_LEN);
    static bool baq_lmem_warning_printed = false;
#else
    constexpr bool baq_use_lmem = false;
    static bool baq_lmem_warning_printed = true;
#endif

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

    timer<system> baq_setup, baq_hmm, baq_postprocess;


#define GAP_OPEN_PROBABILITY (pow(10.0, (-40.0)/10.))
#define GAP_EXTENSION_PROBABILITY 0.1

// note: the lmem path is effectively broken due to the varying band width
#define ENABLE_LMEM_PATH 0

#if ENABLE_LMEM_PATH
// maximum read size for the lmem kernel
#define LMEM_MAX_READ_LEN 151
#define LMEM_MAT_ROW_SIZE (3 * MIN_BAND_WIDTH2 + 6)
