Jump to content
Compvision.ru
erly

CUDA Фильтрация данных

Recommended Posts

Коллеги, кто имеет опыт в CUDA, подскажите, пожалуйста.

Возможно ли в принципе в кернел-фукнции выполнить фильтрацию данных? Так, чтобы размер данных на выходе функции был переменным.
Пример задачи: На входе массив значений float, в результирующий массив надо записать только те значения, которые превышают заданный порог.

Как такая задача решается? Не могу сообразить, как заставить треды выполнить такую синхронизированную запись в выходной массив. Тут получается, что индекс в выходном массиве априорно не известен, зависит от реультатов работы остальных нитей.

Share this post


Link to post
Share on other sites

Техника называется reduction, когда сначала все потоки пишут свои значения в результат, потом половина из них пишет валидные значения, затем ещё половина и т.д. Пока не останется один поток, определяющий финальный размер результата.

Share this post


Link to post
Share on other sites

Спасибо, большое. Видел такое раньше, но не понял до конца значит. Пошел изучать..

Share this post


Link to post
Share on other sites

Почитал про свертку еще. Так и не понял, как ее использовать в описанной выше задаче. Свертка во всех примерах используется для операций типа подсчета суммы всех элементов массива, т.е. когда размер выходного массива заранее известен - фиксированный, или кратен размеру входного массива, или строго зависит от него.

Если не сложно, опиши немного подробнее как применить свертку для фильтрации данных.

Share this post


Link to post
Share on other sites

Я про свертку не писал - редукция же. Каждый второй поток складывает полезные результаты в свой кусочек памяти и запоминает сколько и где, потом каждый четвёртый за двумя предыдущими и т.д. Теоретически, это должно сработать быстро.

Share this post


Link to post
Share on other sites

Ну да, свертка.. редукция - это ж одно и то же. Я про них и говорю. Для суммирования они подходят - итеративно общий результат так получать удобно. А в моей фильтрации по идее итерировать не надо: можно за один проход все вычислить, результат зависит только от одного элемента. Зависимости между тредами возникают только из-за адреса, в который результат надо записать. Например: задача отфильтровать значения по порогу 100. Для входного массива [15, 131, 618] результат будет [131, 618]. То есть вторая нить (которой досталось отфильтровать значение 131) должна как-то понять, что писать в выходной массив надо по индексу 0, т.к. первая нить (фильтрующая 15) в выходной массив свое значение не пропускает.

Смотрю сейчас на atomic functions. Думаю, что на них надо делать, но не понимаю, как они с абсолютным проседанием производительности борятся.

Ну либо все еще не улавливаю всех возможностей редукции..

Share this post


Link to post
Share on other sites

Может Trust заюзать ? Помножить на 0 все что отбрасываем, отсортировать и обрезать вектор по первому нулю. 

Share this post


Link to post
Share on other sites

Спасибо, на него тоже посмотрю. Thrust по любому на базовых CUDA примитивах построен, чудес наверное не стоит от него ждать.

Share this post


Link to post
Share on other sites
__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

Корень можно не извлекать, лучше в квадрат возвести радиус.

Ну и, как мне видится, атомный доступ весь эффект сводит на нет.

Может в параллель только флаги расставить, а скомпоновать массив на CPU ?

 

Share this post


Link to post
Share on other sites

Ну, я ж специально на атомных функциях делал. Вообще, производительность получилась очень хорошая, как ни странно. Надо придумать, почему так получилось..

Как вариант, такое обоснование:
1. Отфильтрованных точек существенно меньше, чем число точек в исходном массиве (примерно 50к против 900к), поэтому атомная функция вызывалась не очень часто.
2. Компоновка выходного массива выполняется параллельно и с хорошей слитностью адресов памяти.

Возможно, как раз второй пункт имеет решающее значение для производительности. Потому что в непосредственно фильтрации ничего ресурсоемкого нет.

Share this post


Link to post
Share on other sites

Create an account or sign in to comment

You need to be a member in order to leave a comment

Create an account

Sign up for a new account in our community. It's easy!

Register a new account

Sign in

Already have an account? Sign in here.

Sign In Now


  • Recently Browsing   0 members

    No registered users viewing this page.

×