ত্রুটি কোড: E3000

বিভাগ: কম্পাইল সময়: স্পার্সকোর বরাদ্দ ব্যর্থতা

এই ত্রুটিটি তখন ঘটে যখন XLA:SparseCore কম্পাইলার বর্তমান SparseCore প্রোগ্রামের জন্য প্রয়োজনীয় নির্দিষ্ট মেমরি স্পেসে একটি সংলগ্ন ব্লক মেমরি বরাদ্দ করতে অক্ষম হয়।

নমুনা ত্রুটি বার্তা:

INTERNAL:Failed to run pass pipeline. Hlo-Op: result.1:279:1: error: 'memref.alloca' op current allocation offset upper bound (140704 words) exceeds the legitimate user allocatable offset upper bound (131071 words) in memory space 201 when allocating 23440 words. result.1:279:1: note: see current operation: %232 = "memref.alloca"() <{operandSegmentSizes = array<i32: 0, 0>}> : () -> memref<23440xf32, 201>

XLA ব্যাকএন্ড: TPU

সংক্ষিপ্ত বিবরণ

স্পার্সকোর (SC) হল স্পার্স ওয়ার্কলোডের জন্য একটি বিশেষায়িত প্রসেসর। এটি ডেটা চলাচল দক্ষতার সাথে পরিচালনা করার জন্য নির্দিষ্ট মেমরি হায়ারার্কির উপর নির্ভর করে। XLA কম্পাইলার হার্ডওয়্যার সীমা এবং ব্যবহারকারী-সংজ্ঞায়িত আকারের উপর ভিত্তি করে স্ট্যাটিকভাবে বাফার আকার এবং বরাদ্দ করার চেষ্টা করে। এই ত্রুটিটি এই বরাদ্দ পর্যায়ে মেমরির বাইরে (OOM) অবস্থা নির্দেশ করে।

ত্রুটি বার্তাটি সাধারণত একটি মেমোরি স্পেস আইডি নির্দিষ্ট করে। নীচে সাধারণ মেমোরি স্পেস এবং তাদের পূর্ণসংখ্যা এনকোডিংগুলি দেওয়া হল:

মেমোরি স্পেস আইডি নাম বিবরণ
0 স্মেম স্থানীয় স্কেলার মেমোরি। স্কেলার রেজিস্টার এবং নিয়ন্ত্রণ প্রবাহের জন্য ব্যবহৃত হয়।
২০১ টাইলস্পেম টাইল-নির্দিষ্ট স্ক্র্যাচপ্যাড মেমোরি। একটি নির্দিষ্ট SC টাইলের জন্য দ্রুত, স্থানীয় SRAM উপলব্ধ।
২০২ স্পেম শেয়ার্ড স্ক্র্যাচপ্যাড মেমোরি। HBM ল্যাটেন্সি লুকানোর জন্য সুবিধাবাদীভাবে ডেটা (ইনপুট, আউটপুট, ইন্টারমিডিয়েট) মঞ্চস্থ করতে ব্যবহৃত হয়।
২০৩ এইচবিএম উচ্চ ব্যান্ডউইথ মেমোরি। টেবিল, হিপ এবং স্ট্যাক এম্বেড করার জন্য ব্যবহৃত বৃহৎ, ভাগ করা মেমোরি।
২০৪ পতাকা সিঙ্ক করুন সমন্বয়ের জন্য ব্যবহৃত সিঙ্ক্রোনাইজেশন প্রিমিটিভ।

SC এবং এর মেমরি শ্রেণিবিন্যাস সম্পর্কে গভীরভাবে জানতে, SparseCore ডকুমেন্টেশনটি দেখুন।

ডিবাগিং

রেজোলিউশন নির্ভর করে কোন মেমোরি স্পেস বরাদ্দ করা হয়নি তার উপর।

পরিস্থিতি ১. এইচবিএম বরাদ্দ ব্যর্থতা

মেমোরি স্পেস আইডি: ২০৩

এই ত্রুটিটি ঘটে যদি SparseCore প্রোগ্রাম দ্বারা অনুরোধ করা একটি একক অস্থায়ী বরাদ্দ উপলব্ধ HBM-এ ফিট করার জন্য খুব বেশি হয়। স্ট্যান্ডার্ড এম্বেডিং ওয়ার্কলোড এবং SC অফলোডেড কালেক্টিভগুলিতে অত্যন্ত বড় প্রতি-কোর পার্টিশন বা ভুল শার্ডিং স্পেসিফিকেশন কম্পাইলারকে বিশাল বাফার অনুরোধ করতে বাধ্য করতে পারে।

প্রস্তাবিত পদক্ষেপ:

  • শেয়ারিং পরীক্ষা করুন: নিশ্চিত করুন যে আপনার এম্বেডিং টেবিল এবং SC ইনপুট/আউটপুট টেনসরগুলি সঠিকভাবে পার্টিশন/শার্ড করা আছে। যদি একটি একক কোর খুব বেশি ডেটার জন্য দায়ী হয়, তাহলে বরাদ্দ ব্যর্থ হতে পারে।
  • সীমা সামঞ্জস্য করুন: max_ids_per_partition এবং max_unique_ids_per_partition পর্যালোচনা করুন। যদি এগুলি অপ্রয়োজনীয়ভাবে বেশি সেট করা হয়, তাহলে কম্পাইলার প্রয়োজনের চেয়ে বেশি মেমরি সংরক্ষণ করে। limits কীভাবে টেবিলে অনুবাদ করে তা দেখুন।

দৃশ্যপট ২। অভ্যন্তরীণ মেমরি ব্যর্থতা

মেমোরি স্পেস আইডি: ০, ২০১, ২০২, ২০৪

Smem , TileSpmem , Spmem , অথবা Sync Flag- এ বরাদ্দ ব্যর্থতা সাধারণত কম্পাইলার বাগ বা বরাদ্দ কৌশলের সীমাবদ্ধতার কারণে ঘটে, যেখানে কম্পাইলার সমস্ত মেমরির প্রয়োজনীয়তা পূরণ করতে ব্যর্থ হয়।

প্রস্তাবিত পদক্ষেপ:

  1. ব্যর্থ XLA অপারেশনটি আলাদা করুন: ব্যর্থতার কারণ হিসেবে নির্দিষ্ট SC HLO বা Mosaic কার্নেল সনাক্ত করতে, মধ্যবর্তী কম্পাইলার উপস্থাপনা তৈরি করুন:
    • SparseCore MLIR ডাম্প করুন: --xla_sc_dump_mlir_to=/path/to/dump ফ্ল্যাগ সেট করুন। এটি SparseCore প্রোগ্রামের MLIR তৈরি করে, যা আপনাকে দেখতে দেয় যে কোন বরাদ্দের আকার ত্রুটি বার্তার সাথে মেলে।
    • ডাম্প মোজাইক LLO: কাস্টম কার্নেলের জন্য, মোজাইক দ্বারা নির্গত সমস্ত নিম্ন স্তরের অপ্টিমাইজার (LLO) প্রোগ্রাম পরিদর্শন করতে --xla_mosaic_dump_to=/path/to/dump ব্যবহার করুন।
  2. স্ক্র্যাচের আকার হ্রাস করুন (প্যালাস ব্যবহারকারী): যদি কোনও মোজাইক কার্নেলের মধ্যে ব্যর্থতা দেখা দেয়, তাহলে আপনার scratch_shapes কনফিগারেশনটি পর্যালোচনা করুন। নিশ্চিত করুন যে আপনার pltpu.SMEM অনুরোধগুলি আপনার নির্দিষ্ট TPU প্রজন্মের জন্য হার্ডওয়্যার স্পেসিফিকেশনের মধ্যে ফিট করে।
  3. সমষ্টিগত অফলোড অক্ষম করুন: যদি SC অফলোড করা সমষ্টিগত ক্রিয়াকলাপ থেকে ত্রুটি দেখা দেয়, তাহলে SC অফলোডিং বৈশিষ্ট্যগুলি অক্ষম করার চেষ্টা করুন:
    • --xla_tpu_enable_sparse_core_collective_offload_all_gather=false
    • --xla_tpu_enable_sparse_core_collective_offload_all_reduce=false
  4. একটি বাগ রিপোর্ট করুন: যদি উপরের ধাপগুলি সমস্যার সমাধান না করে, তাহলে সম্ভবত এটি একটি কম্পাইলার বাগ। অনুগ্রহ করে একটি বাগ রিপোর্ট ফাইল করুন।