আমাদের আগের পোস্টে ভেক্টর-অ্যাড কার্নেলকে nvcc থেকে ওয়ার্প পর্যন্ত অনুসরণ করা হয়েছিল। এবার, আমরা একটি গুরুত্বপূর্ণ SASS নির্দেশ (গ্লোবাল লোড) কে RTX4090-এর হার্ডওয়্যারের মধ্য দিয়ে অনুসরণ করি। পরীক্ষিত CUDA কার্নেলটিতে দুটি লাইন রয়েছে: __global__ void vadd(const float* a, const float* b, float* c, int n) {int i = blockIdx.x * blockDim.x + threadIdx.x; if (i < n) c[i] = a[i] + b[i];}। সংকলিত SASS দেখায়: /*0080*/ IMAD.WIDE R4, R6, R7, c[0x0][0x168] ; // &b[i] /*00a0*/ LDG.E R4, [R4.64] ; // b[i]। এই নির্দেশগুলি ভেক্টর b-এর উপাদানগুলি গ্লোবাল মেমরি থেকে রেজিস্টারে লোড করে।
একটি LDG.E 32টি লেনের প্রতিটিতে চার বাইট চায়। এটি পরিবেশন করতে প্রয়োজন চারটি 32-বাইট সেক্টর, একটি ক্যাশ লাইন, একটি ঠিকানা অনুবাদ, একটি ক্রসবার অতিক্রম, ছত্রিশটি L2 স্লাইসের একটি, এবং সর্বত্র মিস হলে, একটি DRAM চিপে একটি অ্যাক্টিভেট এবং চারটি কলাম রিড। আমাদের ওয়ার্পটি SM-এর চারটি সাব-পার্টিশনের একটিতে আরও এগারোটি আবাসিক ওয়ার্পের সাথে বাস করে। নির্দেশটি লোড/স্টোর ইউনিটে (LSU) ইস্যু করা হয়, যা অপকোড, সক্রিয় লেন মাস্ক, গণনা করা ঠিকানা এবং রেজিস্টার নম্বর কোয়ালেসারে পাঠায়।
কোয়ালেসার নির্ধারণ করে যে ওয়ার্প দ্বারা অনুরোধ করা 128 বাইটের জন্য 4টি ধারাবাহিক সেক্টর অনুরোধ প্রয়োজন। L1 ক্যাশে প্রবেশ করে, অনুরোধটি ভার্চুয়ালি অ্যাড্রেসড, 4-ওয়ে সেট-অ্যাসোসিয়েটিভ ক্যাশে পাঠানো হয়। লুকআপ চারটি স্লটের সেটে ট্যাগ তুলনা করে। যেহেতু এটি প্রথমবারের লোড, এটি মিস করে এবং আরও নিচে নামে। একটি L1 হিট প্রায় 15.4 ন্যানোসেকেন্ডে ফেরত আসে — 40 চক্র। ভার্চুয়াল মেমরির SM ছাড়ার আগে অনুবাদ প্রয়োজন। SM-এর TLB-তে ষোলটি সাম্প্রতিক অনুবাদ থাকে। মিসের খরচ প্রায় 4.4 ন্যানোসেকেন্ড — এগারোটি চক্র।
অনুবাদের পরে, প্রতিটি 128-বাইট লাইনের জন্য একটি অনুরোধ ক্রসবার অতিক্রম করে তার ফিজিক্যাল অ্যাড্রেসের ফাংশন দ্বারা নির্বাচিত 36টির মধ্যে একটি 2 MiB L2 স্লাইসের দিকে এগিয়ে যায়। সমস্ত স্লাইস সমান্তরালে পরিবেশন করতে পারে, যা সমষ্টিগত ব্যান্ডউইথকে একটি একক স্লাইসের 36 গুণ করে তোলে।
উৎস: Hacker News · সারাংশ: HeadlinesBriefing