Обработка переноса блоков в сегментированной сумме в CUDA ⇐ C++

[Anonymous]

Я пытаюсь реализовать сегментированную сумму в CUDA, используя сокращение, без использования атомов. Я могу правильно рассчитать частичные суммы для каждого родителя в блоке. Однако, когда размер сегмента для конкретного родительского элемента превышает размер блока, я сталкиваюсь с проблемами памяти при использовании очень больших массивов.

Код: Выделить всё

__global__ void kernel_a(int* result, int* array, int* parents, int parents_length, int result_length, int* partial) {
extern __shared__ int shared[];

int idx = threadIdx.x;
int thread_id = blockIdx.x * blockDim.x + idx;

if (thread_id < parents_length) {
shared[idx] = array[thread_id];
} else {
shared[idx] = 0;
}
__syncthreads();

for (int stride = 1; stride < blockDim.x; stride *= 2) {
int val = 0;
if (idx >= stride && thread_id < parents_length && parents[thread_id] == parents[thread_id - stride]) {
val = shared[idx - stride];
}
shared[idx] += val;
__syncthreads();
}

if (thread_id < parents_length) {
int parent = parents[thread_id];
if (idx == blockDim.x - 1 || thread_id == parents_length - 1 || parents[thread_id] != parents[thread_id + 1]) {
partial[blockIdx.x * result_length + parent] = shared[idx];
}
}

}

// for handling block boundary carry over
__global__ void kernel_b(int* result, int* array, int* parents, int parents_length, int result_length, int* partial) {
int thread_id = blockIdx.x * blockDim.x + threadIdx.x;

if (thread_id < result_length) {
int sum = 0;
int blocks = (parents_length + blockDim.x - 1) / blockDim.x;
for (int i = 0; i < blocks; ++i) {
sum += partial[i * result_length + thread_id];
}
result[thread_id] = sum;
}

}

У меня есть этот подход, который работает правильно до определенного размера array_size, но после этого ему не хватает памяти. Теперь я понимаю, что проблема заключается в частичном_размере = outlength * ((parents_length + block_size - 1) //block_size. Как я могу получить доступ к частичной сумме для каждого родительского элемента в блоке, чтобы получить окончательную сегментированную сумму эффективно использовать память?


Подробнее здесь: https://stackoverflow.com/questions/786 ... um-in-cuda