OpenCL :: Exploring the 1st Dimension (سری ناقص) – سایر – 2 مه 2023

سلام.

این آزمایش یک فضای “کار” یک بعدی را با دستورات اصلی mql5 OpenCL در دسترس استقرار می دهد.

وظیفه کشف نحوه نگاشت حافظه محلی به گروه های کاری است.

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

نمی‌دانم معنایی دارد یا نه، اما این چیزی است که سعی می‌کنم پیدا کنم.

همچنین باید همیشه و در هر لحظه به خاطر داشته باشید که این ممکن است یک واکنش خاص این سخت افزار باشد، بنابراین باید انعطاف پذیری برای دیگران (شما) وجود داشته باشد تا خودتان آن را آزمایش کنید.

OpenCL دارای اندیس های داخلی است که می توان آنها را از هسته فراخوانی کرد و آنها را ارائه می کند

• شناسه جهانی مورد کار، برای یک بعد
• شناسه محلی مورد کار، برای یک بعد
• شناسه گروه مورد کار، برای یک بعد

    string kernel="__kernel void memtests(__global int* global_id,"
                                         "__global int* local_id,"
                                         "__global int* group_id){"
                                         "global_id[get_global_id(0)]=get_global_id(0);"
                                         "local_id[get_global_id(0)]=get_local_id(0);"
                                         "group_id[get_global_id(0)]=get_group_id(0);}";

کاری که این کار انجام می دهد این است که 3 آرایه حافظه جهانی عدد صحیح global_id , local_id , group_id را دریافت می کند و آنها را با شاخص مربوطه در موقعیت جهانی پر می کند. به عنوان مثال اگر 10 پرتقال در 2 کیسه داشته باشیم، شاخص کیسه را به شاخص آرایه خطی پرتقال ها اختصاص می دهیم.
ما می گوییم، نارنجی[0] در Bag0 و نارنجی است[9] در Bag1 است، ما از شاخص پرتقال در کیسه استفاده نمی کنیم (نارنجی[0] در Bag0 و نارنجی است[4] در Bag1 است) که چیزی در مورد نحوه چیدمان پرتقال ها به ما نمی گوید!

بنابراین از get_global_id(0) برای ذخیره همه شناسه ها استفاده می شود.

ما همه کارهای آماده سازی کسل کننده و آزاردهنده OpenCl را انجام می دهیم، نحوه اجرای برنامه، زمانی که فراخوانی می شود، زمینه های هسته را بافر می کند و یک پارامتر برای تعداد آیتم ها یا پرتقال هایی که باید برای آزمایش پرتاب شوند، ارائه می دهیم.

این کد به نظر می رسد:

#property version   "1.00"

int OnInit()
  {

  EventSetMillisecondTimer(33);

   return(INIT_SUCCEEDED);
  }

void OnDeinit(const int reason)
  {

   
  }
void OnTimer(){
  EventKillTimer();
  int ctx=CLContextCreate(CL_USE_GPU_DOUBLE_ONLY);
  if(ctx!=INVALID_HANDLE){
    string kernel="__kernel void memtests(__global int* global_id,"
                                         "__global int* local_id,"
                                         "__global int* group_id){"
                                         "global_id[get_global_id(0)]=get_global_id(0);"
                                         "local_id[get_global_id(0)]=get_local_id(0);"
                                         "group_id[get_global_id(0)]=get_group_id(0);}";
    string errors="";
    int prg=CLProgramCreate(ctx,kernel,errors);
    if(prg!=INVALID_HANDLE){
    ResetLastError();
    int ker=CLKernelCreate(prg,"memtests");
    if(ker!=INVALID_HANDLE){
    int items=2560;
    int global_ids[];ArrayResize(global_ids,items,0);
    ArrayFill(global_ids,0,items,0);
    int local_ids[];ArrayResize(local_ids,items,0);
    ArrayFill(local_ids,0,items,0);    
    int group_ids[];ArrayResize(group_ids,items,0);
    ArrayFill(group_ids,0,items,0);        
    int global_id_handle=CLBufferCreate(ctx,items*4,CL_MEM_WRITE_ONLY);
    int local_id_handle=CLBufferCreate(ctx,items*4,CL_MEM_WRITE_ONLY);
    int group_id_handle=CLBufferCreate(ctx,items*4,CL_MEM_WRITE_ONLY);
    CLSetKernelArgMem(ker,0,global_id_handle);
    CLSetKernelArgMem(ker,1,local_id_handle);
    CLSetKernelArgMem(ker,2,group_id_handle);
    uint offsets[]={0};
    uint works[]={items};
    CLExecute(ker,1,offsets,works);
    while(CLExecutionStatus(ker)!=CL_COMPLETE){Sleep(10);}
    Print("Kernel finished");
    CLBufferRead(global_id_handle,global_ids,0,0,items);
    CLBufferRead(local_id_handle,local_ids,0,0,items);
    CLBufferRead(group_id_handle,group_ids,0,0,items);
    
    int f=FileOpen("OCL\log.txt",FILE_WRITE|FILE_TXT);
    for(int i=0;i<items;i++){
       FileWriteString(f,"GLOBAL.ID["+IntegerToString(i)+"]="+IntegerToString(global_ids[i])+" : LOCAL.ID["+IntegerToString(i)+"]="+IntegerToString(local_ids[i])+" : GROUP.ID["+IntegerToString(i)+"]="+IntegerToString(group_ids[i])+"n");
       }
    FileClose(f);
    int kernel_local_mem_size=CLGetInfoInteger(ker,CL_KERNEL_LOCAL_MEM_SIZE);
    int kernel_private_mem_size=CLGetInfoInteger(ker,CL_KERNEL_PRIVATE_MEM_SIZE);
    int kernel_work_group_size=CLGetInfoInteger(ker,CL_KERNEL_WORK_GROUP_SIZE);
    Print("Kernel local mem ("+kernel_local_mem_size+")");
    Print("Kernel private mem ("+kernel_private_mem_size+")");
    Print("Kernel work group size ("+kernel_work_group_size+")");

    CLKernelFree(ker);
    CLBufferFree(global_id_handle);
    CLBufferFree(local_id_handle);
    CLBufferFree(group_id_handle);
    }else{Print("Cannot create kernel");}
    CLProgramFree(prg);
    }else{Alert(errors);}
    CLContextFree(ctx);
    }
  else{
    Print("Cannot create ctx");
    }
  }
void OnTick()
  {

   
  }

پارامتر آیتم ها و صادرات شاخص هایی که دریافت کرده ایم به یک فایل برجسته شده است.

همچنین این کد 3 مقدار را چاپ می کند:

1. اندازه حافظه محلی هسته (من 1 میگیرم)
2. اندازه حافظه خصوصی هسته (0 می گیرم)
3. اندازه گروه کاری هسته (من 256 می گیرم)

پس بیایید با پرتاب 100 مورد شروع کنیم.

GLOBAL.ID[0]=0 : LOCAL.ID[0]=0 : GROUP.ID[0]=0
GLOBAL.ID[1]=1 : LOCAL.ID[1]=1 : GROUP.ID[1]=0
GLOBAL.ID[2]=2 : LOCAL.ID[2]=2 : GROUP.ID[2]=0
GLOBAL.ID[3]=3 : LOCAL.ID[3]=3 : GROUP.ID[3]=0
...
GLOBAL.ID[96]=96 : LOCAL.ID[96]=96 : GROUP.ID[96]=0
GLOBAL.ID[97]=97 : LOCAL.ID[97]=97 : GROUP.ID[97]=0
GLOBAL.ID[98]=98 : LOCAL.ID[98]=98 : GROUP.ID[98]=0
GLOBAL.ID[99]=99 : LOCAL.ID[99]=99 : GROUP.ID[99]=0

خروجی فایل به این شکل است (در بالا) و می بینید که به هیچ وجه بار کار را تقسیم نکرده است.

اولین چیزی که به ذهن می رسد این است که “آیا مکانیزم تصمیم گیری داخلی دارد؟” برای هر دو دستور wrapper و mql5 بومی زمانی که دستورات OpenCL را فراخوانی می کنند. چگونه خواهیم دانست؟ بیایید هسته را کمی «سنگین‌تر» کنیم و همچنین یک مقدار شناور استخراج کنیم.

بیایید محدوده ای از مقادیر را از 2.6- تا 2.6 ایجاد کنیم و از آن بخواهیم tanh آن x برابر را محاسبه کند و همچنین آرایه را در لحظه تغییر دهد.

بنابراین ما یک نشانگر شناور جهانی به آرگومان های هسته اضافه می کنیم

__global float* _tangent,

سپس یک آرایه دوتایی در برنامه خود ایجاد می کنیم و آن را با مقادیر تصادفی در محدوده -2.6 تا 2.6 پر می کنیم.

    float tangents[];ArrayResize(tangents,items,0);
    float range=5.2;
    for(int i=0;i<ArraySize(tangents);i++){
       float r=(((float)MathRand())/((float)32767.0)*range)-2.6;
       tangents[i]=r;
       }   

سپس یک دسته بافر برای openCL ایجاد کنید، در این حالت حافظه خواندن نوشتن است نه فقط نوشتن. (خواندن نوشتن برای دستگاه نه میزبان (ما))

int tangents_handle=CLBufferCreate(ctx,items*4,CL_MEM_READ_WRITE);

و سپس بافر را به آرگومان های هسته پیوند دهید

CLSetKernelArgMem(ker,3,tangents_handle);

aaand (من 2 ساعت را در اینجا تلف کردم زیرا فراموش کردم بافر را دوباره بخوانم 🤣) بنابراین، فراموش نکنید که وقتی داده‌ها را می‌خواهید، بافر خواندن را فراخوانی کنید.

CLBufferRead(tangents_handle,tangents,0,0,items);

aa و در انتها بافر رایگان اضافه کنید در غیر این صورت به نظر می رسد که در حال ذخیره (هوشمندانه) مقادیر است

CLBufferFree(tangents_handle);

جالب است، اکنون باید با ارائه یک عدد تکرار به عنوان آرگومان، کار را کمی دشوارتر کنیم.
این محاسبه tanh را حلقه می کند و برای هر نتیجه ما tanh شناور / تکرار مماس را جمع می کنیم.
مرتب

در این مورد باید یک آرگومان در هسته اضافه کنیم، آرگومان را به هسته پیوند داده و مقدار را تنظیم کنیم:

"int iterations){"
int iterations=100;
CLSetKernelArg(ker,4,iterations);

و محاسبات هسته، بیایید فهرست OpenCL را جستجو کنیم

                                         "float sum=(float)0.0;"
                                         "float of=(float)_tangent[get_global_id(0)];"
                                         "for(int i=0;i<iterations;i++){"
                                         "sum+=((float)tanh(of-sum))/((float)iterations);"
                                         "}"
                                         "_tangent[get_global_id(0)]=sum;"

• مجموع را روی 0.0 تنظیم کنید.
• مقدار اولیه را در متغیر mem خصوصی “of” تنظیم کنید
• حلقه به تکرار
• tanh “از” منهای مجموع تا کنون تقسیم بر تکرارها را جمع کنید.
• آرایه مماس را پر کنید

GLOBAL.ID[0]=0 : LOCAL.ID[0]=0 : GROUP.ID[0]=0 : T[0]=0.7702
GLOBAL.ID[1]=1 : LOCAL.ID[1]=1 : GROUP.ID[1]=0 : T[1]=0.0282
GLOBAL.ID[2]=2 : LOCAL.ID[2]=2 : GROUP.ID[2]=0 : T[2]=0.9934
GLOBAL.ID[3]=3 : LOCAL.ID[3]=3 : GROUP.ID[3]=0 : T[3]=2.2652
GLOBAL.ID[4]=4 : LOCAL.ID[4]=4 : GROUP.ID[4]=0 : T[4]=-2.2026
...
GLOBAL.ID[96]=96 : LOCAL.ID[96]=96 : GROUP.ID[96]=0 : T[96]=-1.7437
GLOBAL.ID[97]=97 : LOCAL.ID[97]=97 : GROUP.ID[97]=0 : T[97]=-1.1011
GLOBAL.ID[98]=98 : LOCAL.ID[98]=98 : GROUP.ID[98]=0 : T[98]=0.4125
GLOBAL.ID[99]=99 : LOCAL.ID[99]=99 : GROUP.ID[99]=0 : T[99]=1.8560

بسیار خوب به نظر من کار می کند و ما هنوز هیچ شکافی نداریم. بیایید تکرارها را به 10000 برسانیم اما موارد 100 باقی می مانند.

هنوز یک گروه

GLOBAL.ID[99]=99 : LOCAL.ID[99]=99 : GROUP.ID[99]=0 : T[99]=0.0905

بیایید 10 میلیون تکرار 100 مورد، که ممکن است برای دقت شناور مشکل ساز باشد؟ اجازه بدید ببینم

نه، هنوز یک گروه است، اگرچه gpu کمی تاخیر داشت.

GLOBAL.ID[99]=99 : LOCAL.ID[99]=99 : GROUP.ID[99]=0 : T[99]=-2.4797

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

خوب، حالا بیایید تکرارها را به 1000 برگردانیم و آزمایش را با مقادیر مختلف شروع کنیم.

بیایید 200 مورد را به آن بریزیم

هنوز یک گروه

GLOBAL.ID[0]=0 : LOCAL.ID[0]=0 : GROUP.ID[0]=0 : T[0]=1.5756
GLOBAL.ID[1]=1 : LOCAL.ID[1]=1 : GROUP.ID[1]=0 : T[1]=-1.1957
GLOBAL.ID[2]=2 : LOCAL.ID[2]=2 : GROUP.ID[2]=0 : T[2]=0.6411
...
GLOBAL.ID[198]=198 : LOCAL.ID[198]=198 : GROUP.ID[198]=0 : T[198]=0.5839
GLOBAL.ID[199]=199 : LOCAL.ID[199]=199 : GROUP.ID[199]=0 : T[199]=-1.5742

حالا بیایید 257 مورد را پرتاب کنیم!

GLOBAL.ID[0]=0 : LOCAL.ID[0]=0 : GROUP.ID[0]=0 : T[0]=2.0035
GLOBAL.ID[1]=1 : LOCAL.ID[1]=1 : GROUP.ID[1]=0 : T[1]=-0.0069
GLOBAL.ID[2]=2 : LOCAL.ID[2]=2 : GROUP.ID[2]=0 : T[2]=-0.8145
GLOBAL.ID[3]=3 : LOCAL.ID[3]=3 : GROUP.ID[3]=0 : T[3]=1.7880
...
GLOBAL.ID[255]=255 : LOCAL.ID[255]=255 : GROUP.ID[255]=0 : T[255]=0.2042
GLOBAL.ID[256]=256 : LOCAL.ID[256]=256 : GROUP.ID[256]=0 : T[256]=1.7910

هنوز شکافی وجود ندارد.

258 چطور؟ ما اولین تقسیم خود را دریافت می کنیم!

GLOBAL.ID[0]=0 : LOCAL.ID[0]=0 : GROUP.ID[0]=0 : T[0]=-1.2919
GLOBAL.ID[1]=1 : LOCAL.ID[1]=1 : GROUP.ID[1]=0 : T[1]=-1.2212
...
GLOBAL.ID[127]=127 : LOCAL.ID[127]=127 : GROUP.ID[127]=0 : T[127]=-1.2222
GLOBAL.ID[128]=128 : LOCAL.ID[128]=128 : GROUP.ID[128]=0 : T[128]=1.9752
GLOBAL.ID[129]=129 : LOCAL.ID[129]=0 : GROUP.ID[129]=1 : T[129]=1.0197
...
GLOBAL.ID[255]=255 : LOCAL.ID[255]=126 : GROUP.ID[255]=1 : T[255]=1.9462
GLOBAL.ID[256]=256 : LOCAL.ID[256]=127 : GROUP.ID[256]=1 : T[256]=-1.9560
GLOBAL.ID[257]=257 : LOCAL.ID[257]=128 : GROUP.ID[257]=1 : T[257]=-0.9829

بار کار را به دو گروه تقسیم کرد!

بسیار خوب، پس ما باید معیار تقسیم را پیدا کنیم، بیایید تعدادی اعدادی را که 512، 1024، 2048 و غیره دوست دارد بریزیم و نتایج را ثبت کنیم.

512: دوباره به دو گروه تقسیم می شود:

GLOBAL.ID[0]=0 : LOCAL.ID[0]=0 : GROUP.ID[0]=0 : T[0]=-0.3564
GLOBAL.ID[1]=1 : LOCAL.ID[1]=1 : GROUP.ID[1]=0 : T[1]=2.3337
...
GLOBAL.ID[255]=255 : LOCAL.ID[255]=255 : GROUP.ID[255]=0 : T[255]=-2.4480
GLOBAL.ID[256]=256 : LOCAL.ID[256]=0 : GROUP.ID[256]=1 : T[256]=2.3620
...
GLOBAL.ID[510]=510 : LOCAL.ID[510]=254 : GROUP.ID[510]=1 : T[510]=-2.2709
GLOBAL.ID[511]=511 : LOCAL.ID[511]=255 : GROUP.ID[511]=1 : T[511]=-0.3056

همچنین توجه داشته باشید که شناسه محلی چقدر راحت است و ایندکس را در هر گروه به ما می گوید!

1024 : آها الان به 4 گروه تقسیم شد ! بنابراین حداکثر اندازه گروه برای این دستگاه 256 مورد است؟

GLOBAL.ID[0]=0 : LOCAL.ID[0]=0 : GROUP.ID[0]=0 : T[0]=-0.7910
GLOBAL.ID[1]=1 : LOCAL.ID[1]=1 : GROUP.ID[1]=0 : T[1]=-0.7287
...
GLOBAL.ID[255]=255 : LOCAL.ID[255]=255 : GROUP.ID[255]=0 : T[255]=0.2203
GLOBAL.ID[256]=256 : LOCAL.ID[256]=0 : GROUP.ID[256]=1 : T[256]=1.4999
..
GLOBAL.ID[511]=511 : LOCAL.ID[511]=255 : GROUP.ID[511]=1 : T[511]=0.1762
GLOBAL.ID[512]=512 : LOCAL.ID[512]=0 : GROUP.ID[512]=2 : T[512]=-0.0072
...
GLOBAL.ID[767]=767 : LOCAL.ID[767]=255 : GROUP.ID[767]=2 : T[767]=-2.0688
GLOBAL.ID[768]=768 : LOCAL.ID[768]=0 : GROUP.ID[768]=3 : T[768]=-2.0622
...
GLOBAL.ID[1022]=1022 : LOCAL.ID[1022]=254 : GROUP.ID[1022]=3 : T[1022]=2.2044
GLOBAL.ID[1023]=1023 : LOCAL.ID[1023]=255 : GROUP.ID[1023]=3 : T[1023]=-0.6644

بیایید به آن مقدار 768 بدهیم، آیا مقادیر حتی گروهی را دوست دارد؟ (به خاطر تعداد هسته هایش؟)

GLOBAL.ID[0]=0 : LOCAL.ID[0]=0 : GROUP.ID[0]=0 : T[0]=1.8908
...
GLOBAL.ID[255]=255 : LOCAL.ID[255]=255 : GROUP.ID[255]=0 : T[255]=0.0147
GLOBAL.ID[256]=256 : LOCAL.ID[256]=0 : GROUP.ID[256]=1 : T[256]=-1.5271
...
GLOBAL.ID[511]=511 : LOCAL.ID[511]=255 : GROUP.ID[511]=1 : T[511]=2.3339
GLOBAL.ID[512]=512 : LOCAL.ID[512]=0 : GROUP.ID[512]=2 : T[512]=-0.8512
...
GLOBAL.ID[767]=767 : LOCAL.ID[767]=255 : GROUP.ID[767]=2 : T[767]=-0.1783

خیر , مشکلی وجود ندارد 3 گروه 256 موردی هر کدام . باشه

آیا می توانیم در اینجا به نتیجه ای برسیم؟ ارزش است:

int kernel_work_group_size=CLGetInfoInteger(ker,CL_KERNEL_WORK_GROUP_SIZE);

به ما اطلاع دهید از حداکثر موارد کاری که یک گروه کاری می تواند داشته باشد، زیرا، این همان کاری است که gpu زمانی که هیچ دستورالعملی ندارد، به تنهایی انجام می دهد؟

و اگر این درست است، در مورد این مقدار در اینجا چطور؟

int device_max_work_group_size=CLGetInfoInteger(ctx,CL_DEVICE_MAX_WORK_GROUP_SIZE);

در واقع اجازه دهید آن را در کد اضافه کنم و ببینم چه چیزی را برمی گرداند:

1024 می گوید . 1024 گروه کاری یا 1024 مورد در گروه های کاری در مجموع؟

خب بیایید بفهمیم

اگر بخواهیم 1025 گروه کاری (برای این دستگاه) بگیریم، به 1025*256 آیتم نیاز داریم، یعنی 262400 مورد.

اجازه بدید ببینم . این یک فایل a** بزرگ است … اما خوشبختانه ما فقط به ردیف آخر نیاز داریم.

خوب ما 1025 گروه کاری دریافت می کنیم … باشه

GLOBAL.ID[262399]=262399 : LOCAL.ID[262399]=255 : GROUP.ID[262399]=1024 : T[262399]=-0.1899

آیا ما به این موضوع اشتباه نزدیک شده ایم؟ آیا لازم نیست 1025 گروه کاری به صورت موازی اجرا شوند درست است؟

اگر این نشان دهد که 1024 گروه کاری می توانند به طور همزمان کار کنند چه می شود زیرا این دستگاه دارای 1024 واحد محاسباتی است؟

حتی اگر 1 واحد محاسباتی را گزارش می دهد (این دستورات اطلاعاتی را نیز به آن اضافه می کنم تا بتوانید مقایسه کنید) و حتی اگر کیت ابزار cuda 192 هسته مختلف و 32 تاب را گزارش می دهد.

این احساس می شود:

• انویدیا نصف ** آن را ویرایش کرد
• OpenCL را نصف ** آن را ویرایش کنید
• Mql5 نصف آن را انجام داد

و ما مانده ایم که نارنجک را در دست گرفته ایم تا بفهمیم خوشه انگور است یا نه.

اما ما نصف چیزهایی را دوست نداریم که یک** یا بدون** انجام می دهیم!

بنابراین، چگونه می توانیم بگوییم که f* در حال وقوع است؟

ما باید زمان اجرا را اندازه گیری کنیم، اما چه؟

ما باید بریدگی قابل توجهی را در زمان اجرا کشف کنیم که نشان می‌دهد هسته‌ها در حال تعویض گروه‌ها هستند، و برای پیچیده‌تر کردن کارها، این باید همه‌کاره باشد، بنابراین اگر آن را اجرا کنید باید نشانه‌ای دریافت کنید که در مقایسه با نشانه‌های من می‌توانیم نتیجه‌گیری کنیم یا به آن نزدیک شویم. فعالیت در دستگاه
نه دستگاه من یا دستگاه شما، بلکه به طور کلی. (در صورت امکان)

بیایید همچنین توجه داشته باشیم که اگر یک سازنده به طور کامل همکاری نمی کند، نتیجه دیگری جز 1/2*a**-کردن آن وجود ندارد. پس انصافا اینجا mq و chronos تقصیری نداره . بنابراین از این نظر باید آن را نیز به 0.5*a** آماده کنم. 😂

خوب حالا باید یک پرانتز غول پیکر باز کنیم و متأسفانه کار دیگری انجام دهیم.
معیار کنونی از این نظر مشکل دارد که از حافظه زیاد استفاده می کند.
اگر می‌خواهیم «قطع» هسته‌ها ظاهر شود، باید از «محاسبات» بیشتر از «واکشی» استفاده کنیم.
این تست همچنین زمانی اجرا می‌شود که بخواهیم دائماً اجرا شود، مشکل دوم ما اگر قطع باشد
نزدیک به راه اندازی مجدد حلقه ما متوجه آن نخواهیم شد!

حالا شما به چیزی که من فکر می کنم فکر می کنید: “چرا این را در سطوح پایین تر، nvidia یا amd حل نکنیم”؟ و احتمالاً پاسخ این است که “ما این همه پول R+D را خرج نکردیم تا Khronos بیرون بیاید و از آن سود ببریم”، یا، برای ساده کردن آن، احتمالا “از Cuda استفاده کنید، یا، از Hpi استفاده کنید”.

کدهای کمتر یا کمتر تایپ کردن با تایپ بیشتر، علاوه بر این، شکایت کردن شما را ثروتمند می کند تنها در صورتی که وکیل یا سیاستمدار باشید… 😂

بیایید دنبالش برویم حتی اگر شکست بخورد

معیار 2: برش زمان اجرا با اندازه گروه

class kernel_info{
      public:
  int offset;
  int handle;
ulong start_microSeconds;
ulong end_microSeconds;
      kernel_info(void){reset();}
     ~kernel_info(void){reset();}
 void reset(){
      offset=-1;
      handle=INVALID_HANDLE;
      start_microSeconds=0;
      end_microSeconds=0;
      }
 void setup(int _hndl,ulong _start,int _offset){
      handle=_hndl;
      start_microSeconds=_start;
      offset=_offset;
      }
 void stop(ulong _end){
      end_microSeconds=_end;
      }
};

kernel_info KERNELS[];

int OnInit()
  {

  EventSetMillisecondTimer(33);

   return(INIT_SUCCEEDED);
  }

void OnDeinit(const int reason)
  {

   
  }
void OnTimer(){
  EventKillTimer();
  int ctx=CLContextCreate(CL_USE_GPU_DOUBLE_ONLY);
  if(ctx!=INVALID_HANDLE){
    string kernel="__kernel void bench(__global int* global_id,"
                                      "__global int* local_id,"
                                      "__global int* group_id,"
                                      "__global float* _tangent,"
                                      "int iterations){"
                                      "float sum=(float)0.0;"
                                      "float of=(float)_tangent[get_global_id(0)];"
                                      "for(int i=0;i<iterations;i++){"
                                      "sum+=((float)tanh(of-sum))/((float)iterations);"
                                      "}"
                                      "sum=(float)0.12345;"
                                      "_tangent[get_global_id(0)]=0.123;"
                                      "global_id[get_global_id(0)]=get_global_id(0);"
                                      "local_id[get_global_id(0)]=get_local_id(0);"
                                      "group_id[get_global_id(0)]=get_group_id(0);}";
    string errors="";
    int prg=CLProgramCreate(ctx,kernel,errors);
    if(prg!=INVALID_HANDLE){
    ResetLastError();
    
      int kernels_to_deploy=5;
    
      int iterations=1000;
      ArrayResize(KERNELS,kernels_to_deploy,0);
      bool deployed=true;
      for(int i=0;i<kernels_to_deploy;i++){
         KERNELS[i].handle=CLKernelCreate(prg,"bench");
         if(KERNELS[i].handle==INVALID_HANDLE){deployed=false;}
         }
      
      if(deployed){
      Print("Deployed all kernels!");
      for(int i=0;i<kernels_to_deploy;i++){
         if(KERNELS[i].handle!=INVALID_HANDLE){Print("Kernel ["+i+"] Valid");}
         else{Print("Kernel ["+i+"] InValid");}
         }     
      }else{
      Print("Cannot deploy all kernels!");
      for(int i=0;i<kernels_to_deploy;i++){
         if(KERNELS[i].handle!=INVALID_HANDLE){Print("Kernel ["+i+"] Valid");}
         else{Print("Kernel ["+i+"] InValid");}
         }
      }
    
      for(int i=0;i<kernels_to_deploy;i++){
         if(KERNELS[i].handle!=INVALID_HANDLE){
           CLKernelFree(KERNELS[i].handle);
           }
         }  
    CLProgramFree(prg);
    }else{Alert(errors);}
    CLContextFree(ctx);
    }
  else{
    Print("Cannot create ctx");
    }
  }

  long timer_ended=GetTickCount();
  
    long diff=timer_ended-timer_started;
    if(timer_ended<timer_started){diff=UINT_MAX-timer_started+timer_ended;}
    Print("Time to load and unload "+IntegerToString(kernels_to_deploy)+" kernels = "+IntegerToString(diff)+"ms");
    ExpertRemove();

2023.05.02 20:11:41.352 blog_kernel_times_benchmark (USDJPY,H1) Time to load and unload 5 kernels = 94ms

2023.05.02 20:12:15.704 blog_kernel_times_benchmark (USDJPY,H1) Time to load and unload 50 kernels = 78ms

2023.05.02 20:13:16.359 blog_kernel_times_benchmark (USDJPY,H1) Time to load and unload 50 kernels = 94ms

عالی بدون تاخیر وجود دارد. بیایید عملیات انبوه 5000 هسته را انجام دهیم. اگر 78 میلی‌ثانیه دریافت کنیم، چراغ سبز برای معیار دریافت می‌کنیم.

(هر چند ممکن است تاخیر چاپی وجود داشته باشد:D پس فقط مواردی را که مستقر نشده اند چاپ کنید! و وضعیت)

2023.05.02 20:15:35.724 blog_kernel_times_benchmark (USDJPY,H1) Deployed all kernels!
2023.05.02 20:15:35.746 blog_kernel_times_benchmark (USDJPY,H1) Time to load and unload 5000 kernels = 94ms

عالی . پس بیا بریم !

اکنون باید محاسبات را در هسته به این صورت تبدیل کنیم: مقدار را دریافت کنید، مقدار دیوانه‌واری را روی آن محاسبه کنید، ارزش بدهید.

ما یک شرط داریم که باید برای تکرارها رعایت شود:

میلی‌ثانیه‌ای که برای اجرای یک واحد کاری طول می‌کشد باید قدر معینی بزرگ‌تر از فاصله تایمر باشد تا بتوانیم آن را اندازه‌گیری کنیم!

بنابراین اگر “معیار” خود تکرارهای بهینه را پیدا کند، تا زمانی که تکرارهایی که ارسال می‌کند در “زمان‌هایی” بزرگ‌تر از بازه زمانی باشد، وارد یک حلقه می‌شود.

اجازه دهید کاری را که در اینجا انجام می‌دهیم تکرار کنم: می‌خواهیم تغییر «تغییر» را در عناصر پردازش «تصور کنیم»، یعنی لحظه‌ای که واحدهای شلوغ تخلیه می‌شوند و داده‌های جدید را به طور مؤثر دریافت می‌کنند، یا من فکر می‌کنم که تعداد پردازش‌ها را به ما می‌دهد. عناصر (من حدس می زنم برابر باشد

CL_DEVICE_MAX_WORK_GROUP_SIZE

چی ؟

اجازه دهید توضیح دهم که فکر می کنم چه چیزی می تواند در اینجا اتفاق بیفتد:

ما یک واحد محاسباتی با 10 هسته فرعی پردازشی داریم (عناصر پردازش)

اگر 10 هسته را به طور همزمان برای اجرا ارسال کنم، به طور کلی 150 میلی ثانیه زمان اجرا دریافت می کنم، به این معنی که حداقل زمان ثبت شده من در آیتم های اطلاعات هسته که از حداکثر زمان ثبت شده کم می شود، 150 میلی ثانیه خواهد بود.

حالا اگر 11 هسته را اجرا کنم و 300 میلی ثانیه پیدا کنم، در مقابل اگر 170 میلی ثانیه پیدا کنم، چیزی درست به من می گوید.

درست ؟ من ممکن است اشتباه کنم

اجازه بدید ببینم ! این جالبه ! 😍

اما به قسمت 2 ادامه می یابد