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

[erly 0]

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

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

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

[Nuzhny 243]

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

[erly 0]

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

[erly 0]

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

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

[Nuzhny 243]

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

[erly 0]

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

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

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

[Smorodov 578]

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

[erly 0]

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

[erly 0]

__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 - размер входного массива. Не меняется.

[Smorodov 578]

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

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

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

 

[erly 0]

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

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

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