Tło projektu Step1
Niedawno opublikowałem na Githubie oprogramowanie open-source o nazwie "Yunhan Grill Master". Adres repozytorium kodu:
.
To oprogramowanie do grillowania zostało zaprojektowane z uwzględnieniem układu RK3399, w tym zarówno dużych, jak i małych architektur rdzeni CPU i GPU z obsługą OpenCL. Program może utrzymywać pełne obciążenie CPU i GPU procesora przez dłuższy czas za pośrednictwem Pthread, OpenMP i OpenCL oraz drukować w czasie rzeczywistym temperaturę procesora, stan online rdzenia procesora, częstotliwość dużego rdzenia procesora, częstotliwość małego rdzenia procesora, prędkość obliczeniową procesora i prędkość obliczeniową GPU na terminalu co 5 sekund. Można również częściowo skomentować funkcje w kodzie, aby można je było zastosować do układów procesora, które nie obsługują OpenCL lub mają małą liczbę rdzeni. Na przykład, podczas modyfikowania kodu pod kątem kompatybilności sprzętowej, zrozumienie podstawowych pojęć, takich jak " co to jest GND w obwodzie elektronicznym " może pomóc uniknąć błędów w konfiguracjach związanych z obwodem.
Step2 Testbed
Społeczność Leez udostępniła komputer jednopłytkowy P710 oparty na układzie RK3399. Grupa dokonała już przeglądu konfiguracji sprzętowej tego małego komputera; zobacz "Leez-P710: Wypalanie na wszystkich cylindrach".
Ja również miałem szczęście wypróbować Leez P710. Poniżej znajduje się mój testowy produkt. Widać, że specjalnie zaprojektowałem dla niego kanały chłodzące.
Poniżej znajduje się obudowa, którą zaprojektowałem i wydrukowałem w 3D. Jest to małe pudełko z portami i kanałami powietrznymi pozostawionymi na bocznych ściankach.
Po zainstalowaniu białych nóżek, pudełko może stać, tak jak mała obudowa.
Dzisiejszy test zostanie przeprowadzony na Leez P710.
Krok 3 Konfiguracja środowiska
Oprogramowanie firmware Debiana wydane przez społeczność Leez jest już dostarczane z GCC i biblioteką OpenMP.
Jest jednak jedna rzecz, na którą należy zwrócić uwagę: GPU w Leez P710 to Mali T860, ale nie ma pliku nagłówkowego OpenCL w firmware Debiana wydanym przez społeczność Leez, więc pojawią się problemy podczas kompilacji kodu. Możemy przejść do repozytorium rockchip-linux na GitHub, aby pobrać bibliotekę obsługi GPU Mali za pomocą polecenia:
git clone https://github.com/rockchip-linux/libmali.git
Następnie umieszczamy include w katalogu /usr i bezpośrednio łączymy foldery.
Krok 4 Kompilacja i uruchomienie
Pobierz kod źródłowy z repozytorium wspomnianego na początku tego artykułu.
Umieść icy.c i przetestuj. Skopiuj pliki .cl w folderze "GrillMaster" do wygodnego katalogu, takiego jak /home/test.
Skompiluj polecenia w następujący sposób:
gcc -o icy.c -lpthread -lOpenCL -fopenmp
Wykonaj
. /ickey
Po uruchomieniu programu poniższy obraz wyświetla informacje drukowane w czasie rzeczywistym z terminala. Widzimy, że temperatura wynosi około 70 stopni. Sześć rdzeni procesora (od 0 do 5) jest online, z mniejszą częstotliwością rdzenia A53 wynoszącą 1,42 GHz i dużą częstotliwością rdzenia A72 wynoszącą 1,8 GHz. Nie ma tragedii, aby jeden rdzeń miał trudności z wieloma otaczającymi go rdzeniami.
.
Wśród nich częstotliwość rdzenia CPU zmienia się dynamicznie, spadając niekiedy do 1,42. Ponieważ wykorzystujemy wszystkie rdzenie procesora, temperatura jest nieco wysoka, co uruchamia aktywny downclocking.
Wartości "cpu used = xxx" i "ocl used = xxx" w sekcji Powyższy wykres jest wynikiem uruchomienia różnych algorytmów na CPU i GPU i pomiaru czasu (dokładne algorytmy podano w sekcji 5). Można zaobserwować, że CPU i GPU znacznie zwalniają, gdy temperatura jest zbyt wysoka. CPU wykonywał mnożenie i dzielenie zmiennoprzecinkowe w celu obliczenia liczby pi, podczas gdy GPU wykorzystywał sortowanie bąbelkowe do sortowania danych całkowitych.
Kontynuując wypalanie maszyny, końcowa temperatura ustabilizowała się na poziomie około 73 stopni, z niewielkimi wahaniami. Po dynamicznym obniżeniu taktowania procesora w celu redukcji ciepła, częstotliwość wzrosła, następnie temperatura wzrosła, a dynamiczne obniżenie taktowania nastąpiło ponownie. Częstotliwość rdzenia A72 spadła do 1,2 GHz, a temperatura wyzwalacza wynosiła około 72,8 stopnia. Miało to również wpływ na szybkość obliczeń, przy czym obliczenia zmiennoprzecinkowe procesora były o około 15% wolniejsze po obniżeniu częstotliwości taktowania.
Wyjaśnienie zasady działania programu Step5
W funkcji podstawowej otwierane są trzy wątki.
Wątek 1 odczytuje i drukuje temperaturę procesora, częstotliwość rdzeni procesora i status online rdzeni procesora w określonych odstępach czasu (aby sprawdzić, czy fiasko jednego rdzenia mającego trudności z wieloma otaczającymi go rdzeniami się wydarzy), a także drukuje na terminalu prędkości pracy uzyskane przez wątki 2 i 3.
Wątek 2 wypełnia CPU poprzez umieszczenie kodu obliczeniowego pi wewnątrz pętli while(1).
Wątek 3 ładuje procesor graficzny, wstawiając kod sortowania bąbelkowego do pętli while(1).
1. int main()
2. {
3. pthread_t tid_term.
4. char* p_term = NULL;
5. pthread_create(&tid_term, NULL, thread_term, NULL); pthread_t tid_term = NULL.
6.
7. pthread_t tid_cpu;
8. char* p_cpu = NULL;
9. pthread_create(&tid_cpu, NULL, thread_cpu, NULL);
10.
11. pthread_t tid_ocl;
12. char* p_ocl = NULL;
13. pthread_create(&tid_ocl, NULL, thread_ocl, NULL);
14.
15. sl<ickey>eep(1);
16.
17. pthread_join(tid_term, (void **)&p_term);
18. pthread_join(tid_cpu, (void **)&p_cpu);
19. pthread_join(tid_ocl, (void **)&p_ocl);;;
20. return 0;
21. }
Uruchomienie wątku 1 w celu odczytania i wydrukowania informacji o temperaturze
1. void *thread_term(void *arg)
2. {
3. int fd.
4. while (1)
5. {
6. fd = open(TEMP_PATH, O_RDONLY);
7. char buf[20];
8. read(fd, buf, 20); double temp; char buf[20]; read(fd, buf, 20)
9. double temp; temp = atoi(buf)
10. read(fd, buf, 20); double temp; temp = atoi(buf) / 1000.0; printf("temperatura: %.0")
11. printf("temperatura: %.1lf\n",temp);
12. close(fd); system("cat /sys /sys"); system("cat /sys")
13.
14. system("cat /sys/devices/system/cpu/online"); fd = open(CPU0_0)
15.
16. fd = open(CPU0_PATH, O_RDONLY); char buf1[20];; fd = open(CPU0_PATH, O_RDONLY)
17. fd = open(CPU0_PATH, O_RDONLY); char buf1[20];
18. read(fd, buf1, 20); temp = atoi(buf1, 20); read(fd, buf1, 20)
19. temp = atoi(buf1) / 1000000.0;
20. printf("A53 Freq: %.2lf\n", temp);
21. close(fd);
22.
23. fd = open(CPU4_PATH, O_RDONLY);
24. fd = open(CPU4_PATH, O_RDONLY); char buf2[20];
25. read(fd, buf2, 20); temp = atoi(buf2, 20); read(fd, buf2, 20)
26. temp = atoi(buf2) / 1000000.0;
27. printf("A72 Freq: %.2lf\n", temp);
28. close(fd);
29. printf("\n"); if (iscpu == 1, temp); printf("A72 Freq.
30.
31. if (iscpu == 1)
32. printf("cpu used == 1); if (iscpu == 1)
33. printf("cpu used = %lf s\n", timecpu); if (iscpu == 1) {
34. iscpu = 0; }
35. }
36. if (isocl == 1)
37. {
38. printf("ocl used = %lf s\n", timeocl); isocl = 0; } if (isocl == 1) {
39. timeocl); isocl = 0; }
40. }
41. sl<ickey>eep(5);
42. }
43.
44.}
Funkcja grilla OpenMP dla wątku 2 do uruchomienia (aby zapełnić procesor)
1. void *thread_cpu(void *arg)
2. {
3. while (1)
4. {
5. double s = 1;
6. double pi = 0;
7. double i = 1.0;
8. double i = 1.0; double n = 1.0; double n = 1.0; double n = 1.0
9. double dt; double start_time; double start_time; double start_time; double start_time
10. double start_time; int cnt = 0; int
11. start_time = microtime(); int cnt = 0; start_time = microtime(); start_time = microtime()
12. start_time = microtime(); #pragma omp parallel for num_threads(6)
13. #pragma omp parallel for num_threads(6)
14. for (cnt = 0; cnt<100000000; cnt++)
15. {
16. pi += i; n = n + 2; #pragma omp parallel for num_threads(6)
17. n = n + 2; s = -s; n
18. s = -s.
19. i = s / n; }
20. }
21. pi = 4 * pi; dt = microtime() - start_time; }
22. dt = microtime() - start_time; timecpu = dt; }
23. timecpu = dt; iscpu = 1; }
24. iscpu = 1; }
25. }
26. }
Funkcja grilla OpenCL do uruchomienia dla wątku 3 (aby zapełnić GPU)
1. void *thread_ocl(void *arg)
2. {
3. int array_a[20] = { 0, 1, 8, 7, 6, 2, 3, 5, 4, 9,17,19,15,10,18,16,14,13,12,11 };
4. int array_b[20] = { 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,0,0,0,0,0,0,0,0,0,0,0,0,0 };
5. size_t datasize = 20 * sizeof(int);
6. size_t ocl_string_size; char *ocl_string; size_t ocl_string_size
7. char *ocl_string; double start_time, dt, dt, ocl_string_size
8. char *ocl_string; double start_time, dt, dt_err; start_time = microtime()
9. start_time = microtime(); dt_err = microtime()
10. dt_err = microtime() - start_time; ocl_string = (char *ocl_string)
11.
12. ocl_string = (char *)malloc(400);
13. //ocl_string = (char *)malloc(20);
14.
15. cl_platform_id platform_id.
16. cl_device_id device_id;
17. cl_context context; cl_command_queue command
18. cl_command_queue command_queue; cl_mem buffer_a, cl_command_queue; cl_command_queue command_queue
19. cl_mem buffer_a, buffer_b; cl_program program; cl_program_b, buffer_b
20. cl_program program; cl_kernel kernel
21. cl_kernel kernel; cl_event kernelEvent; cl_event kernelEvent
22. cl_event kernelEvent; clGetPlatformIDs; clGetPlatformIDs; clGetPlatformIDs
23.
24. clGetPlatformIDs(1, &platform_id, NULL); clGetDeviceIDs(1, &platform_id, NULL);
25. clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_GPU, 1, &device_id, NULL); clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_GPU, 1, &device_id, NULL);
26.
27. context = clCreateContext(NULL, 1, &device_id, NULL, NULL, NULL);
28. command_queue = clCreateCommandQueue(context, device_id, 0, NULL);
29.
30. buffer_a = clCreateBuffer(context, CL_MEM_READ_ONLY, datasize, NULL, NULL);
31. buffer_b = clCreateBuffer(context, CL_MEM_READ_ONLY, datasize, NULL, NULL); buffer_b = clCreateBuffer(context, CL_MEM_READ_ONLY, datasize, NULL, NULL);
32.
33.
34. ocl_string_size = get_ocl_string("test.cl", ocl_string);
35. clEnqueueWriteBuffer(command_queue, buffer_a, CL_FALSE, 0, \
36. datasize, array_a, 0, NULL, NULL);
37. clEnqueueWriteBuffer(command_queue, buffer_b, CL_FALSE, 0, \
38. datasize, array_b, 0, NULL, NULL); clEnqueueWriteBuffer(command_queue, buffer_b, CL_FALSE, 0, \
39. program = clCreateProgramWithSource(context, 1, (const char **)&ocl_string, \
40. &ocl_string_size, NULL);
41.
42. clBuildProgram(program, 1, &device_id, NULL, NULL, NULL);
43. kernel = clCreateKernel(program, "test", NULL);
44.
45. clSetKernelArg(kernel, 0, sizeof(cl_mem), &buffer_a);
46. clSetKernelArg(kernel, 1, sizeof(cl_mem), &buffer_b);
47.
48.
49. size_t global_work_size[1] = { 20 };
50. while (1)
51. {
52. start_time = microtime();
53. clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, \
54. global_work_size, NULL, 0, NULL, &kernelEvent);
55. clWaitForEvents(1, &kernelEvent).
56. clEnqueueReadBuffer(command_queue, buffer_b, CL_TRUE, 0, \
57. datasize, array_b, 0, NULL, NULL); clEnqueueReadBuffer(command_queue, buffer_b, CL_TRUE, 0, \
58. dt = microtime() - start_time - dt_err;
59. timeocl = dt.
60. timeocl = dt; timeocl = dt; isocl = 1; }
61. }
62.
63. clReleaseKernel(kernel); clReleaseProgram(program).
64. clReleaseProgram(program).
65. clReleaseCommandQueue(command_queue); clReleaseMemObject(buffer_queue)
66. clReleaseMemObject(buffer_a); clReleaseMemObject(buffer_a)
67. clReleaseMemObject(buffer_b); clReleaseContext(buffer_b); clReleaseContext(buffer_b)
68. clReleaseContext(context).
69. free(ocl_string);
70. }
Ponieważ używamy OpenCL, musimy również napisać plik OpenCL cl, który można nazwać test.cl.
Podsumowanie
Artykuł grupowicza, "leez-p710: on fire", pokazuje również zrzuty ekranu z testu warunków skrajnych przy użyciu StressAppTest.
Wyświetla kilka biegów, gdy duży rdzeń A72 RK3399 jest obniżany do 1,8 GHz, 1,6 GHz, 1,4 GHz, 1,2 GHz i tak dalej.
Ten test wykazał, że gdy częstotliwość rdzenia A72 RK3399 osiągnęła około 73 stopnie, spadła o trzy biegi do 1,2 GHz, a następnie pozostała stabilna. Ogólnie.
Ogólnie rzecz biorąc, wyniki testu "Yunhan Grill Master" nie różniły się znacząco od oczekiwanej sytuacji, chociaż wystąpiło pewne przegrzanie z powodu downclockingu. Nie było to jednak w takim samym stopniu jak w przypadku Raspberry Pi 3.
Wpływ taktowania 1,2 GHz na wydajność jest akceptowalny. W praktyce istnieje bardzo niewiele przypadków, w których zarówno CPU, jak i GPU są w pełni obciążone jednocześnie. Rdzeń A72 może utrzymywać częstotliwość 1,8 GHz przez dłuższy czas.
Niedawno opublikowałem na Githubie oprogramowanie open-source o nazwie "Yunhan Grill Master". Adres repozytorium kodu:
.
To oprogramowanie do grillowania zostało zaprojektowane z uwzględnieniem układu RK3399, w tym zarówno dużych, jak i małych architektur rdzeni CPU i GPU z obsługą OpenCL. Program może utrzymywać pełne obciążenie CPU i GPU procesora przez dłuższy czas za pośrednictwem Pthread, OpenMP i OpenCL oraz drukować w czasie rzeczywistym temperaturę procesora, stan online rdzenia procesora, częstotliwość dużego rdzenia procesora, częstotliwość małego rdzenia procesora, prędkość obliczeniową procesora i prędkość obliczeniową GPU na terminalu co 5 sekund. Można również częściowo skomentować funkcje w kodzie, aby można je było zastosować do układów procesora, które nie obsługują OpenCL lub mają małą liczbę rdzeni. Na przykład, podczas modyfikowania kodu pod kątem kompatybilności sprzętowej, zrozumienie podstawowych pojęć, takich jak " co to jest GND w obwodzie elektronicznym " może pomóc uniknąć błędów w konfiguracjach związanych z obwodem.
Step2 Testbed
Społeczność Leez udostępniła komputer jednopłytkowy P710 oparty na układzie RK3399. Grupa dokonała już przeglądu konfiguracji sprzętowej tego małego komputera; zobacz "Leez-P710: Wypalanie na wszystkich cylindrach".
Ja również miałem szczęście wypróbować Leez P710. Poniżej znajduje się mój testowy produkt. Widać, że specjalnie zaprojektowałem dla niego kanały chłodzące.
Poniżej znajduje się obudowa, którą zaprojektowałem i wydrukowałem w 3D. Jest to małe pudełko z portami i kanałami powietrznymi pozostawionymi na bocznych ściankach.
Po zainstalowaniu białych nóżek, pudełko może stać, tak jak mała obudowa.
Dzisiejszy test zostanie przeprowadzony na Leez P710.
Krok 3 Konfiguracja środowiska
Oprogramowanie firmware Debiana wydane przez społeczność Leez jest już dostarczane z GCC i biblioteką OpenMP.
Jest jednak jedna rzecz, na którą należy zwrócić uwagę: GPU w Leez P710 to Mali T860, ale nie ma pliku nagłówkowego OpenCL w firmware Debiana wydanym przez społeczność Leez, więc pojawią się problemy podczas kompilacji kodu. Możemy przejść do repozytorium rockchip-linux na GitHub, aby pobrać bibliotekę obsługi GPU Mali za pomocą polecenia:
git clone https://github.com/rockchip-linux/libmali.git
Następnie umieszczamy include w katalogu /usr i bezpośrednio łączymy foldery.
Krok 4 Kompilacja i uruchomienie
Pobierz kod źródłowy z repozytorium wspomnianego na początku tego artykułu.
Umieść icy.c i przetestuj. Skopiuj pliki .cl w folderze "GrillMaster" do wygodnego katalogu, takiego jak /home/test.
Skompiluj polecenia w następujący sposób:
gcc -o icy.c -lpthread -lOpenCL -fopenmp
Wykonaj
. /ickey
Po uruchomieniu programu poniższy obraz wyświetla informacje drukowane w czasie rzeczywistym z terminala. Widzimy, że temperatura wynosi około 70 stopni. Sześć rdzeni procesora (od 0 do 5) jest online, z mniejszą częstotliwością rdzenia A53 wynoszącą 1,42 GHz i dużą częstotliwością rdzenia A72 wynoszącą 1,8 GHz. Nie ma tragedii, aby jeden rdzeń miał trudności z wieloma otaczającymi go rdzeniami.
.
Wśród nich częstotliwość rdzenia CPU zmienia się dynamicznie, spadając niekiedy do 1,42. Ponieważ wykorzystujemy wszystkie rdzenie procesora, temperatura jest nieco wysoka, co uruchamia aktywny downclocking.
Wartości "cpu used = xxx" i "ocl used = xxx" w sekcji Powyższy wykres jest wynikiem uruchomienia różnych algorytmów na CPU i GPU i pomiaru czasu (dokładne algorytmy podano w sekcji 5). Można zaobserwować, że CPU i GPU znacznie zwalniają, gdy temperatura jest zbyt wysoka. CPU wykonywał mnożenie i dzielenie zmiennoprzecinkowe w celu obliczenia liczby pi, podczas gdy GPU wykorzystywał sortowanie bąbelkowe do sortowania danych całkowitych.
Kontynuując wypalanie maszyny, końcowa temperatura ustabilizowała się na poziomie około 73 stopni, z niewielkimi wahaniami. Po dynamicznym obniżeniu taktowania procesora w celu redukcji ciepła, częstotliwość wzrosła, następnie temperatura wzrosła, a dynamiczne obniżenie taktowania nastąpiło ponownie. Częstotliwość rdzenia A72 spadła do 1,2 GHz, a temperatura wyzwalacza wynosiła około 72,8 stopnia. Miało to również wpływ na szybkość obliczeń, przy czym obliczenia zmiennoprzecinkowe procesora były o około 15% wolniejsze po obniżeniu częstotliwości taktowania.
Wyjaśnienie zasady działania programu Step5
W funkcji podstawowej otwierane są trzy wątki.
Wątek 1 odczytuje i drukuje temperaturę procesora, częstotliwość rdzeni procesora i status online rdzeni procesora w określonych odstępach czasu (aby sprawdzić, czy fiasko jednego rdzenia mającego trudności z wieloma otaczającymi go rdzeniami się wydarzy), a także drukuje na terminalu prędkości pracy uzyskane przez wątki 2 i 3.
Wątek 2 wypełnia CPU poprzez umieszczenie kodu obliczeniowego pi wewnątrz pętli while(1).
Wątek 3 ładuje procesor graficzny, wstawiając kod sortowania bąbelkowego do pętli while(1).
1. int main()
2. {
3. pthread_t tid_term.
4. char* p_term = NULL;
5. pthread_create(&tid_term, NULL, thread_term, NULL); pthread_t tid_term = NULL.
6.
7. pthread_t tid_cpu;
8. char* p_cpu = NULL;
9. pthread_create(&tid_cpu, NULL, thread_cpu, NULL);
10.
11. pthread_t tid_ocl;
12. char* p_ocl = NULL;
13. pthread_create(&tid_ocl, NULL, thread_ocl, NULL);
14.
15. sl<ickey>eep(1);
16.
17. pthread_join(tid_term, (void **)&p_term);
18. pthread_join(tid_cpu, (void **)&p_cpu);
19. pthread_join(tid_ocl, (void **)&p_ocl);;;
20. return 0;
21. }
Uruchomienie wątku 1 w celu odczytania i wydrukowania informacji o temperaturze
1. void *thread_term(void *arg)
2. {
3. int fd.
4. while (1)
5. {
6. fd = open(TEMP_PATH, O_RDONLY);
7. char buf[20];
8. read(fd, buf, 20); double temp; char buf[20]; read(fd, buf, 20)
9. double temp; temp = atoi(buf)
10. read(fd, buf, 20); double temp; temp = atoi(buf) / 1000.0; printf("temperatura: %.0")
11. printf("temperatura: %.1lf\n",temp);
12. close(fd); system("cat /sys /sys"); system("cat /sys")
13.
14. system("cat /sys/devices/system/cpu/online"); fd = open(CPU0_0)
15.
16. fd = open(CPU0_PATH, O_RDONLY); char buf1[20];; fd = open(CPU0_PATH, O_RDONLY)
17. fd = open(CPU0_PATH, O_RDONLY); char buf1[20];
18. read(fd, buf1, 20); temp = atoi(buf1, 20); read(fd, buf1, 20)
19. temp = atoi(buf1) / 1000000.0;
20. printf("A53 Freq: %.2lf\n", temp);
21. close(fd);
22.
23. fd = open(CPU4_PATH, O_RDONLY);
24. fd = open(CPU4_PATH, O_RDONLY); char buf2[20];
25. read(fd, buf2, 20); temp = atoi(buf2, 20); read(fd, buf2, 20)
26. temp = atoi(buf2) / 1000000.0;
27. printf("A72 Freq: %.2lf\n", temp);
28. close(fd);
29. printf("\n"); if (iscpu == 1, temp); printf("A72 Freq.
30.
31. if (iscpu == 1)
32. printf("cpu used == 1); if (iscpu == 1)
33. printf("cpu used = %lf s\n", timecpu); if (iscpu == 1) {
34. iscpu = 0; }
35. }
36. if (isocl == 1)
37. {
38. printf("ocl used = %lf s\n", timeocl); isocl = 0; } if (isocl == 1) {
39. timeocl); isocl = 0; }
40. }
41. sl<ickey>eep(5);
42. }
43.
44.}
Funkcja grilla OpenMP dla wątku 2 do uruchomienia (aby zapełnić procesor)
1. void *thread_cpu(void *arg)
2. {
3. while (1)
4. {
5. double s = 1;
6. double pi = 0;
7. double i = 1.0;
8. double i = 1.0; double n = 1.0; double n = 1.0; double n = 1.0
9. double dt; double start_time; double start_time; double start_time; double start_time
10. double start_time; int cnt = 0; int
11. start_time = microtime(); int cnt = 0; start_time = microtime(); start_time = microtime()
12. start_time = microtime(); #pragma omp parallel for num_threads(6)
13. #pragma omp parallel for num_threads(6)
14. for (cnt = 0; cnt<100000000; cnt++)
15. {
16. pi += i; n = n + 2; #pragma omp parallel for num_threads(6)
17. n = n + 2; s = -s; n
18. s = -s.
19. i = s / n; }
20. }
21. pi = 4 * pi; dt = microtime() - start_time; }
22. dt = microtime() - start_time; timecpu = dt; }
23. timecpu = dt; iscpu = 1; }
24. iscpu = 1; }
25. }
26. }
Funkcja grilla OpenCL do uruchomienia dla wątku 3 (aby zapełnić GPU)
1. void *thread_ocl(void *arg)
2. {
3. int array_a[20] = { 0, 1, 8, 7, 6, 2, 3, 5, 4, 9,17,19,15,10,18,16,14,13,12,11 };
4. int array_b[20] = { 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0,0,0,0,0,0,0,0,0,0,0,0,0,0 };
5. size_t datasize = 20 * sizeof(int);
6. size_t ocl_string_size; char *ocl_string; size_t ocl_string_size
7. char *ocl_string; double start_time, dt, dt, ocl_string_size
8. char *ocl_string; double start_time, dt, dt_err; start_time = microtime()
9. start_time = microtime(); dt_err = microtime()
10. dt_err = microtime() - start_time; ocl_string = (char *ocl_string)
11.
12. ocl_string = (char *)malloc(400);
13. //ocl_string = (char *)malloc(20);
14.
15. cl_platform_id platform_id.
16. cl_device_id device_id;
17. cl_context context; cl_command_queue command
18. cl_command_queue command_queue; cl_mem buffer_a, cl_command_queue; cl_command_queue command_queue
19. cl_mem buffer_a, buffer_b; cl_program program; cl_program_b, buffer_b
20. cl_program program; cl_kernel kernel
21. cl_kernel kernel; cl_event kernelEvent; cl_event kernelEvent
22. cl_event kernelEvent; clGetPlatformIDs; clGetPlatformIDs; clGetPlatformIDs
23.
24. clGetPlatformIDs(1, &platform_id, NULL); clGetDeviceIDs(1, &platform_id, NULL);
25. clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_GPU, 1, &device_id, NULL); clGetDeviceIDs(platform_id, CL_DEVICE_TYPE_GPU, 1, &device_id, NULL);
26.
27. context = clCreateContext(NULL, 1, &device_id, NULL, NULL, NULL);
28. command_queue = clCreateCommandQueue(context, device_id, 0, NULL);
29.
30. buffer_a = clCreateBuffer(context, CL_MEM_READ_ONLY, datasize, NULL, NULL);
31. buffer_b = clCreateBuffer(context, CL_MEM_READ_ONLY, datasize, NULL, NULL); buffer_b = clCreateBuffer(context, CL_MEM_READ_ONLY, datasize, NULL, NULL);
32.
33.
34. ocl_string_size = get_ocl_string("test.cl", ocl_string);
35. clEnqueueWriteBuffer(command_queue, buffer_a, CL_FALSE, 0, \
36. datasize, array_a, 0, NULL, NULL);
37. clEnqueueWriteBuffer(command_queue, buffer_b, CL_FALSE, 0, \
38. datasize, array_b, 0, NULL, NULL); clEnqueueWriteBuffer(command_queue, buffer_b, CL_FALSE, 0, \
39. program = clCreateProgramWithSource(context, 1, (const char **)&ocl_string, \
40. &ocl_string_size, NULL);
41.
42. clBuildProgram(program, 1, &device_id, NULL, NULL, NULL);
43. kernel = clCreateKernel(program, "test", NULL);
44.
45. clSetKernelArg(kernel, 0, sizeof(cl_mem), &buffer_a);
46. clSetKernelArg(kernel, 1, sizeof(cl_mem), &buffer_b);
47.
48.
49. size_t global_work_size[1] = { 20 };
50. while (1)
51. {
52. start_time = microtime();
53. clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL, \
54. global_work_size, NULL, 0, NULL, &kernelEvent);
55. clWaitForEvents(1, &kernelEvent).
56. clEnqueueReadBuffer(command_queue, buffer_b, CL_TRUE, 0, \
57. datasize, array_b, 0, NULL, NULL); clEnqueueReadBuffer(command_queue, buffer_b, CL_TRUE, 0, \
58. dt = microtime() - start_time - dt_err;
59. timeocl = dt.
60. timeocl = dt; timeocl = dt; isocl = 1; }
61. }
62.
63. clReleaseKernel(kernel); clReleaseProgram(program).
64. clReleaseProgram(program).
65. clReleaseCommandQueue(command_queue); clReleaseMemObject(buffer_queue)
66. clReleaseMemObject(buffer_a); clReleaseMemObject(buffer_a)
67. clReleaseMemObject(buffer_b); clReleaseContext(buffer_b); clReleaseContext(buffer_b)
68. clReleaseContext(context).
69. free(ocl_string);
70. }
Ponieważ używamy OpenCL, musimy również napisać plik OpenCL cl, który można nazwać test.cl.
Podsumowanie
Artykuł grupowicza, "leez-p710: on fire", pokazuje również zrzuty ekranu z testu warunków skrajnych przy użyciu StressAppTest.
Wyświetla kilka biegów, gdy duży rdzeń A72 RK3399 jest obniżany do 1,8 GHz, 1,6 GHz, 1,4 GHz, 1,2 GHz i tak dalej.
Ten test wykazał, że gdy częstotliwość rdzenia A72 RK3399 osiągnęła około 73 stopnie, spadła o trzy biegi do 1,2 GHz, a następnie pozostała stabilna. Ogólnie.
Ogólnie rzecz biorąc, wyniki testu "Yunhan Grill Master" nie różniły się znacząco od oczekiwanej sytuacji, chociaż wystąpiło pewne przegrzanie z powodu downclockingu. Nie było to jednak w takim samym stopniu jak w przypadku Raspberry Pi 3.
Wpływ taktowania 1,2 GHz na wydajność jest akceptowalny. W praktyce istnieje bardzo niewiele przypadków, w których zarówno CPU, jak i GPU są w pełni obciążone jednocześnie. Rdzeń A72 może utrzymywać częstotliwość 1,8 GHz przez dłuższy czas.
