erly 0 Report post Posted November 3, 2019 Коллеги, кто имеет опыт в CUDA, подскажите, пожалуйста. Возможно ли в принципе в кернел-фукнции выполнить фильтрацию данных? Так, чтобы размер данных на выходе функции был переменным. Пример задачи: На входе массив значений float, в результирующий массив надо записать только те значения, которые превышают заданный порог. Как такая задача решается? Не могу сообразить, как заставить треды выполнить такую синхронизированную запись в выходной массив. Тут получается, что индекс в выходном массиве априорно не известен, зависит от реультатов работы остальных нитей. Share this post Link to post Share on other sites
Nuzhny 243 Report post Posted November 3, 2019 Техника называется reduction, когда сначала все потоки пишут свои значения в результат, потом половина из них пишет валидные значения, затем ещё половина и т.д. Пока не останется один поток, определяющий финальный размер результата. Share this post Link to post Share on other sites
erly 0 Report post Posted November 3, 2019 Спасибо, большое. Видел такое раньше, но не понял до конца значит. Пошел изучать.. Share this post Link to post Share on other sites
erly 0 Report post Posted November 4, 2019 Почитал про свертку еще. Так и не понял, как ее использовать в описанной выше задаче. Свертка во всех примерах используется для операций типа подсчета суммы всех элементов массива, т.е. когда размер выходного массива заранее известен - фиксированный, или кратен размеру входного массива, или строго зависит от него. Если не сложно, опиши немного подробнее как применить свертку для фильтрации данных. Share this post Link to post Share on other sites
Nuzhny 243 Report post Posted November 4, 2019 Я про свертку не писал - редукция же. Каждый второй поток складывает полезные результаты в свой кусочек памяти и запоминает сколько и где, потом каждый четвёртый за двумя предыдущими и т.д. Теоретически, это должно сработать быстро. Share this post Link to post Share on other sites
erly 0 Report post Posted November 4, 2019 Ну да, свертка.. редукция - это ж одно и то же. Я про них и говорю. Для суммирования они подходят - итеративно общий результат так получать удобно. А в моей фильтрации по идее итерировать не надо: можно за один проход все вычислить, результат зависит только от одного элемента. Зависимости между тредами возникают только из-за адреса, в который результат надо записать. Например: задача отфильтровать значения по порогу 100. Для входного массива [15, 131, 618] результат будет [131, 618]. То есть вторая нить (которой досталось отфильтровать значение 131) должна как-то понять, что писать в выходной массив надо по индексу 0, т.к. первая нить (фильтрующая 15) в выходной массив свое значение не пропускает. Смотрю сейчас на atomic functions. Думаю, что на них надо делать, но не понимаю, как они с абсолютным проседанием производительности борятся. Ну либо все еще не улавливаю всех возможностей редукции.. Share this post Link to post Share on other sites
Smorodov 578 Report post Posted November 5, 2019 Может Trust заюзать ? Помножить на 0 все что отбрасываем, отсортировать и обрезать вектор по первому нулю. Share this post Link to post Share on other sites
erly 0 Report post Posted November 5, 2019 Спасибо, на него тоже посмотрю. Thrust по любому на базовых CUDA примитивах построен, чудес наверное не стоит от него ждать. Share this post Link to post Share on other sites
erly 0 Report post Posted November 5, 2019 __global__ void TubeFilter(float *in, float *out, float *limits, uint *out_size, uint in_size) { __shared__ uint save_count; // count of successfully filtered items in current block if (threadIdx.x == 0) { save_count = 0; // initialize by some one thread } __syncthreads(); uint read_idx = blockIdx.y * blockDim.x * blockDim.y * gridDim.x + blockIdx.x * blockDim.x * blockDim.y + threadIdx.x; if (read_idx >= in_size) { return; } uint r3 = read_idx * 3; // each thread works with all 3 dimensions of a point float x = in[r3]; float y = in[r3 + 1]; float z = in[r3 + 2]; bool save_flag = false; // true if the filter conditions are complied uint save_idx = 0; // the thread will write to this index of out array if (sqrt(x*x + y*y) <= limits[0] && z >= limits[1] && z <= limits[2]) { save_flag = true; save_idx = atomicAdd(&save_count, 1); } __syncthreads(); __shared__ uint first; // first write position of current block if (threadIdx.x == 0) { // increase the common size of out array in global memory first = atomicAdd(out_size, save_count); } if (save_flag) { uint w3 = (first + save_idx) * 3; out[w3] = x; out[w3 + 1] = y; out[w3 + 2] = z; } } Вроде получилось ядро фильтрации. На двух atomicAdd, один из них для значения в shared памяти, второй - для global памяти. Проревьюйте, пожалуйста, кому интересно. Здесь реализован фильтр координат, попадающих в цилиндрическую область заданного радиуса и высоты. in - входной массив 3D координат - float x, y, z координаты. out - выходной массив 3D координат. В него пишутся только точки, удовлетворяющие фильтру. limits - параметры фильтрации: [0] - радиус, [1] - минимальное значение по оси Z, [2] - максимальное значение по Z. out_size - текущий размер выходного массива. Аллоцирован в global памяти. Изначально установлен в host-функции в 0, затем увеличивается в каждом блоке на количество отфильтрованных точек. in_size - размер входного массива. Не меняется. Share this post Link to post Share on other sites
Smorodov 578 Report post Posted November 6, 2019 Корень можно не извлекать, лучше в квадрат возвести радиус. Ну и, как мне видится, атомный доступ весь эффект сводит на нет. Может в параллель только флаги расставить, а скомпоновать массив на CPU ? Share this post Link to post Share on other sites
erly 0 Report post Posted November 6, 2019 Ну, я ж специально на атомных функциях делал. Вообще, производительность получилась очень хорошая, как ни странно. Надо придумать, почему так получилось.. Как вариант, такое обоснование: 1. Отфильтрованных точек существенно меньше, чем число точек в исходном массиве (примерно 50к против 900к), поэтому атомная функция вызывалась не очень часто. 2. Компоновка выходного массива выполняется параллельно и с хорошей слитностью адресов памяти. Возможно, как раз второй пункт имеет решающее значение для производительности. Потому что в непосредственно фильтрации ничего ресурсоемкого нет. Share this post Link to post Share on other sites