Применение ядра на 3-х мерные данные

Доброго времени суток.
Вот уже месяц ломаю себе голову над организацией решения одной задачи. Имеются 3-х мерные данные (куб 32*32*32, значения типа: 0.00392157 0.00392157 0.00392157 0.00392157 0.00784314). Необходимо посчитать градиент по ним, используя ядро (Собеля в данном случае).
Ядро произвольного размера (генерируется на CPU). Т.е для того, что бы применить ядро нужно в каждой точке исходной матрицы взять подматрицы размерностью = размерности ядра Собеля и почленно перемножить с соответствующими компонентами ядра, сложить значения, поделить на их количество. Это будет производная по одному направлению (остальные производные найти несложно).
Так, каким образом организовать ядро, не могу сообразить. Главная задача , которую нужно распараллелить, как мне кажется - сложение элементов после их перемножения (т.е если у нас ядро Собеля 9*9*9 = 729 операций сложения для каждого элемента исходной матрицы 32*32*32 = 32768).
Нашел пример "редукции" здесь: . Здесь для одного массива (у меня их получится 32768). Можно добавить blockIdx.y для перемещения между массивами: вместо

  1. reduce<<<( BLOCK_SIZE, 1, 1 ), (8,0,0)>>>
  2. //Сделать:
  3. reduce<<<( BLOCK_SIZE, 32768, 1 ), (8,0,0)>>>

Уперся в стену: В алгоритме редукции сложение элементов массива - пирамидообразное. Т.е все элементы массива в процессе сложения меняются :

Это недопустимо, так как одни и те же элементы используются в разных подматрицах. Единственным решением мне кажется создавать двумерные массивы: [32768][729]. Т.е вместо того, чтобы брать соседние элементы в исходной матрице, представить исходную матрицу как массив подматриц матриц. И для каждой такой подматрицы выполнять сложение элементов. Тут получается оперирование огромными объемами памяти, что плохо.
Сейчас у меня все организовано в виде одномерных массивов, где доступ к элементам осуществляется так: [x+x_size*y+x_size*y_size*z] это трудно для понимания (в коде происходит частый пересчет индексов), зато нет проблем с транспортированием данных на девайс и обратно. Код не выкладываю, он не рабочий. Требуется реорганизация.
Подскажите пожалуйста как лучше организовать решение такой задачи.

Forums: 

Смотря на какой карте

Смотря на какой карте считать, если Ферми то тут лучше делать шару 32*32*9 и там уже каждой нити давать считать свертку матриц 9*9*9, КПД будет много выше, чем при редукции имхо.
Получится 729 умножений с накоплением для каждого вновь считанного числа, что есть гуд...
Если на G200 то тут сложнее, тк к шару не влезет такой большой массив, будет много обращений в глобальную память.