HLO থেকে Thunks

এই ডকুমেন্টে একটি XLA হাই লেভেল অপ্টিমাইজার (HLO) মডিউলের প্রাথমিক অবস্থা থেকে চূড়ান্ত এক্সিকিউটেবল অবস্থা পর্যন্ত যাত্রার রূপরেখা দেওয়া হয়েছে। কখনও কখনও আমরা "মডিউল" বাদ দিই এবং এটিকে "HLO" হিসাবে উল্লেখ করি।

HLO থেকে Thunks ডায়াগ্রাম

প্রি-অপ্টিমাইজেশন এইচএলও

আমরা প্রি-অপ্টিমাইজেশন HLO মডিউল দিয়ে শুরু করি। প্রি-অপ্টিমাইজেশন HLO তে এমন কোনও অপারেশন ( ops ) থাকে না যা XLA এর অভ্যন্তরীণ হিসাবে বিবেচিত হয়, যেমন fusion বা bitcast । এই পর্যায়ে Ops এর কোনও লেআউট নেই, অথবা যদি থাকে তবে এটি উপেক্ষা করা হবে। প্রি-অপ্টিমাইজেশন HLO সাধারণত TensorFlow এবং JAX এর মতো উচ্চ-স্তরের ফ্রেমওয়ার্ক দ্বারা উত্পাদিত হয়। XLA ফ্ল্যাগ -xla_dump_to ব্যবহার করার সময়, প্রি-অপ্টিমাইজেশন HLO ফাইলের নাম "before_optimizations.txt" প্রত্যয় সহ একটি ফাইলে ডাম্প করা হয়।

HLO মডিউল অপ্টিমাইজ করুন

XLA:GPU পাইপলাইন পাসের একটি ক্রম চালিয়ে প্রি-অপ্টিমাইজেশন HLO কে অপ্টিমাইজড HLO তে রূপান্তরিত করে। পাসগুলিকে শব্দার্থগতভাবে একসাথে গ্রুপ করা যেতে পারে এবং নিম্নলিখিত ক্রমে চালানো যেতে পারে:

এর মধ্যে শার্ডি পার্টিশনারের মতো পাস বা এসপিএমডি শার্ডিংয়ের জন্য পাস অন্তর্ভুক্ত রয়েছে।

অপ্টিমাইজেশন পাস

এর মধ্যে বৈধকরণ পাস এবং সরলীকরণ পাস উভয়ই অন্তর্ভুক্ত থাকতে পারে।

যৌথ অপ্টিমাইজেশন পাস

অপ্টিমাইজেশন পাসের মতো, কিন্তু যৌথ অপারেশনের উপর ফোকাস করে।

লেআউট অ্যাসাইনমেন্ট পাস

প্রতিটি HLO অপশনকে একটি লেআউট বরাদ্দ করা হয় যা নির্দেশ আকৃতির একটি অংশ। লেআউটটি নিয়ন্ত্রণ করে যে টেনসরটি মেমরিতে শারীরিকভাবে কীভাবে স্থাপন করা হয়।

লেআউট সহ একটি আকৃতির উদাহরণ:

f32[10,20,30]{2,0,1}

এলিমেন্ট টাইপের পরে, আকৃতির লজিক্যাল ডাইমেনশন থাকে, তারপরে মাইনর থেকে মেজর ক্রমে লেআউট বিন্যাসের ক্রমানুসারে পরিবর্তন হয়। এই উদাহরণে, সবচেয়ে মাইনর ডাইমেনশন হল 30, দ্বিতীয় মাইনর ডাইমেনশন হল 10, এবং মেজর ডাইমেনশন হল 20।

লেআউট অ্যাসাইনমেন্টের লক্ষ্য হল একটি লোভী কৌশল ব্যবহার করে প্রয়োজনীয় ভৌত স্থানান্তরের সংখ্যা কমানো। এটি নির্দিষ্ট লেআউট সীমাবদ্ধতা দিয়ে শুরু হয় (যেমন, cuDNN/cuBLAS লাইব্রেরিগুলি ধারাবাহিক মাত্রা আশা করে) এবং লেআউটগুলিকে "নিচে" এবং তারপর HLO গ্রাফে "উপরে" প্রচার করে। লেআউট প্রচারের শেষে, কিছু নির্দেশাবলীর পরস্পরবিরোধী লেআউট থাকতে পারে, একটি অপারেন্ড থেকে প্রচারিত, অন্যটি ব্যবহারকারী থেকে প্রচারিত। এই দ্বন্দ্ব সমাধানের জন্য, একটি copy HLO নির্দেশ সন্নিবেশ করা হয় যা অপারেন্ড লেআউট থেকে নির্দেশ লেআউটে লেআউট পরিবর্তন করে।

লেআউট স্বাভাবিকীকরণ পাস

যেহেতু ভৌত আকৃতি বের করা কিছুটা কঠিন, তাই লেআউট নরমালাইজেশন আকৃতিটি এমনভাবে পুনর্লিখনের চেষ্টা করে যাতে এটি ডিফল্ট লেআউট {rank-1, rank-2, …, 0} ব্যবহার করে। উপরের উদাহরণে, নরমালাইজড শেপ হবে f32[20,10,30]{2,1,0} । লেআউট পরিবর্তনকারী কপি অপশনগুলি transpose এবং bitcast সংমিশ্রণে পুনর্লিখন করা হয়। বর্তমানে আমরা সমস্ত অপশন স্বাভাবিক করতে পারছি না, তবুও কিছু অপশন রয়েছে যার অ-ডিফল্ট লেআউট থাকতে পারে, বিশেষ করে gather এবং dot । নরমালাইজড অপশন এবং নন-নরমালাইজড অপশনের মধ্যে সীমানায় bitcast অপশন থাকবে যা একটি ট্রান্সপোজকে প্রতিনিধিত্ব করবে, অর্থাৎ একটি ট্রান্সপোজ যার একটি লেআউট নির্ধারিত থাকে যা এটিকে শারীরিকভাবে একটি নো-অপ করে তোলে।

লেআউট নরমালাইজেশন কিছু ইমপ্লিসিট ট্রান্সপোজকে স্পষ্ট করে তোলে যা গুরুত্বপূর্ণ কারণ কোডজেন একটি ডেডিকেটেড ইমিটারের সাহায্যে স্পষ্ট ট্রান্সপোজ পরিচালনা করতে পারে। উদাহরণস্বরূপ, একটি রিশেপকে টেকনিক্যালি অপারেন্ড এবং ফলাফলের মধ্যে একটি ভিন্ন ভৌত লেআউট থাকতে দেওয়া হয় (যেমন ভিন্ন র‍্যাঙ্কের কারণে)। লেআউট নরমালাইজেশন পাসের অংশ হিসাবে চালিত ReshapeDecomposer পাস একটি রিশেপকে transpose , রিশেপ bitcast এবং transpose ক্রমানুসারে রূপান্তরিত করে।

পোস্ট লেআউট অ্যাসাইনমেন্ট অপ্টিমাইজেশন পাস

এখানে সবচেয়ে গুরুত্বপূর্ণ পাসগুলি হল ট্রাইটন ফিউশন (GEMM ফিউশন + সফটম্যাক্স/লেয়ারনর্ম ফিউশন) অথবা লাইব্রেরি কলগুলিতে পুনর্লিখন। এই ধাপে অটোটিউনিংও চলে, যেখানে XLA বিভিন্ন ইমিটারের মধ্যে নির্বাচন করে, কনভলিউশন বা ডটের জন্য সেরা অ্যালগরিদম নির্বাচন করে, ট্রাইটন ইমিটার দ্বারা পরিচালিত ফিউশনের জন্য সেরা টাইলিং খুঁজে বের করে ইত্যাদি।

ফিউশন পাস

দুটি প্রধান পাস হল PriorityFusion এবং Multi-Output ফিউশন।

PriorityFusion এ, আমরা খরচ মডেল দ্বারা পরিচালিত ফিউশন তৈরি করি। ফিউজ করার সময়, যদি সমস্ত ব্যবহারকারীর মধ্যে অপটি মিশ্রিত করা সম্ভব হয়, তাহলে আমরা একাধিক ব্যবহারকারীর সাথে অপটি ডুপ্লিকেট করার অনুমতি দেব। সম্ভব হলে আমরা বিদ্যমান ট্রাইটন সফটম্যাক্স ফিউশনগুলিকেও প্রসারিত করার অনুমতি দেব।

Multi-Output ফিউশন হল একটি পৃথক পাস যা একটি অপারেন্ড ভাগ করে এমন অপস/ফিউশনগুলিকে ফিউজ করার অনুমতি দেয়। এটি অতিরিক্ত আউটপুট(গুলি) যোগ করে ডুপ্লিকেশন ছাড়াই ব্যবহারকারীদের মধ্যে অপারেন্ড/অপারেন্ড ফিউশনগুলিকে ফিউজ করতে পারে, যাতে ফিউজ করা অপের অন্যান্য ব্যবহারকারীরা এই আউটপুটগুলিতে পুনঃনির্দেশিত হতে পারে। এই পাসটি HLO গ্রাফে চক্রগুলি প্রবেশ না করার জন্য সতর্ক থাকতে হবে।

মাল্টি-আউটপুট ফিউশনের পরে, সাধারণ সাবএক্সপ্রেশন এলিমিনেশন ( HloCSE পাস) চলে, যা পূর্বে ডুপ্লিকেট করা অপশনগুলিকে একই ফিউশনে পরিণত করলে পুনরায় একত্রিত করার সম্ভাবনা বেশি।

বেশ কিছু পোস্ট-ফিউশন পাস

সমষ্টিগত সম্পর্কিত বেশ কয়েকটি পাস (যেমন সেগুলিকে অ্যাসিঙ্কে পরিণত করা, অথবা সমষ্টিগতের একটি নির্দিষ্ট আপেক্ষিক ক্রম প্রয়োগ করা)।

অবশেষে আমরা CopyInsertion চালাই যেখানে কপি যোগ করা হয় যাতে ইন-প্লেস অপারেশনগুলি অন্য কোথাও প্রয়োজনীয় ডেটা ওভাররাইট না করে।

অপ্টিমাইজেশনের শেষে, যদি "after_optimizations.txt" ফাইলের নামের প্রত্যয়যুক্ত ফাইলে -xla_dump_to ফ্ল্যাগ ব্যবহার করা হয়, তাহলে অপ্টিমাইজ করা HLO ডাম্প করা হয়। যদি আপনি মধ্যবর্তী পাসের পরে HLO ডাম্প করতে চান যা আসলে HloModule পরিবর্তন করে, তাহলে আপনি -xla_dump_hlo_pass_re=.* ফ্ল্যাগ (অথবা নির্দিষ্ট পাসের মধ্যে সীমাবদ্ধ রাখার জন্য একটি নির্দিষ্ট রেগুলার এক্সপ্রেশন) ব্যবহার করতে পারেন।

সময়সূচী

একটি শিডিউল ছাড়া একটি HLO মডিউলে অপশন প্রক্রিয়াকরণের ক্রমানুসারে কিছুটা স্বাধীনতা থাকে। অপারেন্ড/ফলাফল সম্পর্ক এবং নিয়ন্ত্রণ নির্ভরতা সম্পর্কিত যেকোনো টপোলজিক্যাল সাজানো বৈধ। শিডিউলিং কোন নির্দিষ্ট ক্রম ব্যবহার করতে হবে তা নির্ধারণ করে। এই পর্যায়ে প্রধান উদ্বেগ হল সর্বোচ্চ মেমোরি খরচ যা টেনসরের জীবনকালের উপর নির্ভর করে। প্রাথমিক ধাপে, আমরা বিভিন্ন শিডিউলার অ্যালগরিদম চেষ্টা করি এবং সেই সময়সূচী নির্বাচন করি যা সর্বোচ্চ মেমোরি খরচ কমিয়ে আনবে। মনে রাখবেন যে এই মুহুর্তে আমরা এখনও কোনও ফিজিক্যাল বাফার নিয়ে কাজ করি না (এটি "বাফার অ্যাসাইনমেন্ট" এ ঘটবে) এবং মেমোরি ব্যবহার অনুকরণ করি।

তারপর LatencyHidingScheduler পাস রান করে এবং কম্পিউট-কমিউনিকেশন ওভারল্যাপ সর্বাধিক করার চেষ্টা করে। কিন্তু এটি আবার মেমরির ব্যবহার বাড়িয়ে দিতে পারে।

পরিশেষে, যদি সর্বোচ্চ মেমোরি খরচ আমাদের কাছে উপলব্ধ মেমোরির পরিমাণের চেয়ে বেশি হয়, তাহলে আমরা HloRematerialization চালাই। এই পাসটি কর্মক্ষমতার মূল্যে মেমোরির ব্যবহার কমানোর চেষ্টা করে, যেমন কিছু ফিউশন বিভক্ত হতে পারে এবং কিছু অপশনকে ছোট বাফার লাইফটাইম দেওয়ার জন্য ডুপ্লিকেট করা যেতে পারে। যদি রিমেটেরিয়ালাইজেশন ঘটে, তাহলে মডেলের দিকে মেমোরির প্রয়োজনীয়তা কমানোর উপায়গুলি অনুসন্ধান করা উপকারী হতে পারে (যেমন, ছোট ব্যাচ আকার ব্যবহার করে)।

বাফার অ্যাসাইনমেন্ট

LLVM IR-এ নামানোর ঠিক আগে, আমরা বাফার অ্যাসাইনমেন্ট পাসগুলি চালাই যা HLO গ্রাফের প্রতিটি নির্দেশের জন্য বাফার স্লাইস বরাদ্দ করবে। বাফার অ্যাসাইনমেন্টটি বেশ কয়েকটি ধাপে সম্পন্ন হয়:

  1. HloDataflowAnalysis নির্দেশাবলীতে HloValues ​​(মূলত লজিক্যাল বাফার) বরাদ্দ করে। ইন-প্লেস অপশনের জন্য, একটি অপারেন্ডের HloValue পুনরায় ব্যবহার করা যেতে পারে। একটি অপ একাধিক HloValue সংজ্ঞায়িত করতে পারে (যেমন একটি টুপল ফলাফল আকৃতি সহ)।

  2. HloAliasAnalysis অ্যালিয়াসিং অপারেশনের জন্য বাফারগুলিকে একত্রিত করার চেষ্টা করে এবং HloValue থেকে HloBuffer পর্যন্ত একটি ম্যাপিং গণনা করে।

  3. BufferAssignment HloBuffers এর ম্যাপিং গণনা করে যাতে একটি বড় বাফারের ভিতরে স্লাইসগুলিকে এমনভাবে বাফার করা যায় যাতে একই বাফার স্লাইসটি বিভিন্ন HloBuffers এর জন্য ব্যবহার করা না হয় যাদের লাইফ টাইম ওভারল্যাপিং থাকে। যেসব অপশনের নাম হতে পারে, তাদের ক্ষেত্রে সামান্য ওভারল্যাপ থাকা ঠিক আছে (একটি HloBuffer এর শেষ সময় অন্য HloBuffer এর শুরুর সময়ের সাথে মিলে যেতে পারে)। -xla_dump_to ফ্ল্যাগ ব্যবহার করার সময়, বাফার অ্যাসাইনমেন্ট সম্পর্কে কিছু তথ্য "after_optimizations-buffer-assignment.txt" নামক প্রত্যয় সহ একটি ফাইলে ডাম্প করা হয়।

থাঙ্কস

একটি HLO গ্রাফ অপ্টিমাইজ এবং নির্ধারিত হওয়ার পরে, এটি একটি নির্দিষ্ট ব্যাকএন্ডের (CPU বা GPU) জন্য থাঙ্কের একটি রৈখিক ক্রমানুসারে নামানো হয়।

XLA-তে, একটি Thunk হল একটি স্বয়ংসম্পূর্ণ কাজের ইউনিটের একটি বিমূর্তকরণ যা রানটাইম সম্পাদন করে। এটি একটি কম্পাইল করা কার্নেল লঞ্চ, নির্দিষ্ট অপারেশন, লাইব্রেরি কল, নিয়ন্ত্রণ-প্রবাহ গঠন, যৌথ যোগাযোগ ইত্যাদি হতে পারে। একটি Thunk Sequence একটি নির্দিষ্ট ব্যাকএন্ডের জন্য সম্পূর্ণ এক্সিকিউটেবলকে প্রতিনিধিত্ব করে।

থাঙ্ক নির্গমন

একটি নির্ধারিত HLO গণনাকে একটি থাঙ্ক সিকোয়েন্সে রূপান্তর করার প্রক্রিয়াটিকে "থাঙ্ক নির্গমন" বলা হয়। এটি প্রতিটি ব্যাকএন্ডে একটি ডেডিকেটেড ইমিটার ক্লাস দ্বারা পরিচালিত হয়।

GPU ব্যাকএন্ডের জন্য, এটি IrEmitterUnnested দ্বারা পরিচালিত হয়। EmitHloComputation একটি গণনায় HLO নির্দেশাবলীর নির্ধারিত তালিকার মধ্য দিয়ে পুনরাবৃত্তি করে এবং একটি বিশেষায়িত Emit... পদ্ধতিতে প্রেরণ করে (যেমন, EmitFusion , EmitConvolutionThunk , EmitWhile )। এই প্রতিটি পদ্ধতি উপযুক্ত Thunk অবজেক্ট(গুলি) তৈরি করে এবং সেগুলিকে thunk ক্রমানুসারে যুক্ত করে।

CPU ব্যাকএন্ডের জন্য, ThunkEmitter এই ভূমিকা পালন করে এবং একইভাবে সংগঠিত হয়। Final ThunkSequence CpuExecutable এ এমবেড করা আছে।

মনে রাখবেন যে একটি HLO মডিউলের এন্ট্রি গণনার প্রতিটি নির্দেশ চূড়ান্ত থাঙ্ক ক্রমের no ( kTuple , kConstant , ..), এক, অথবা একাধিক (উদাহরণস্বরূপ sort instruction) থাঙ্কের সাথে সঙ্গতিপূর্ণ হতে পারে।

কমান্ড বাফার: GPU-তে এক্সিকিউশন অপ্টিমাইজ করা

আধুনিক GPU হার্ডওয়্যার আপনাকে GPU অপারেশনের একটি ক্রম (কার্নেল লঞ্চ, মেমোরি কপি ইত্যাদি) একবার রেকর্ড করতে এবং তারপর ন্যূনতম CPU ওভারহেডের মাধ্যমে একাধিকবার ক্রমটি পুনরায় চালানোর অনুমতি দেয়। এটি একটি গুরুত্বপূর্ণ কর্মক্ষমতা অপ্টিমাইজেশন, বিশেষ করে অনেক ছোট, দ্রুত-লঞ্চিং কার্নেল সহ ওয়ার্কলোডের জন্য। XLA CUDA গ্রাফ বা HIP গ্রাফের একটি বিমূর্ত রূপ হিসেবে কমান্ড বাফার ব্যবহার করে। মূল ইন্টারফেসটি GpuCommandBuffer এ সংজ্ঞায়িত করা হয়েছে।

CommandBufferThunk দ্বারা একটি কমান্ড বাফারকে একটি thunk ক্রমানুসারে উপস্থাপন করা হয়।

ইমিটারটি সরাসরি HLO নির্দেশাবলী থেকে এই থাঙ্ক তৈরি করে না। পরিবর্তে, এটি CommandBufferConversionPass দ্বারা করা হয় যা ThunkSequence-এ চলে।

এই পাসটি সামঞ্জস্যপূর্ণ thunks-এর (যেমন, KernelThunk s এবং GemmThunk s-এর একটি সিরিজ) সংলগ্ন সাব-সিকোয়েন্স সনাক্ত করে। এরপর এটি পাওয়া সাব-সিকোয়েন্সটিকে একটি একক CommandBufferThunk দিয়ে প্রতিস্থাপন করে। নতুন thunkটি মূল thunks-এর লজিককে হালকা CommandBufferCmd অবজেক্টের তালিকা হিসেবে ধারণ করে। যখন একটি CommandBufferThunk একটি প্রদত্ত GPU স্ট্রীমে প্রথমবারের মতো কার্যকর করে, তখন এটি তার কমান্ডের ক্রমকে একটি হার্ডওয়্যার কমান্ড বাফারে "রেকর্ড" করে। পরবর্তী সমস্ত কার্যকরকরণে, এটি রেকর্ড করা ক্রমটিকে "রিপ্লে" করার জন্য GPU-তে কেবল একটি কমান্ড জারি করে। এটি প্রতিটি পৃথক কার্নেল চালু করার CPU ওভারহেড এড়ায়।

কার্যকরযোগ্য

XLA কম্পাইলেশন পাইপলাইনের চূড়ান্ত পণ্য হল একটি স্বয়ংসম্পূর্ণ, প্ল্যাটফর্ম-নির্দিষ্ট Executable । এই অবজেক্টটি CPU বা GPU এর মতো একটি টার্গেট ডিভাইসে কম্পাইল করা প্রোগ্রাম চালানোর জন্য প্রয়োজনীয় সমস্ত তথ্য ধারণ করে। এটি কম্পাইলার এবং রানটাইমের মধ্যে সেতু। PJRT এর মতো আধুনিক রানটাইমগুলি সামান্য উচ্চ-স্তরের অ্যাবস্ট্রাকশন ব্যবহার করে ( PjRtExecutable দেখুন), কিন্তু এগুলি শেষ পর্যন্ত একটি ব্যাকএন্ড-নির্দিষ্ট এক্সিকিউটেবলকে আবৃত করে।

একটি Executable সংকলনের সময় তৈরি হওয়া বেশ কয়েকটি গুরুত্বপূর্ণ তথ্য থাকে। যদিও সঠিক বিষয়বস্তু ব্যাকএন্ড অনুসারে পরিবর্তিত হয়, তবে সাধারণত এতে অন্তর্ভুক্ত থাকে:

  • কম্পাইল করা কোড: এটি হল নিম্ন-স্তরের মেশিন কোড যা ডিভাইসে চলবে। সিপিইউগুলির জন্য, এটি সাধারণত এক বা একাধিক অবজেক্ট ফাইল। জিপিইউগুলির জন্য, এটি হল PTX বা HSACO ফর্ম্যাটে কম্পাইল করা ডিভাইস কোড, যা রানটাইমের সময় জিপিইউতে লোড করা হয়।

  • এক্সিকিউশন প্ল্যান (ThunkSequence): রানটাইম লজিকের মূল অংশ। এটি Thunk অবজেক্টের একটি রৈখিক ক্রম। প্রতিটি thunk কাজের একটি একক ইউনিটকে প্রতিনিধিত্ব করে, যেমন একটি কার্নেল চালু করা, একটি লাইব্রেরি ফাংশন কল করা (যেমন, cuBLAS), অথবা নিয়ন্ত্রণ প্রবাহ পরিচালনা করা। রানটাইম এই ক্রমটির মাধ্যমে পুনরাবৃত্তি করে প্রোগ্রামটি কার্যকর করে।

  • মেমোরি লেআউট (BufferAssignment): BufferAssigner দ্বারা তৈরি এই গুরুত্বপূর্ণ মেটাডেটাটি গণনার জন্য সম্পূর্ণ মেমোরি লেআউট বর্ণনা করে। এটি প্রতিটি বাফারের আকার এবং প্যারামিটার, আউটপুট এবং অস্থায়ী মানের জন্য মেমোরি কীভাবে বরাদ্দ এবং পুনঃব্যবহার করা হয় তা নির্দিষ্ট করে। রানটাইম ডিভাইস মেমোরি বরাদ্দ করতে এবং প্রতিটি থাঙ্কে সঠিক পয়েন্টার পাস করতে এটি ব্যবহার করে।

  • (ঐচ্ছিক) HLO মডিউল: ডিবাগিং এবং প্রোফাইলিংয়ের জন্য, এক্সিকিউটেবল প্রায়শই চূড়ান্ত, অপ্টিমাইজ করা HloModule-এর একটি রেফারেন্স ধরে রাখে যেখান থেকে এটি কম্পাইল করা হয়েছিল।

প্রতিটি নির্দিষ্ট ব্যাকএন্ডের জন্য কম্পাইলার দ্বারা চূড়ান্ত এক্সিকিউটেবল তৈরির কাজটি সাজানো হয়। কম্পাইলার বাস্তবায়নের RunBackend পদ্ধতি হল কম্পাইলেশন প্রক্রিয়ার চূড়ান্ত ধাপ, যা সমস্ত কম্পাইল করা আর্টিফ্যাক্টগুলিকে একটি এক্সিকিউটেবল অবজেক্টে প্যাকেজ করে। GpuCompiler এবং CpuCompiler যথাক্রমে GPU এবং CPU লক্ষ্য করে।

যখন একজন ব্যবহারকারী একটি এক্সিকিউটেবলে Execute... কল করে, তখন রানটাইম মেমরি বরাদ্দ করার জন্য BufferAssignment ব্যবহার করে, এবং তারপর কম্পাইল করা কোড ব্যবহার করে ডিভাইসে অপারেশন শুরু করার জন্য ThunkSequence ব্যবহার করে।