Передача аргументів через __локову пам'ять в OpenCL - пам'яті, opencl

Мене бентежить __локальна пам'ять у OpenCL тут. Я прочитав деякі специфікації, які говорять, що потік даних повинен бути від хоста до __Глобальний, а потім __Місцевий. Але я також бачу деякі функції ядра, як це:

__kernel void foo(__local float * a)

Мені було цікаво, як дані передавалися безпосередньо в __local пам'ять таким чином?

Дякую.

Відповіді:

Неможливо заповнити локальний буфер на стороні хоста. Тому ви повинні слідувати хосту потоку -> __global -> __local.

Локальний буфер можна створити на хості, а потім передається як параметр ядра або на сторону gpu всередині ядра. Створення локального буфера на стороні хоста дає перевагу вирішити питання про його розмір перед запуском ядра, що може бути важливим, якщо розмір локального буфера повинен бути іншим при кожному запуску ядра.

Місцева пам’ять не видно ні для кого, крім aодна робоча група і може бути виділена, оскільки робоча група надсилається апаратним забезпеченням для багатьох архітектур. Апаратне забезпечення, яке може змішувати кілька робочих груп з різних ядер на кожній CU, дозволить компоненту планування обробляти локальну пам’ять для кожної з виданих груп. Він не існує до запуску групи і не існує після завершення роботи групи. Розмір цієї області - це те, що ви вказуєте, як зазначають інші відповіді.

Результатом цього є єдиний спосіб для багатьохархітектури для заповнення локальної пам'яті з хосту полягали б у тому, щоб компілятор вставляв код ядра, який копіював би дані з глобальної пам'яті. Беручи до уваги, що в якості основи, це не гірше з точки зору продуктивності, щоб програміст робив це вручну, і дає більше контролю над тим, що саме відбувається. Ви не потрапите в ситуацію, коли компілятор завжди генерує код копіювання і закінчує копіювання більше, ніж було насправді необхідним, оскільки API не давав зрозуміти, яка пам'ять копіюється, а яка ні.

Підсумовуючи, ви не можете заповнити локальну пам’ять ні в одномуавтоматизований спосіб. На практиці вам рідко захочеться, тому що це вручну дає можливість лише помістити результат першого етапу в локальний, видаливши зайві операції копіювання, або перетворити дані на шляху до локальних, дозволяючи заповнення або транспонування даних для усунення банківських конфліктів тощо.

Як сказав @doqtor, розмір локальної пам'яті для параметра ядра можна вказати за допомогою clSetKernelArg дзвінки

На щастя, OpenCL 1.2+ підтримує VLA (масив змінної довжини), параметр ядра локальної пам'яті більше не потрібен.