OpenCL :: کاوش در بعد اول (بخش 2 معیار اشتباه و اصلاح) – سایر – 2 مه 2023

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

بسیار خوب، پس کاری که ما در اینجا انجام خواهیم داد مانند قبل است، اما این بار از یک نوع دوگانه استفاده کنید

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

    string kernel="__kernel void bench(__global double* _tangent,"
                                      "int iterations){"
                                      "double sum=(double)0.0;"
                                      "double of=(double)_tangent[get_global_id(0)];"
                                      "for(int i=0;i<iterations;i++){"
                                      "sum+=((double)tanh(of-sum))/((double)iterations);"
                                      "}"
                                      "_tangent[get_global_id(0)]=sum;}";

ما یک آرایه دوتایی با اندازه آن مطابق با تعداد هسته هایمان ایجاد می کنیم و همیشه مراقب هستیم که آن را با دقت نمونه برداری کنیم، بنابراین به 5 هسته برمی گردیم و این را مستقر می کنیم:

      
        double tangents[];
        ArrayResize(tangents,kernels_to_deploy,0);
        double range=5.2;
        for(int i=0;i<ArraySize(tangents);i++){
           double r=(((double)MathRand())/((double)32767.0)*range)-2.6;
           tangents[i]=r;
           }         
        int tangents_id=CLBufferCreate(ctx,kernels_to_deploy*8,CL_MEM_READ_WRITE);
        
          bool args_set=true;
          for(int i=0;i<kernels_to_deploy;i++){
             ResetLastError();
             if(!CLSetKernelArgMem(KERNELS[i].handle,0,tangents_id)){
               Print("Cannot assign buffer to kernel("+i+") error #"+IntegerToString(GetLastError()));
               args_set=false;
               }else{
               CLSetKernelArg(KERNELS[i].handle,1,iterations);
               }
             }
         if(args_set){
         Print("All arguments for all kernels set!");
         }else{
         Print("Cannot setup kernel args!");
         }

ما بافر را ایجاد می کنیم و سپس باید آن را به همه هسته ها متصل کنیم درست است؟ بیایید ببینیم آیا می توانیم این کار را انجام دهیم!

اما، فراموش نکنید که باید بافر را نیز تخلیه کنیم، پس این را بعد از حلقه unload اضافه کنید

(ما id مماس ها را از حلقه خارج می کنیم، اکنون در یک توزیع معمولی، این در داخل یک ساختار پیچیده و مدیریت می شود، اما ما در حال آزمایش هستیم، بنابراین نیازی به این نیست تا بتوان ماموریت ماه را فرود آورد!)

    CLBufferFree(tangents_id);
    CLProgramFree(prg);

ما آن را به آنجا اضافه می کنیم، آن را اجرا می کنیم و با یک بار اجرا 2 چیز را متوجه می شویم!

2023.05.02 20:49:49.762 blog_kernel_times_benchmark (USDJPY,H1) All arguments for all kernels set!
2023.05.02 20:49:49.776 blog_kernel_times_benchmark (USDJPY,H1) Time to load and unload 5 kernels = 94ms

خوب است، حالا بیایید 50 هسته انجام دهیم، می خواهیم تاخیرهای استقرار را اندازه گیری کنیم

2023.05.02 20:50:51.875 blog_kernel_times_benchmark (USDJPY,H1) Deployed all kernels!
2023.05.02 20:50:51.875 blog_kernel_times_benchmark (USDJPY,H1) All arguments for all kernels set!
2023.05.02 20:50:51.891 blog_kernel_times_benchmark (USDJPY,H1) Time to load and unload 50 kernels = 93ms

دوست داشتنی و 5000 هسته؟

خوب کمی تاخیر وجود دارد اما به نظر می رسد خوب است

2023.05.02 20:52:03.356 blog_kernel_times_benchmark (USDJPY,H1) All arguments for all kernels set!
2023.05.02 20:52:03.373 blog_kernel_times_benchmark (USDJPY,H1) Time to load and unload 5000 kernels = 110ms

حالا بالاخره بریم سر اصل مطلب!

ما باید به سرعت کرنل ها را روشن کنیم و بارهای کاری و افست آنها را تنظیم کنیم و یک فاصله زمانی 1 میلی ثانیه ای تنظیم کنیم، و همچنین نباید دوباره این قسمت از عملکرد تایمر را وارد کنیم در غیر این صورت با مشکل مواجه خواهیم شد. بنابراین . Bool indication kernelsRunning=false; 😊

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

بیایید همچنین یک نشانه تکمیل شده را در شی kernel_info خود اضافه کنیم، و من دستگیره را از تنظیمات حذف می کنم، زیرا از آن استفاده نکردم، این یک تابع خوب mql5، CLExecuteKernelList، مانند CommandQueue در api OpenCL اصلی خواهد بود.

وقتی همه چیز کامل شد، زمان ها را محاسبه می کنیم، اما بعداً به آن می پردازیم، بنابراین ابتدا دوباره به 5 هسته کاهش می دهیم.

اکنون کلاس ما به این صورت است:

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

دستورات تایمر موجود را با دروازه می کنیم

if(!kernelsRunning)
{
}

و ما در اینجا بسیار مراقب هستیم، باید پیش‌بینی کنیم که هسته‌ها پس از پایان تست، تخلیه شوند یا آزمایش اصلا شروع نشود، بنابراین:

یک متغیر exitNow را در بالا اضافه می کنیم که در صورت شکست تست یا پایان تست روی true تنظیم می شود.

همه زمینه ها تبدیل به متغیرهای دامنه جهانی می شوند …

ما برخی از موارد را از بخش قدیمی حذف می کنیم، نگران نباشید من آن را همانطور که در فایل منبع بود ذخیره کردم … بنابراین تایمر ما اکنون به این شکل است:

  bool exitNow=false;
  if(!kernelsRunning)
  {
  EventKillTimer();
  ctx=CLContextCreate(CL_USE_GPU_DOUBLE_ONLY);
  
  int kernels_to_deploy=5;  
  tangents_id=INVALID_HANDLE;
  if(ctx!=INVALID_HANDLE){
    string kernel="__kernel void bench(__global double* _tangent,"
                                      "int iterations){"
                                      "double sum=(double)0.0;"
                                      "double of=(double)_tangent[get_global_id(0)];"
                                      "for(int i=0;i<iterations;i++){"
                                      "sum+=((double)tanh(of-sum))/((double)iterations);"
                                      "}"
                                      "_tangent[get_global_id(0)]=sum;}";
    string errors="";
    prg=CLProgramCreate(ctx,kernel,errors);
    if(prg!=INVALID_HANDLE){
    ResetLastError();
    
      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!");    
      
        double tangents[];
        ArrayResize(tangents,kernels_to_deploy,0);
        double range=5.2;
        for(int i=0;i<ArraySize(tangents);i++){
           double r=(((double)MathRand())/((double)32767.0)*range)-2.6;
           tangents[i]=r;
           }         
        tangents_id=CLBufferCreate(ctx,kernels_to_deploy*8,CL_MEM_READ_WRITE);
        
          bool args_set=true;
          for(int i=0;i<kernels_to_deploy;i++){
             ResetLastError();
             if(!CLSetKernelArgMem(KERNELS[i].handle,0,tangents_id)){
               Print("Cannot assign buffer to kernel("+i+") error #"+IntegerToString(GetLastError()));
               args_set=false;
               }else{
               CLSetKernelArg(KERNELS[i].handle,1,iterations);
               }
             }
         if(args_set){
         Print("All arguments for all kernels set!");
         
         
         
         }else{
         Print("Cannot setup kernel args!");
         exitNow=true;
         }
      }else{
      Print("Cannot deploy all kernels!");
      exitNow=true;
      }
    
    }else{Alert(errors);exitNow=true;}
    }
  else{
    Print("Cannot create ctx");
    exitNow=true;
    }
  }
  
  
  
  
    if(exitNow){
    if(tangents_id!=INVALID_HANDLE){CLBufferFree(tangents_id);}
    for(int i=0;i<ArraySize(KERNELS);i++){
       if(KERNELS[i].handle!=INVALID_HANDLE){CLKernelFree(KERNELS[i].handle);}
       }
    if(prg!=INVALID_HANDLE){CLProgramFree(prg);}
    if(ctx!=INVALID_HANDLE){CLContextFree(prg);}
    Print("DONE");
    ExpertRemove();
    }

باشه حالا بیا فکر کنیم…

بیایید ابتدا به تکمیل کار بپردازیم، این بخش آسان است

پس از تکمیل، در حال حاضر، ما خارج شده و تایمر را متوقف می کنیم.

توجه داشته باشید که ما هنوز چیزی را “راه اندازی” نکرده ایم، ممکن است هنگام انجام این کار، خطاهای بی شماری دریافت کنیم.

  else if(!Busy&&kernelsRunning){
  Busy=true;
  
  
    bool still_running=false;
    for(int i=0;i<ArraySize(KERNELS);i++){
    if(!KERNELS[i].completed){
      if(CLExecutionStatus(KERNELS[i].handle)==CL_COMPLETE){
      KERNELS[i].completed=true;
      }else{still_running=true;}
      }
    }
  
    if(!still_running){
      EventKillTimer();
      exitNow=true;
      }
  
  if(!exitNow){Busy=false;}
  }

به نظر من به اندازه کافی ساده و درست به نظر می رسد:

• وارد لیست هسته ها می شویم
• اگر چیزی تمام شده است، آن را روی تکمیل تنظیم می کنیم
• اگر نه، پرچمی را که هنوز در حال اجراست روشن می کنیم
• از حلقه خارج شوید
• اگر چیزی هنوز در حال اجرا نیست، تایمر را بکشید
• اکنون خروج را روشن کنید
• نشانگر مشغول بودن را خاموش نکنید

اوه، فراموش کردم زمان پایان را اندازه گیری کنم! این را در قسمت تکمیل شده اضافه کنید

KERNELS[i].stop(GetMicrosecondCount());

و این فراخوان اجرایی است:

         
           uint offsets[]={0};
           uint works[]={1};
           for(int i=0;i<ArraySize(KERNELS);i++){
              offsets[0]=i;
              CLExecute(KERNELS[i].handle,1,offsets,works);
              KERNELS[i].setup(GetMicrosecondCount(),i);
              }
           kernelsRunning=true;
           EventSetMillisecondTimer(1);

بیایید ببینیم چه اتفاقی می افتد ، من اکنون نمی توانم چیزی را ببینم ، اگرچه خارج شد. بنابراین .. که با 1000 تکرار در 5 هسته بود

حال وظیفه این است که زمان اجرای هر کرنل را بالاتر از بازه تایمر … pffft ..

برای این کار باید یافته های خود را در یک فایل خروجی بگیریم!

      int f=FileOpen("OCL\kernel_bench.txt",FILE_WRITE|FILE_TXT);
      if(f!=INVALID_HANDLE){
        for(int i=0;i<ArraySize(KERNELS);i++){
           ulong micros=KERNELS[i].end_microSeconds-KERNELS[i].start_microSeconds;
           if(KERNELS[i].completed){
             FileWriteString(f,"K["+IntegerToString(i)+"] completed in ("+IntegerToString(micros)+")microSecondsn");
             }
           else
             {
             FileWriteString(f,"K["+IntegerToString(i)+"] not completedn");
             }
           }
        FileClose(f);
        }

این را به بلوک خروج اضافه می کنیم و منتظر می مانیم و می بینیم.

و وویلا

K[0] completed in (87334)microSeconds
K[1] completed in (87320)microSeconds
K[2] completed in (87300)microSeconds
K[3] completed in (87279)microSeconds
K[4] completed in (87261)microSeconds

حالا اینها یعنی چه؟ هیچ چیز آنها باید زیر آستانه اجرای ما باشند. اجازه بدید ببینم

یک میکروثانیه … 1000000م ثانیه است، یا یک ثانیه 1000000 میکروثانیه است، بنابراین آنچه در اینجا می بینیم 87 میلی ثانیه است و ما به فاصله 1 میلی ثانیه دسترسی داریم، بسیار خوب. من به آن اعتماد ندارم زیرا ممکن است برای حلقه نیز تاخیر ایجاد شود.

بنابراین … بیایید کالک ها را سنگین تر کنیم (تکرارهای بیشتر) من یک میلیون تکرار می فرستم. در حال حاضر، اینها در همان زمان کم و بیش تمام می شود

من همچنین mt5 را خاموش می‌کنم و برای هر بار اجرا مجدد آن را راه‌اندازی می‌کنم، نمی‌دانم که آیا کش وجود دارد یا خیر، اما می‌خواهم از آن اجتناب کنم.

-فکر می‌کنم باید تست را تا زمانی که وضعیت هسته‌ها در حال اجراست یا در صف اجراست یا smth ادامه دارد ادامه دهم-

به نظر می رسد که گیر کرده است … من انتظار داشتم 80 ثانیه اجرا شود، اکنون 5 دقیقه است … 15 دقیقه بسیار خوب، یک کارگزار. اجازه دهید چند مورد را آنجا اضافه کنید … @#%!#!%$@^

  
    bool still_running=false;
    int running_total=0;
    int completed_total=0;
    int queued_total=0;
    int submitted_total=0;
    int unknown_total=0;
    for(int i=0;i<ArraySize(KERNELS);i++){
    if(!KERNELS[i].completed){
      ENUM_OPENCL_EXECUTION_STATUS status=CLExecutionStatus(KERNELS[i].handle);
      if(status==CL_COMPLETE){
      completed_total++;
      KERNELS[i].completed=true;
      KERNELS[i].stop(GetMicrosecondCount());
      }else if(status==CL_RUNNING){running_total++;still_running=true;}
      else if(status==CL_QUEUED){queued_total++;}
      else if(status==CL_SUBMITTED){submitted_total++;}
      else if(status==CL_UNKNOWN){unknown_total++;}
      }else{
      completed_total++;
      }
    }
    string message="Running("+IntegerToString(running_total)+")n";
           message+="Completed("+IntegerToString(completed_total)+")n";
           message+="Queued("+IntegerToString(queued_total)+")n";
           message+="Submitted("+IntegerToString(submitted_total)+")n";
           message+="Unknown("+IntegerToString(unknown_total)+")n";
    Comment(message);

تغییر حلقه انتظار async به این … بیایید ببینیم چرا f*** شکست می خورد…

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

حلقه اکنون به این تغییر می کند:

  
    bool still_running=false;
    int running_total=0;
    int completed_total=0;
    int queued_total=0;
    int submitted_total=0;
    int unknown_total=0;
    for(int i=0;i<ArraySize(KERNELS);i++){
    if(!KERNELS[i].completed){
      ENUM_OPENCL_EXECUTION_STATUS status=CLExecutionStatus(KERNELS[i].handle);
      if(status==CL_COMPLETE){
      completed_total++;
      KERNELS[i].completed=true;
      KERNELS[i].stop(GetMicrosecondCount());
      }else if(status==CL_RUNNING){running_total++;still_running=true;}
      else if(status==CL_QUEUED){queued_total++;still_running=true;}
      else if(status==CL_SUBMITTED){submitted_total++;still_running=true;}
      else if(status==CL_UNKNOWN){unknown_total++;still_running=true;}
      }else{
      completed_total++;
      }
    }
    string message="Running("+IntegerToString(running_total)+")n";
           message+="Completed("+IntegerToString(completed_total)+")n";
           message+="Queued("+IntegerToString(queued_total)+")n";
           message+="Submitted("+IntegerToString(submitted_total)+")n";
           message+="Unknown("+IntegerToString(unknown_total)+")n";
    Comment(message);
  
    if(!still_running){
      EventKillTimer();
      exitNow=true;
      }

به 1000 تکرار کاهش یافت، فکر می کنم دیدم که یک به یک از هسته ها عبور می کند. بیایید تکرارهای x10 را اضافه کنیم و ببینیم.

همان 100k تکرار … همان . خوب آیا مشکلی با دقت اعشاری یا چیزی وجود دارد و ما نمی توانیم 1 میلیون را بزنیم؟

ما برویم، بله، یک ناشناخته باقی مانده است و 4 هسته کامل با یک میلیون تکرار، اما چرا؟

در آنجا گیر کرده است اما خوشبختانه به نظر نمی رسد مشکلی در دستگاه ایجاد کند!

اما چرا آنجا آویزان است؟ ، اگرچه زمان های 1000 10000 و 100000 تقریباً فوری بودند، بنابراین اگر زمان را بالاتر از بازه مورد نیاز خود قرار نمی دهیم، بیایید کمی تغییر کنیم، بیایید عملیات ها را در هسته ترکیب نکنیم فقط یک تن s*** را محاسبه کنیم و سپس آن را پاس کنید.

فکر نمی کنم مهم باشد (اگر مهم است و می دانید به من اطلاع دهید)

بنابراین ما += را حذف می کنیم اما اکنون مشکل این است که مقدار tanh را از حافظه پنهان خود ارائه می دهد …. بنابراین …. اجازه دهید این را به یک اضافه تبدیل کنیم … 😛

"sum=tanh(of)+((double)iterations)/((double)100.0);"

این calc است، بیایید دوباره برای 1 میلیون تکرار اجرا کنیم … aaand بله یک مشکل دقیق یا چیزی وجود داشت.

حالا بیایید زمان ها را بررسی کنیم.

K[0] completed in (370644)microSeconds
K[1] completed in (479982)microSeconds
K[2] completed in (604963)microSeconds
K[3] completed in (729959)microSeconds
K[4] completed in (839271)microSeconds

قطعاً یک عمل صف در اینجا در حال انجام است. بیایید هسته ها را به 50 برسانیم.

دوباره گیر کرد . 19 تکمیل شده 31 ناشناخته. باشه باید ببینیم خطا چیه

یک بررسی خطا در هنگام اجرا اضافه شد، بدون bueno. مشکل وجود ندارد، بنابراین اگر فرض کنیم از Sub->queued->running->unknown یا تکمیل شده باشد، باید در ناشناخته ظاهر شود.

بیایید ببینیم اسناد در مورد آن چه می گویند:

هوا، جالب … 😊 باشه . اجازه می دهد خطا وضعیت را نیز تعیین کند.

2023.05.02 22:14:21.458 blog_kernel_times_benchmark (USDJPY,H1) Unknown status for kernel(44) #5101
2023.05.02 22:14:21.458 blog_kernel_times_benchmark (USDJPY,H1) ----: unknown OpenCL error 65536

خطای ناشناخته … kay #5101 است …. خطای داخلی، بسیار خوب، بسیار روشنگر …

این به ما می‌گوید که این کار به وضوح کار نمی‌کند، بنابراین ما نمی‌توانیم به این روش محک بزنیم و به نوعی «api» درست است زیرا از ما می‌پرسد چرا روی زمین سعی می‌کنیم کاری را انجام دهیم که خودش قرار است انجام دهد.

من در اینجا استراحت می کنم اما وبلاگ ها را منتشر می کنم، ناقص بودن آنها (در حال حاضر) ممکن است به کسی یک یا دو ایده بدهد.

من همچنین 2 منبعی را که تا کنون استفاده کرده ام پیوست می کنم.

—- باشه —–

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

چه چیزی تغییر می کند:

• ما به ایده اصلی باز می گردیم.
• و ما فقط یک اندازه گیری زمان را روی کل موضوع انجام می دهیم.

زیرا آنچه در این صفحه می خوانید همان کاری است که لفاف OpenCL انجام می دهد، بنابراین تنها کاری که ما باید انجام دهیم این است که آن را برانیم یا هر چیزی که می خواهید آن را صدا بزنید تا تا شود.

ما اندازه گیری خواهیم کرد که چند گروه کاری (که در یک واحد محاسباتی کار می کنند) طول می کشد تا اندازه گیری زمان یک بار دو بار تا شود و غیره.

و در این مورد از “fold” منظور ما برعکس است، مانند، وقتی مشاهده می کنیم که افزایش آشکاری در زمان محاسبات وجود دارد، می دانیم که عددی که برای پردازش ارسال می کنیم به معنای چیزی است، یا آستانه ای که به تازگی از آن عبور کرده ایم به معنای چیزی است. .

آیا آسان تر نخواهد بود اگر بتوانید GPU خود را تولید کنید و از آن بپرسید “بله، چند هسته دارید که می توانم به طور همزمان از آنها استفاده کنم، چند آیتم دریافت می کنند و چقدر حافظه در آن وجود دارد؟” و در واقع جواب داد؟ بله، شما، mq، khronos و من احتمالاً به این فکر کرده ایم.

در مورد من منتظرم ببینم چه زمانی و کجا ارقام 192cores و 32warp در محاسبات ظاهر می شوند.

بنابراین، بنچمارک 2 به شکل زیر خواهد بود:

ما از 1D استفاده می کنیم

ما همچنان به ثبت شناسه گروه ادامه می دهیم! برای نگه داشتن برگه ها در گروه های ایجاد شده

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

<توجه داشته باشید، تلاش قبلی ممکن است مشکلی نداشته باشد، اما چیزهایی وجود دارد که من با آنها آشنا نیستم یا چیزی که نادیده گرفته ام در حال شکستن آن است>