ত্রুটি কোড: E1001

বিভাগ: কম্পাইল সময়: স্কোপড ভিমেম ওওএম

এই ত্রুটিটি নির্দেশ করে যে প্রোগ্রামটির জন্য বরাদ্দকৃত পরিমাণের চেয়ে বেশি স্কোপড ভেক্টর মেমরি (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 ফ্ল্যাগের মাধ্যমে পুরো প্রোগ্রামের জন্য বিশ্বব্যাপী

এবং

এই ত্রুটিগুলি সাধারণত কম্পাইলারের অভ্যন্তরীণ কোনো বাগের কারণে অথবা কোনো কাস্টম কার্নেল তার বরাদ্দ সীমা অতিক্রম করার কারণে ঘটে থাকে।

ব্যবহারযোগ্য সীমা {#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" ...

কার্নেলটি পুনরায় টিউন করুন

যদি ত্রুটিটি কোনো কাস্টম কার্নেল থেকে উদ্ভূত হয়, তাহলে কার্নেলের মেমরির প্রয়োজনীয়তা কমাতে নিম্নলিখিত কৌশলগুলো ব্যবহার করুন:

  • ব্লক সাইজ সমন্বয় করুন: Scoped Vmem-এর ব্যবহার কমাতে আপনার কার্নেল কনফিগারেশনে ব্লক সাইজ (টাইল সাইজ) হ্রাস করুন।
  • প্রতি-কার্নেল স্কোপড ভিমেম লিমিট সেট করুন: vmem_limit_bytes প্যারামিটার ব্যবহার করে সেই নির্দিষ্ট কার্নেলের জন্য প্রয়োজনীয় মেমরির পরিমাণ স্পষ্টভাবে অনুরোধ করুন।
  • মেমরি কালারিং পরিবর্তন করুন: pallas.tpu.with_memory_space_constraint ব্যবহার করে Vmem-এ কার্নেলের ইনপুট/আউটপুটগুলিকে সুস্পষ্টভাবে কালার বা সীমাবদ্ধ করুন। খেয়াল রাখবেন যেন Vmem-এর অনেক বেশি ইনপুট/আউটপুট কালার করা না হয়, কারণ এর ফলে সামগ্রিকভাবে Vmem OOM (আউট অফ মেমরি) হতে পারে।
  • Vmem সীমা সমন্বয় করুন: যদি কার্নেল-নির্দিষ্ট রিটিউনিং কঠিন হয় অথবা সমস্যাটি অনেকগুলো কার্নেলকে প্রভাবিত করে, তাহলে আপনি --xla_tpu_scoped_vmem_limit_kib ফ্ল্যাগটি ব্যবহার করে গ্লোবাল Vmem সীমা সমন্বয় করতে পারেন।