Осваиваю OpenCL. О том, как работают группы, объясняется на пальцах и картинках, но не уверен, правильно ли я понимаю, как оно управляется из кода. Поэтому у меня несколько вопросов. Для того чтобы посмотреть, что там происходит, я сделал следующее (да, использую Boost.Compute, но думаю интерфейсы у всех примерно одинаковые): 1. Завел буфер на 1024*1024 элементов. 2. Установил global work и local work размеры, и выполнил программу: Код (C++): std::size_t const size = 1024 * 1024; compute::vector<std::size_t> device_data(size, context); kernel.set_arg(0, device_data); std::size_t global_work_size[] = { size }; std::size_t local_work_size[] = { 256 }; queue.enqueue_nd_range_kernel(kernel, 1, 0, global_work_size, local_work_size); queue.finish(); Сама программа выглядит так: Код (C): kernel void process(global size_t* output) { output[get_global_id(0)] = get_group_id(0); } То есть в каждый элемент буфера записывается номер группы, в которой выполнялась задача, номер которой (задачи) совпадает с индексом элемента в буфере. Посчитал статистику, и увидел, что всего групп было 4096, по 256 item-ов в каждой. Теперь вопросы: 1. Правильно ли я понимаю, что barrier(CLK_GLOBAL_MEM_FENCE) будет работать только внутри группы? То есть если добавить его в мой пример, то синхронизироваться будут не все 1M айтемов, а группы по 256 штук, а сами группы будут работать как придется? 2. Правильно ли я понимаю, что параллельно работают вот эти 256 айтемов, а группы - в каком-то там порядке по принципу тредпула? То есть последовательно-параллельно по мере готовности видеокарты? 3. Я поставил local work size = 256 потому что у меня CL_DEVICE_MAX_WORK_GROUP_SIZE = 256. Это и есть максимальная "параллельность" моей видеокарты, или все-таки сами группы (по 256 параллельных айтемов в каждой) могут тоже работать параллельно? Если да, то как узнать, сколько групп максимального размера могут работать одновременно?
Для начала картинка для понимания структуры групп и айтемов Эта картинка дает понять, какой параметр зависит от железа. Твое число 256 - это кол-во workitems в одной группе. Тоже зависит от железа. Но тут ограничение сделано из-за ограничения кол-ва локальный (по сути shared) памяти. Эта память видна всем потокам внутри одной группы. Кол-во групп ты можешь посчитать - 1024*1024/256 = 4096 Для пользователя доступна возможность опираться на синхронизацию по workitems внутри одной группы (локально) так и по workitems глобально Считается, что параллельно выполняется группа. Но это не так. Реально параллельно выполняется Wavefront У АМД он 64 на сколько я помню . У Nvidia это называется warp и их 32. Обычно , для простого пользователя, на wavefront по-барабану. По этому считают что все потоки выполняются параллельно. А синхронизируют их локально или глобально. Теперь по вопросам Нет. Такая синхронизация будет ждать все треды в испольнительной системе. (для синхронизации глобальной памяти) то есть пока все треды не запишут \ сделают все свои дела, до этого участка кода Будет синхронизироваться 1м тредов. Думаю на этот вопрос я ответил описанием выше. резюмируя : wavefront работает параллельно (по настоящему параллельно, как simd блоки в CPU) а остальное уже по мере готовности. С точки зрения пользователя - он пишет программу так, что у него все треды параллельны (имеется ввиду, что нет четкого порядка следования. Начать выполнятся может любая из групп) Это не максимальная параллельность. Это просто зависит от локально доступной памяти. (64кб). То есть по сути у тебя каждый айтем может писать в локальную память (общую) которую видно только одной группе. Принцип программирования заключается не в том, чтоб понять сколько могут работать одновременно. Это нужно оставить производителям видеокарт. Пусть они заботятся об параллельности. Пользователю нужно правильно и грамотно распределить общие ресурсы которых мало.
TermoSINteZ, спасибо за разъяснения Да, я понимаю. Вопрос тут просто из любопытства. И еще вопрос: в чем смысл dimensions? То есть, в чем смысл - примерно ясно, но что оно дает практически. Посмотрел обсуждения, и например предлагается делать умножение матриц через 2 dimensions. Почему это лучше, чем 1 dimension и работать с матрицей как с одномерным массивом?
_DEN_, Ну потому что меньше кода писать. У тебя треды имеют по сути свои индексы. Через геттеры получаешь нужный "адрес треда" и уже четко знаешь с чем работаешь. Но по факту разницы нет. Можно и так и так. Основной принцип - меньше вычислений - быстрее кернел. Максимальное кл-во элементов все равно больше некоторого X не будет. то есть если допустим максимус 1024, 1024, 64. (по 3м измерениям). То это будет эквивалентно максимуму по по 1 измерению равному 1024*1024*64
