Thread blocks и симметричные мультипроцессоры

Доброго дня.
Есть у меня несколько непоняток с CUDA, может поможете разрешить.
К примеру в документации сказано, что Thread Block может содержать максимум 512 тредов и выполняется всегда на одном мультипроцессоре. Косвенно в пользу этого факта говорит и возможность синхронизации тредов внутри одного блока с помощью __syncthreads(). Верно я утверждаю? И верно ли из этого следующее утверждения - все warp создаются в пределах одного блока и значит выполняются на одном мультипроцессоре. Т.е. вероятно планировщик выбирает следующий warp за 4 такта выполняет его за 4 такта - по такту на каждый тред. Затем выбирает другой warp (возможно тот же).
Тогда таким образом, для загрузки всех мультипроцессоров необходимо создать как минимум столько блоков, сколько этих самых мультипроцессоров. Т.е. в моем случае с GTX 280 (где 30 мультипроцессоров и 240 потоковых процессоров) - необходимо минимум 30 блоков? И могу ли в этом случае я ожидать линейного прироста производительности? Просто в моем случае я получаю прирост только раза в 3 по сравнению с загрузкой одного блока.

Forums: 

Блоков надо бы еще больше

Блоков надо бы еще больше (т.к. в ситуации, когда что-то ждет синхронизации мультипроцессор может простаивать), мануал рекомендует "100, а лучше 1000".

А столь незначительный прирост может быть связан с неудачным паттерном доступа к памяти, либо же со слишком простым kernel (все лимитируется, например, той же памятью или скоростью загрузки-выгрузки)

планировщик выдаёт команду на

планировщик выдаёт команду на исполнение за 2 такта. За 4 она исполняется на 8-ми полосном конвейере. Таким образом реализовано dual-issue.

прирост производительности при масштабировании числа блоков от 1 до 30 на GTX280 в самом деле должен быть линейным.