বিভাগ: কম্পাইল সময়: স্কোপড ভিমেম ওওএম
এই ত্রুটিটি নির্দেশ করে যে প্রোগ্রামটির জন্য বরাদ্দকৃত পরিমাণের চেয়ে বেশি স্কোপড ভেক্টর মেমরি (Vmem) প্রয়োজন।
নমুনা ত্রুটি বার্তা:
RESOURCE_EXHAUSTED: Ran out of memory in memory space vmem while allocating on stack for %my-custom-kernel = bf16[2048,4096]{1,0:T(8,128)(2,1)} custom-call(...) ...
XLA ব্যাকএন্ড: TPU
সংক্ষিপ্ত বিবরণ
টিপিইউ-তে ভেক্টর মেমরি (VMEM) থাকে, যা একটি লোকাল স্ক্র্যাচপ্যাড মেমরি এবং এটি শুধুমাত্র টেনসরকোর (TC) দ্বারা ব্যবহৃত হয়। কম্পাইলার বিভিন্ন ধরণের অ্যালোকেশনের জন্য Vmem পরিচালনা করে:
- ইন্সট্রাকশন-স্কোপড অ্যালোকেশন: একটিমাত্র HLO ইন্সট্রাকশন নির্বাহ করার সময় Vmem-এ অস্থায়ী সঞ্চয়। এর মধ্যে অপারেন্ড স্প্যান বাফার (যেমন ডাবল বাফারিং-এর জন্য) এবং রেজিস্টার স্পিল অন্তর্ভুক্ত।
- প্রোগ্রাম-স্কোপড অ্যালোকেশন: যে অ্যালোকেশনগুলো একটিমাত্র HLO ইন্সট্রাকশনের আওতার বাইরে থাকে। এগুলো সাধারণত HLO টেম্পোরারি এবং অন্তর্বর্তী ফলাফল, যা HLO ইন্সট্রাকশনের ইনপুট এবং/অথবা আউটপুট হিসেবে ব্যবহৃত হয়।
যখন ইন্সট্রাকশন-স্কোপড অ্যালোকেশন কোনো একটি ইন্সট্রাকশনের জন্য নির্ধারিত অ্যালোকেশন সীমা অতিক্রম করে, তখন একটি কম্পাইল টাইম স্কোপড ভিমেম ওওএম (OOM) ঘটে। এই সীমাটি নিয়ন্ত্রিত হয়।
- --xla_tpu_scoped_vmem_limit_kib ফ্ল্যাগের মাধ্যমে পুরো প্রোগ্রামের জন্য বিশ্বব্যাপী
এবং
- vmem_limit_bytes প্যারামিটারের মাধ্যমে কাস্টম কার্নেলের জন্য।
এই ত্রুটিগুলি সাধারণত কম্পাইলারের অভ্যন্তরীণ কোনো বাগের কারণে অথবা কোনো কাস্টম কার্নেল তার বরাদ্দ সীমা অতিক্রম করার কারণে ঘটে থাকে।
ব্যবহারযোগ্য সীমা {#usable_limits}
বিভিন্ন TPU হার্ডওয়্যারের Vmem সাইজ ভিন্ন ভিন্ন হয়। এছাড়াও, বিভিন্ন হার্ডওয়্যার বিভিন্ন কাজের জন্য ভিন্ন ভিন্ন পরিমাণে Vmem রিজার্ভ করতে পারে, যার ফলে Scoped Vmem-এ বরাদ্দযোগ্য পরিমাণ সীমিত হয়ে যায়।
নিচে বিভিন্ন হার্ডওয়্যারের সাধারণ সীমাবদ্ধতাগুলো দেওয়া হলো।
| হার্ডওয়্যার | সীমা (মানব-বান্ধব) | সীমা (কিলোবাইট) | সীমা (বাইট) |
|---|---|---|---|
| টিপিইউ ভি২ | ১৬ এমবি | ১৬৩৮৪ | ১৬৭৭৭২১৬ |
| টিপিইউ ভি৩ | ১৬ এমবি | ১৬৩৮৪ | ১৬৭৭৭২১৬ |
| টিপিইউ ভি৪ | ১৬ এমবি | ১৬৩৮৪ | ১৬৭৭৭২১৬ |
| টিপিইউ ভি৪আই | ১৬ এমবি | ১৬৩৮৪ | ১৬৭৭৭২১৬ |
| টিপিইউ ভি৫ই | ১২৮ এমবি | ১৩১০৭২ | ১৩৪২১৭৭২৮ |
| টিপিইউ ভি৫পি | ৬৩.৯ মেগাবাইট | ৬৫৪৭২ | ৬৭০৪৩৩২৮ |
| টিপিইউ ভি৬ই | ১২৭.৯ মেগাবাইট | ১৩১০০৮ | ১৩৪১৫২১৯২ |
| টিপিইউ ৭x | ৬৩.৯ মেগাবাইট | ৬৫৪৭২ | ৬৭০৪৩৩২৮ |
ডিবাগিং
ত্রুটিটি কাস্টম কার্নেল নাকি স্ট্যান্ডার্ড HLO থেকে উদ্ভূত হয়েছে তা শনাক্ত করতে ত্রুটির বার্তাটি মনোযোগ সহকারে বিশ্লেষণ করুন। কাস্টম কার্নেলের কারণে সৃষ্ট ত্রুটির সিগনেচারটি নিম্নরূপ হওয়া উচিত:
Ran out of memory in memory space vmem while allocating on stack for %my-custom-call = <output-shape> custom-call(<params>), custom_call_target="tpu_custom_call" ...
- কাস্টম কার্নেল স্কোপড Vmem OOM : যদি ত্রুটিটি একটি কাস্টম কার্নেলকে নির্দেশ করে → কার্নেল রিটিউন- এ যান।
- নন-কার্নেল Vmem সংক্রান্ত সমস্যা : যদি কোনো নন-কাস্টম-কার্নেল অপারেশনের কারণে Vmem OOM ঘটে, তবে এটি সম্ভবত একটি অভ্যন্তরীণ কম্পাইলার বাগ। অনুগ্রহ করে একটি HLO ডাম্প সহ XLA-তে একটি বাগ রিপোর্ট ফাইল করুন ।
কার্নেলটি পুনরায় টিউন করুন
যদি ত্রুটিটি কোনো কাস্টম কার্নেল থেকে উদ্ভূত হয়, তাহলে কার্নেলের মেমরির প্রয়োজনীয়তা কমাতে নিম্নলিখিত কৌশলগুলো ব্যবহার করুন:
- ব্লক সাইজ সমন্বয় করুন: Scoped Vmem-এর ব্যবহার কমাতে আপনার কার্নেল কনফিগারেশনে ব্লক সাইজ (টাইল সাইজ) হ্রাস করুন।
- প্রতি-কার্নেল স্কোপড ভিমেম লিমিট সেট করুন: vmem_limit_bytes প্যারামিটার ব্যবহার করে সেই নির্দিষ্ট কার্নেলের জন্য প্রয়োজনীয় মেমরির পরিমাণ স্পষ্টভাবে অনুরোধ করুন।
- মেমরি কালারিং পরিবর্তন করুন: pallas.tpu.with_memory_space_constraint ব্যবহার করে Vmem-এ কার্নেলের ইনপুট/আউটপুটগুলিকে সুস্পষ্টভাবে কালার বা সীমাবদ্ধ করুন। খেয়াল রাখবেন যেন Vmem-এর অনেক বেশি ইনপুট/আউটপুট কালার করা না হয়, কারণ এর ফলে সামগ্রিকভাবে Vmem OOM (আউট অফ মেমরি) হতে পারে।
- Vmem সীমা সমন্বয় করুন: যদি কার্নেল-নির্দিষ্ট রিটিউনিং কঠিন হয় অথবা সমস্যাটি অনেকগুলো কার্নেলকে প্রভাবিত করে, তাহলে আপনি --xla_tpu_scoped_vmem_limit_kib ফ্ল্যাগটি ব্যবহার করে গ্লোবাল Vmem সীমা সমন্বয় করতে পারেন।