Перейти к содержимому
Compvision.ru

Recommended Posts

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

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

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

Поделиться сообщением


Ссылка на сообщение
Поделиться на других сайтах

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

Поделиться сообщением


Ссылка на сообщение
Поделиться на других сайтах

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

Поделиться сообщением


Ссылка на сообщение
Поделиться на других сайтах

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

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

Поделиться сообщением


Ссылка на сообщение
Поделиться на других сайтах

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

Поделиться сообщением


Ссылка на сообщение
Поделиться на других сайтах

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

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

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

Поделиться сообщением


Ссылка на сообщение
Поделиться на других сайтах

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

Поделиться сообщением


Ссылка на сообщение
Поделиться на других сайтах

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

Поделиться сообщением


Ссылка на сообщение
Поделиться на других сайтах
__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 - размер входного массива. Не меняется.

Поделиться сообщением


Ссылка на сообщение
Поделиться на других сайтах

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

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

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

 

Поделиться сообщением


Ссылка на сообщение
Поделиться на других сайтах

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

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

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

Поделиться сообщением


Ссылка на сообщение
Поделиться на других сайтах

Создайте учётную запись или войдите для комментирования

Вы должны быть пользователем, чтобы оставить комментарий

Создать учётную запись

Зарегистрируйтесь для создания учётной записи. Это просто!

Зарегистрировать учётную запись

Войти

Уже зарегистрированы? Войдите здесь.

Войти сейчас


  • Сейчас на странице   0 пользователей

    Нет пользователей, просматривающих эту страницу

×