OpenCL :: کاوش در بعد اول (بخش 3: معیار صحیح) – تجزیه و تحلیل و پیش بینی – 3 مه 2023

قسمت 1 را بخوانید

قسمت 2 را بخوانید

تست را چند بار اجرا کنید، در بسیاری از تکرارها زمان اجرای یکسانی دریافت می‌کنیم که به این معنی است که چیزی اشتباه است یا پاسخ‌های حافظه پنهان دریافت می‌کنیم.

و این تابع تایمر است، بیایید ببینیم که آیا کار می کند، انجام می دهد:

2023.05.03 02:45:59.711 blog_benchmark_2 (USDJPY,H1)    Work Items (1) Iterations(1000)
2023.05.03 02:45:59.711 blog_benchmark_2 (USDJPY,H1)    Work Groups (1)
2023.05.03 02:45:59.711 blog_benchmark_2 (USDJPY,H1)    Milliseconds (15)

خوب حالا دنبال چی هستیم؟

هر کاری که به صورت موازی انجام شود در زمان نزدیک خواهد بود درست است؟ بنابراین ما به دنبال آن تاخیر اضافی هستیم که در مقایسه با اجرای کلی کوچک نیست

پس بیایید 2048 مورد را شروع کنیم و تکرارها را افزایش می دهیم تا بتوانیم بارها را متوجه شویم، بیایید دوباره 1 میلیون مورد را امتحان کنیم.

در واقع اجازه می دهد 1 مورد و 1 میلیون تکرار ارسال کنیم.

2023.05.03 02:49:57.474 blog_benchmark_2 (USDJPY,H1)    Work Items (1) Iterations(1000000)
2023.05.03 02:49:57.474 blog_benchmark_2 (USDJPY,H1)    Work Groups (1)
2023.05.03 02:49:57.474 blog_benchmark_2 (USDJPY,H1)    Milliseconds (16)

عالی، 1 مورد، 100 میلیون تکرار

2023.05.03 02:51:17.223 blog_,benchmark_2 (USDJPY,H1)    Work Items (1) Iterations(100000000)
2023.05.03 02:51:17.223 blog_benchmark_2 (USDJPY,H1)    Work Groups (1)
2023.05.03 02:51:17.223 blog_benchmark_2 (USDJPY,H1)    Milliseconds (16)

خوب دوباره کرنل را به این تغییر داد:

    string kernel="__kernel void memtests(__global int* group_id,"
                                         "int iterations){"
                                         "double sum=(double)0.0;"
                                         "double inc=(double)-2.6;"
                                         "double step=(double)0.01;"
                                         "double radi=(double)0.0;"
                                         "double rstep=(double)0.00001;"
                                         "for(int i=0;i<iterations;i++){"
                                         "sum=((double)tanh(inc))*((double)sin(radi));"
                                         "inc+=step;"
                                         "radi+=rstep;"
                                         "if(inc>2.6&&step>0.0){step=(double)-0.01;}"
                                         "if(inc<-2.6&&step<0.0){step=(double)0.01;}"
                                         "if(radi>6.28&&rstep>0.0){step=(double)-0.00001;}"
                                         "if(radi<0.0&&rstep<0.0){step=(double)0.00001;}"                                         
                                         "}"
                                         "group_id[get_global_id(0)]=get_group_id(0);}";

به 10000 تکرار کاهش یافته است، به طور کلی زیاده روی نکنید زیرا ممکن است پاسخ های کش دریافت کنید و در لحظه ای که محاسبه را تغییر می دهید، فرض کنید شکایت gpu خود را خواهید شنید.

بنابراین 10000 تکرار:

2023.05.03 03:07:20.865 blog_benchmark_2 (USDJPY,H1)    Work Items (1) Iterations(10000)
2023.05.03 03:07:20.865 blog_benchmark_2 (USDJPY,H1)    Work Groups (1)
2023.05.03 03:07:20.865 blog_benchmark_2 (USDJPY,H1)    Milliseconds (31)

اجازه می دهد تا 100K افزایش یابد

من دوباره 31 میلی‌ثانیه می‌گیرم، مطمئنم اما این بار بعد از راه‌اندازی مجدد آن را نیز تست می‌کنم.

بنابراین 1 مورد 100 هزار تکرار بدون راه اندازی مجدد کامپیوتر.

2023.05.03 03:11:24.106 blog_benchmark_2 (USDJPY,H1)    Work Items (1) Iterations(100000)
2023.05.03 03:11:24.106 blog_benchmark_2 (USDJPY,H1)    Work Groups (1)
2023.05.03 03:11:24.106 blog_benchmark_2 (USDJPY,H1)    Milliseconds (47)

Ok و 1 مورد 100K تکرار کامپیوتر را راه اندازی مجدد کنید.

2023.05.03 03:13:26.448 blog_benchmark_2 (USDJPY,H1)    Work Items (1) Iterations(100000)
2023.05.03 03:13:26.448 blog_benchmark_2 (USDJPY,H1)    Work Groups (1)
2023.05.03 03:13:26.448 blog_benchmark_2 (USDJPY,H1)    Milliseconds (47)

در همان زمان، بسیار خوب، ما با این هسته خوب هستیم و ببینیم چه کسی تصمیم گرفته است ظاهر شود:

2023.05.03 03:13:26.448 blog_benchmark_2 (USDJPY,H1)    Kernel private mem (40)

اندازه گیری حافظه خصوصی برای اولین بار صفر نشد. باشه پس 40 چیه؟ 40 بایت در هر مورد؟

بیایید هسته را اندازه گیری کنیم، غیر جهانی و غیر محلی و غیر آرگومان را می شماریم تا:

آره 5 5 در 8 40 بایتش دو برابر میشه . خوب پس این هم کار می کند. عالی.

حافظه خصوصی فوق‌العاده سریع است، شما نگران آن نباشید. اینکه چگونه حد را اندازه می‌گیریم، این سوال بعدی است.

بسیار خوب، آیا اکنون می تواند 1 میلیون تکرار طول بکشد؟

2023.05.03 03:18:41.921 blog_benchmark_2 (USDJPY,H1)    Work Items (1) Iterations(1000000)
2023.05.03 03:18:41.921 blog_benchmark_2 (USDJPY,H1)    Work Groups (1)
2023.05.03 03:18:41.921 blog_benchmark_2 (USDJPY,H1)    Milliseconds (203)

بسیار خوب، اما آیا این عدد چیزی است که به ما امکان می دهد “تا” را تشخیص دهیم؟

بیایید ببینیم، بیایید شروع به افزایش موارد کنیم، من لاگ ها را نمی چسبانم و فقط آنها را لیست می کنم

موارد#	گروه ها#	ام‌اس
1	1	203
100	1	203
256	1	188
512	2	297
1024	4	578
2048	8	1125
4096	1	2235

هوم اون پایین چی شد؟ از 8 گروه به 1 گروه رفتیم و صفحه فلش شد

بنابراین 1024 حداکثر آیتم در هر واحد محاسباتی است و 256 حداکثر اندازه گروه است که نشان دهنده 4 1024 فرآیند همزمان و 4 گروه است که می توانند در یک حافظه کار کنند؟ (برای 1 بعدی با فرض اینکه خودمان هیچ تقسیم بندی انجام نمی دهیم)

باید به خاطر داشته باشیم که به آن اجازه می‌دهیم در مورد تقسیم‌بندی‌های این معیار تصمیم بگیرد و 1 گروه را انتخاب کرد، یا در اینجا هیچ کاری موازی انجام نداد و فایلی که ما خارج می‌کنیم این را تأیید می‌کند.

GROUP.ID[4095]=0

آخرین ورودی 0 بود. بنابراین چیزی محاسبه نکرد و احتمالاً اشکال داشت.

با نگاهی به شکل های بالا، فکر نمی کنم همزمانی وجود داشته باشد، اگر 4 گروه به صورت موازی اجرا شوند یا 2 گروه یا 8 گروه، زمان کم و بیش یکسانی را با 1 گروه خواهیم دید درست است؟ بنابراین چیزی که به صورت موازی اجرا می شود، 256 مورد است.

هوم، خب حالا بیایید پارامتر دیگری را در تست خود معرفی کنیم و آن را معیار 3 بنامیم. پارامتر ابعاد محلی در فراخوانی اجرا

bool  CLExecute(
   int          kernel,                   
   uint         work_dim,                 
   const uint&  global_work_offset[],     
   const uint&  global_work_size[],       
   const uint&  local_work_size[]         
   );

اگر این را روی 32 تنظیم کنم چه می شود؟ می توانیم 8 زیر گروه دریافت کنیم.

آن 8 زیر گروه به طور همزمان اجرا می شوند اما اگر i 10 زیر گروه (یعنی 2 گروه) را مستقر کنیم، زمان اجرای بالاتری خواهیم داشت؟

فکر می‌کنم نیازی به آزمایش این موضوع نیست، اما اجازه دهید تأیید کنیم.

فهمیدم دستگاهی دارم که می تواند 256 مورد را همزمان محاسبه کند و به همین دلیل است که عدد صحیح

CL_MAX_WORK_GROUP_SIZE این عدد را نیز برمی گرداند اندازه گروه می تواند 256 باشد زیرا 256 که همزمان اجرا می شوند می توانند حافظه محلی را به اشتراک بگذارند.

آیا می تواند به همین سادگی باشد؟

اگر پارامتر local_work_size را روی 32 تنظیم کنم (منطبق با warp/wavefront) برای # آیتم های مختلف، بیایید نحوه تقسیم گروه ها را آزمایش کنیم.

    uint work_items_per_group[]={32};
    long msStarted=GetTickCount();
    CLExecute(ker,1,offsets,works,work_items_per_group);

فکر نمی کنم نیازی به تغییر چیز دیگری داشته باشم.

بیایید تعداد گروه هایی را که ایجاد می کند ببینیم

(توجه داشته باشید که 4096 واقعاً خراب شد زیرا در اجرای بعدی با خطای زمینه مواجه شدم که نیاز به راه اندازی مجدد داشت)

موارد	گروه ها
1	هیچکدام، البته خطا درست است
32	1 (203 میلی‌ثانیه)
64	2 (203 میلی‌ثانیه)
128	4 (203 میلی‌ثانیه)
256	8 (203 میلی‌ثانیه)
512	16 (297 میلی‌ثانیه)
1024	32 (578 میلی‌ثانیه)
2048	64 (1125 میلی‌ثانیه)
4096	1 (2234 میلی‌ثانیه)

بسیار خب، بدیهی است که

long  CLGetInfoInteger(ctx,CL_DEVICE_MAX_WORK_GROUP_SIZE);

که 1024 را در دستگاه من برمی گرداند، به حداکثر موارد موازی اشاره دارد که می توانند همزمان اجرا شوند.

این بدان معناست که اگر اندازه محلی را روی 1024 تنظیم کنید، 1024 مورد را به صورت موازی اجرا نمی‌کند، آن را به بسته‌های 1024 تقسیم می‌کند و احتمالاً از کار می‌افتد و کند می‌شود، اما اگر آن را روی 1 تنظیم کنید، هر آیتم کاری را به عنوان گروه خودش در نظر می‌گیرد. ، درست ؟

بنابراین برای این تست که نیازی به گروه ندارد، اگر اندازه محلی را روی 1 تنظیم کنم، اگر آیتم ها روی 1024 مرده باشند، سریع ترین اجرا را دریافت می کنم؟

نه، من دوباره مشکل دارم. اگر کل اقلام 256 باشد چه؟

نه , اشکال … ermm مضرب 2 شاید ?

نه، مشکل… پس دوباره چه اتفاقی می افتد؟

چه می شود اگر من 8 مورد و 1 محلی داشته باشم، بنابراین 1 مورد در هر گروه 8 گروه اساساً

هیچ اشکالی وجود ندارد، بنابراین محدودیتی در تعداد گروه ها وجود دارد اما ربطی به این ندارد:

CL_DEVICE_MAX_WORK_GROUP_SIZE

و راه برای کشف محدودیت گروه این است که به دنبال کاهش از افزایش # گروه به 1 و به دنبال آن مشکل صفحه نمایش باشید ???

اما 2048 مورد کار را نیز تحمل کرد، بنابراین، من فکر می کنم که ما شاهد عدم ارتباط بین OpenCL و سخت افزار هستیم.

آیا در هر بعد 8 است، پس 8^3؟ 512؟ یا ریشه مکعب 1024 (حداکثر موارد ? ) hmmm

این ~ 10 است پس حداکثر 10 گروه؟ (به ازای بعد من فرض می کنم)

چگونه می دانستیم؟ 320 مورد با 32 مورد محلی احتمالاً کار خواهد کرد

خوب بود اما 64 مورد هم کار کرد 2048 مورد با 32 محلی … به هر حال .. من خراب شدم یک بار دیگر آن را برمی دارم.

من معیارها را ضمیمه می کنم …

همچنین چیزی برای گرفتن از همه این تست ها وجود دارد، و این ارزش است

CLGetInfoInteger(ker,CL_KERNEL_WORK_GROUP_SIZE);

داره بهت میگه چندتا نمونه های هسته می تواند به صورت موازی اجرا شود این دستگاه (یا به ازای هر واحد محاسباتی اما می توان آن را با cpu آزمایش کرد)