Тяга /CUDA: Какой эффективный способ суммировать элементы по координатам подматрицы?

Question

2D матрица: thrust::device_vector<bool> (строка за строкой).
Координаты подматрицы:

struct {
    uint16_t left;
    uint16_t top;
    uint16_t right;
    uint16_t bottom;
};

Как наиболее эффективно найти сумму матричных элементов, определяемых координатами подматрицы? На данный момент я использую временный ключевой вектор, где все области подматрицы однозначно отмечены, а затем применяются thrust::reduce_by_key на матрицу и ключевую матрицу. Но мне не нравится такой подход. Заранее спасибо.