Jak korzystać z async_work_group_copy w OpenCL?

Question

Chciałbym zrozumieć, jak poprawnie używać wywołania async_work_group_copy() w OpenCL. Rzućmy okiem na uproszczonym przykładzie.

__kernel void test(__global float *x) { 
    __local xcopy[GROUP_SIZE]; 

    int globalid = get_global_id(0); 
    int localid = get_local_id(0); 
    event_t e = async_work_group_copy(xcopy, x+globalid-localid, GROUP_SIZE, 0); 
    wait_group_events(1, &e); 
} 

Odniesienie http://www.khronos.org/registry/cl/sdk/1.0/docs/man/xhtml/async_work_group_copy.html mówi „Wykonaj kopię asynchronicznej elementów num_elements gentype z src do dst Kopia asynchroniczny jest wykonywany przez wszystkich roboczych elementów w pracy grupowej i ta wbudowana funkcja musi więc zostać napotkana przez wszystkie elementy pracy w grupie roboczej wykonującej jądro z tymi samymi wartościami argumentów, w przeciwnym razie wyniki nie zostaną zdefiniowane. "

Ale to nie wyjaśnia moje pytania ...

Chciałbym wiedzieć, jeżeli następujące założenia są poprawne:

1. Wezwanie do async_work_group_copy() musi być wykonany przez wszystkich prac -pis w grupie.
2. Połączenie powinno być takie, że adres źródłowy jest identyczny dla wszystkich elementów roboczych i wskazuje na pierwszy element obszaru pamięci, który ma zostać skopiowany.
3. Ponieważ mój adres źródłowy jest względny w oparciu o globalny identyfikator elementu pracy pierwszego elementu roboczego w grupie roboczej. Więc muszę odjąć lokalny identyfikator, aby adres był taki sam dla wszystkich elementów pracy ...
4. Czy trzeci parametr to naprawdę liczba elementów (nie rozmiar w bajtach)?

pytania Bonus:

a. Czy mogę po prostu użyć bariery (CLK_LOCAL_MEM_FENCE) zamiast wait_group_events() i zignorować zwracaną wartość? Jeśli tak, czy byłoby to prawdopodobnie szybsze?

b. Czy kopia lokalna ma sens również w przypadku przetwarzania na procesorach, czy jest to obciążenie ogólne, ponieważ mimo to współużytkują pamięć podręczną?

Pozdrawiam, Stefan

Answer 1

Jednym z głównych powodów, dla tej funkcji jest umożliwienie istniejącym kompilator kierowcy/kernel skutecznie skopiować pamięć bez dewelopera konieczności dokonywania założenia dotyczące sprzętu.

Opisujesz, jaka pamięć jest kopiowana tak, jakby była kopią jednowątkową, a async_work_group_copy robi to za ciebie przy użyciu sprzętu równoległego.

dla konkretnego pytania:

1. nigdy nie widziałem async_work_group_copy wykorzystywane przez tylko niektóre z elementów pracy w grupie. Zawsze zakładałem, że to dlatego, że to konieczne. Myślę, że blokujący charakter wait_group_events powoduje, że wszystkie elementy pracy stają się częścią kopii.

2. Tak. Adresy źródłowe (i docelowe) muszą być takie same dla wszystkich elementów pracy.

3. Można odjąć lokalny identyfikator, aby uzyskać prawidłowy adres, ale uważam, że opieranie adresu na identyfikatorze GroupID również rozwiązuje ten problem. (get_group_id)

4. Tak. Ostatnim parametrem jest liczba elementów, a nie rozmiar w bajtach.

a. Nie. Oparty na zdarzeniu przekonasz się, że twoja bariera jest niemal natychmiast uderzana przez elementy pracy, a dane nie będą koniecznie kopiowane. Ma to sens, ponieważ niektóre urządzenia opencl mogą nie używać w ogóle jednostek obliczeniowych do rzeczywistej operacji kopiowania.

b. Myślę, że implementacje openp cpu mogą zagwarantować wykorzystanie pamięci podręcznej L1 podczas korzystania z pamięci lokalnej. Jedynym sposobem, aby się upewnić, czy to działa lepiej, jest porównywanie aplikacji z różnymi ustawieniami.