فصل 6: شاخص‌های عملکرد GPU و تحلیل

تحلیل و بهینه‌سازی عملکرد برنامه‌های GPU بسیار مهم است تا بتوان از منابع سخت‌افزاری GPU به طور کارآمد و مؤثر استفاده کرد. در این فصل، ما شاخص‌های عملکرد GPU، ابزارهای پروفایل‌گیری و بهینه‌سازی، تکنیک‌های شناسایی گلوگاه‌های عملکردی و راهبردهای بهبود عملکرد GPU را بررسی خواهیم کرد.

پهنای باند، تأخیر و پهنای باند حافظه

سه شاخص اساسی برای ارزیابی عملکرد GPU عبارتند از پهنای باند، تأخیر و پهنای باند حافظه. درک این شاخص‌ها و پیامدهای آن‌ها برای تحلیل و بهینه‌سازی برنامه‌های GPU ضروری است.

پهنای باند

پهنای باند به تعداد عملیات یا وظایفی اشاره دارد که GPU می‌تواند در یک بازه زمانی مشخص انجام دهد. این شاخص معمولاً بر حسب عملیات‌های نقطه‌ای در ثانیه (FLOPS) یا دستورالعمل در ثانیه (IPS) اندازه‌گیری می‌شود. GPU‌ها برای دستیابی به پهنای باند بالا طراحی شده‌اند و با بهره‌گیری از موازی‌سازی و اجرای همزمان تعداد زیادی از رشته‌ها به این هدف می‌رسند.

پهنای باند نظری بیشینه یک GPU را می‌توان با استفاده از فرمول زیر محاسبه کرد:

پهنای باند بیشینه (FLOPS) = تعداد هسته‌های CUDA × فرکانس ساعت × FLOPS در هر هسته CUDA در هر چرخه

به عنوان مثال، GPU NVIDIA GeForce RTX 2080 Ti دارای 4352 هسته CUDA، فرکانس پایه 1350 مگاهرتز و هر هسته CUDA می‌تواند 2 عملیات نقطه‌ای در هر چرخه (FMA - ضرب-جمع ادغام‌شده) انجام دهد. بنابراین، پهنای باند نظری بیشینه آن برابر است با:

پهنای باند بیشینه (FLOPS) = 4352 × 1350 مگاهرتز × 2 = 11.75 TFLOPS

با این حال، دستیابی به پهنای باند نظری بیشینه در عمل چالش‌برانگیز است، زیرا عوامل مختلفی مانند الگوهای دسترسی به حافظه، واگرایی شاخه و محدودیت‌های منابع در آن دخیل هستند.

تأخیر

تأخیر به زمان مورد نیاز برای اتمام یک عملیات یا وظیفه واحد اشاره دارد. در زمینه GPU‌ها، تأخیر معمولاً با عملیات دسترسی به حافظه مرتبط است. GPU‌ها دارای سیستم حافظه سلسله‌مراتبی هستند و دسترسی به داده‌ها از سطوح مختلف این سلسله‌مراتب حافظه، تأخیرهای متفاوتی را به همراه دارد.اینجا ترجمه فارسی فایل مارک‌داون است:

انواع تأخیرهای معمول برای سطوح مختلف حافظه در یک GPU به شرح زیر است:

ثبات‌ها: 0-1 سیکل
حافظه اشتراکی: 1-2 سیکل
کش L1: 20-30 سیکل
کش L2: 200-300 سیکل
حافظه جهانی (DRAM): 400-800 سیکل

تأخیر می‌تواند تأثیر قابل توجهی بر عملکرد GPU داشته باشد، به ویژه زمانی که بین عملیات وابستگی وجود دارد یا زمانی که رشته‌ها در انتظار بازیابی داده از حافظه هستند. تکنیک‌هایی مانند پنهان‌سازی تأخیر، پیش‌بارگیری و کش‌بندی می‌توانند تأثیر تأخیر بر عملکرد GPU را کاهش دهند.

پهنای باند حافظه

پهنای باند حافظه به نرخی اشاره دارد که در آن داده می‌تواند بین GPU و زیرسیستم حافظه آن منتقل شود. این معمولاً با واحد بایت بر ثانیه (B/s) یا گیگابایت بر ثانیه (GB/s) اندازه‌گیری می‌شود. GPU‌ها از رابط‌های حافظه پرپهنای باند مانند GDDR6 یا HBM2 برای پشتیبانی از ماهیت داده‌بر کاربردهای گرافیکی و محاسباتی استفاده می‌کنند.

پهنای باند حافظه اوج نظری یک GPU را می‌توان با استفاده از فرمول زیر محاسبه کرد:

پهنای باند حافظه اوج (GB/s) = فرکانس ساعت حافظه × عرض باس حافظه ÷ 8

به عنوان مثال، یک GPU NVIDIA GeForce RTX 2080 Ti دارای فرکانس ساعت حافظه 7000 مگاهرتز (موثر) و عرض باس حافظه 352 بیت است. بنابراین، پهنای باند حافظه اوج نظری آن برابر است با:

پهنای باند حافظه اوج (GB/s) = 7000 مگاهرتز × 352 بیت ÷ 8 = 616 گیگابایت بر ثانیه

پهنای باند حافظه عامل بحرانی در عملکرد GPU است، زیرا بسیاری از کاربردهای GPU محدود به حافظه هستند، به این معنی که عملکرد آنها توسط نرخ انتقال داده بین GPU و حافظه محدود می‌شود. بهینه‌سازی الگوهای دسترسی به حافظه، کاهش انتقال داده و بهره‌گیری از سلسله مراتب حافظه می‌تواند به بهبود استفاده از پهنای باند حافظه کمک کند.

ابزارهای پروفایل‌گیری و بهینه‌سازی عملکرد

ابزارهای پروفایل‌گیری و بهینه‌سازی عملکرد برای تحلیل رفتار برنامه‌های کاربردی GPU، شناسایی نقاط ضعف عملکرد و هدایت تلاش‌های بهینه‌سازی ضروری هستند. این ابزارها بینش‌هایی در مورد جنبه‌های مختلف عملکرد GPU مانند زمان اجرای کرنل، دسترسی به حافظه و غیره ارائه می‌دهند.فایل مارک‌داون این است: الگوهای ESS، اشغال و استفاده از منابع

برخی از ابزارهای محبوب پروفایل‌گیری و بهینه‌سازی عملکرد برای GPU‌ها شامل موارد زیر است:

NVIDIA Visual Profiler (nvvp): یک ابزار پروفایل‌گیری گرافیکی که نمای جامعی از عملکرد برنامه‌های GPU ارائه می‌دهد. به توسعه‌دهندگان امکان می‌دهد اجرای کرنل، انتقال حافظه و تماس‌های API را تحلیل کنند و توصیه‌هایی برای بهینه‌سازی ارائه می‌دهد.
NVIDIA Nsight: یک محیط توسعه یکپارچه (IDE) که قابلیت‌های پروفایل‌گیری و اشکال‌زدایی برای برنامه‌های GPU را شامل می‌شود. از زبان‌ها و چارچوب‌های مختلف برنامه‌نویسی مانند CUDA، OpenCL و OpenACC پشتیبانی می‌کند.
NVIDIA Nsight Compute: یک ابزار پروفایل‌گیری مستقل که بر تحلیل عملکرد کرنل GPU متمرکز است. شاخص‌های عملکردی مفصلی مانند پهنای باند دستورالعمل، کارایی حافظه و اشغال را ارائه می‌دهد و به شناسایی گلوگاه‌های عملکردی در سطح کد منبع کمک می‌کند.
AMD Radeon GPU Profiler (RGP): یک ابزار پروفایل‌گیری برای GPU‌های AMD که داده‌های عملکردی را برای برنامه‌های DirectX، Vulkan و OpenCL ضبط و نمایش می‌دهد. به بینش‌هایی در مورد استفاده از GPU، استفاده از حافظه و انسداد خط‌لوله دست می‌یابد.
AMD Radeon GPU Analyzer (RGA): یک ابزار تحلیل استاتیک که کد شیدر GPU را تحلیل می‌کند و پیش‌بینی‌های عملکردی، استفاده از منابع و پیشنهادات بهینه‌سازی را ارائه می‌دهد.

این ابزارها معمولاً با ابزارگذاری کد برنامه GPU، جمع‌آوری داده‌های عملکردی در طول اجرا و ارائه داده‌ها در قالب دوستانه برای کاربر کار می‌کنند. آنها معمولاً نماهای زمانی، شمارنده‌های عملکرد و همبستگی با کد منبع را ارائه می‌دهند تا به توسعه‌دهندگان در شناسایی مشکلات عملکردی و بهینه‌سازی کد کمک کنند.

مثال: پروفایل‌گیری از یک برنامه CUDA با استفاده از NVIDIA Visual Profiler (nvvp)

برنامه CUDA را با فعال‌سازی پروفایل‌گیری بسازید:
```
nvcc -o myapp myapp.cu -lineinfo
```
برنامه را با پروفایل‌گیری اجرا کنید:
```
nvprof ./myapp
```
Visual Profiler را باز کنید:
```
nvvp
```
داده‌های پروفایل‌گیری تولید شده را وارد کنید.Here is the Persian translation of the provided markdown file, with the code comments translated:
تحلیل نمای زمانی، عملکرد هسته، انتقال حافظه و تماس های API.
شناسایی موانع عملکردی و بهینه سازی کد بر اساس توصیه های پروفایلر.

شناسایی موانع عملکردی

شناسایی موانع عملکردی برای بهینه سازی برنامه های GPU بسیار مهم است. موانع عملکردی می توانند از عوامل مختلفی مانند الگوهای دسترسی ناکارآمد به حافظه، اشغال پایین، انشعاب شاخه و محدودیت های منابع ناشی شوند. برخی از تکنیک های رایج برای شناسایی موانع عملکردی عبارتند از:

پروفایل کردن: استفاده از ابزارهای پروفایل برای اندازه گیری زمان اجرای هسته، زمان انتقال حافظه و هزینه API می تواند به شناسایی بخش هایی از برنامه که بیشترین زمان و منابع را مصرف می کنند، کمک کند.
تحلیل اشغال: اشغال به نسبت وارپ های فعال به حداکثر تعداد وارپ های پشتیبانی شده توسط یک GPU اشاره دارد. اشغال پایین می تواند نشانگر عدم استفاده کامل از منابع GPU باشد و ممکن است نیاز به بهینه سازی ابعاد بلوک و شبکه یا کاهش استفاده از ثبات و حافظه مشترک را نشان دهد.
بررسی الگوهای دسترسی به حافظه: الگوهای دسترسی ناکارآمد به حافظه، مانند دسترسی های غیر هماهنگ به حافظه یا دسترسی های مکرر به حافظه جهانی، می تواند تأثیر قابل توجهی بر عملکرد GPU داشته باشد. تحلیل الگوهای دسترسی به حافظه با استفاده از ابزارهای پروفایل می تواند فرصت هایی برای بهینه سازی مانند استفاده از حافظه مشترک یا بهبود محلی سازی داده را شناسایی کند.
بررسی انشعاب شاخه: انشعاب شاخه زمانی رخ می دهد که رشته های درون یک وارپ مسیرهای اجرایی متفاوتی را به دلیل دستورات شرطی انتخاب می کنند. شاخه های متفاوت می توانند به سریالی شدن و کاهش عملکرد منجر شوند. شناسایی و به حداقل رساندن انشعاب شاخه می تواند به بهبود عملکرد GPU کمک کند.
نظارت بر استفاده از منابع: GPU ها منابع محدودی مانند ثبات، حافظه مشترک و بلوک های رشته دارند. نظارت بر استفاده از منابع با استفاده از ابزارهای پروفایل می تواند به شناسایی موانع مربوط به منابع کمک کند و راهنمایی برای تلاش های بهینه سازی مانند کاهش استفاده از ثبات را ارائه دهد.اینجا ترجمه فارسی فایل مارک‌داون است. برای کد، فقط نظرات را ترجمه کرده‌ایم، نه خود کد:

مثال: شناسایی یک گلوگاه دسترسی به حافظه با استفاده از NVIDIA Nsight Compute

پروفایل برنامه CUDA را با استفاده از Nsight Compute انجام دهید:
```
ncu -o profile.ncu-rep ./myapp
```
گزارش پروفایل تولید شده را در Nsight Compute باز کنید.
بخش "Memory Workload Analysis" را تجزیه و تحلیل کنید تا الگوهای دسترسی ناکارآمد به حافظه مانند دسترسی‌های غیر هم‌پوشان یا استفاده زیاد از حافظه جهانی را شناسایی کنید.
الگوهای دسترسی به حافظه را بر اساس بینش‌های ارائه شده توسط Nsight Compute بهینه کنید، مانند استفاده از حافظه اشتراکی یا بهبود محلی داده‌ها.

راهبردهای بهبود عملکرد GPU

پس از شناسایی گلوگاه‌های عملکرد، راهبردهای مختلفی را می‌توان برای بهبود عملکرد GPU به کار گرفت. برخی از راهبردهای بهینه‌سازی رایج عبارتند از:

افزایش موازی‌سازی: اطمینان حاصل کنید که برنامه به تعداد کافی وظایف موازی تجزیه شده است تا از منابع GPU به طور کامل استفاده شود. این ممکن است شامل تنظیم ابعاد بلوک و شبکه، استفاده از جریان‌ها برای اجرای همزمان یا بهره‌گیری از موازی‌سازی سطح وظیفه باشد.
بهینه‌سازی الگوهای دسترسی به حافظه: کارایی دسترسی به حافظه را با کاهش دسترسی‌های به حافظه جهانی، استفاده از حافظه اشتراکی برای داده‌های پرتکرار و اطمینان از دسترسی‌های هم‌پوشان بهبود بخشید. تکنیک‌هایی مانند تقسیم‌بندی حافظه، تبدیل چیدمان داده و حافظه پنهان می‌توانند در بهینه‌سازی عملکرد حافظه کمک کنند.
کاهش واگرایی شاخه: واگرایی شاخه را با بازسازی کد برای اجتناب از شاخه‌های واگرا در داخل یک وارپ به حداقل برسانید. تکنیک‌هایی مانند پیش‌بینی شاخه، شاخه‌بندی وابسته به داده و برنامه‌نویسی سطح وارپ می‌توانند در کاهش تأثیر واگرایی شاخه کمک کنند.
بهره‌گیری از سلسله مراتب حافظه: به طور مؤثر از سلسله مراتب حافظه GPU استفاده کنید با حداکثر کردن استفاده از ثبات‌ها و حافظه اشتراکی برای داده‌های پرتکرار. از حافظه بافر و حافظه ثابت برای داده‌های فقط‌خوان که دارای محلی‌سازی فضایی هستند یا به طور یکنواخت در میان رشته‌ها دسترسی پیدا می‌کنند، استفاده کنید.
همپوشانی محاسبه و حافظه: محاسبات و انتقال داده را به طور همزمان انجام دهید تا پهنای باند حافظه را بهینه کنید. از تکنیک‌هایی مانند پیش‌بارگذاری داده، همپوشانی انتقال و محاسبه و استفاده از جریان‌ها برای همپوشانی استفاده کنید.اینجا ترجمه فارسی فایل مارک‌داون است، با توجه به اینکه کد برنامه‌نویسی نباید ترجمه شود، فقط توضیحات کد به فارسی ترجمه شده‌اند:

انتقال حافظه: پنهان کردن تأخیر انتقال حافظه با همپوشانی محاسبه با انتقالات حافظه با استفاده از جریان‌های CUDA یا صف‌های فرمان OpenCL. این امکان را به GPU می‌دهد تا در حالی که داده‌ها بین حافظه میزبان و دستگاه در حال انتقال هستند، محاسبات را انجام دهد.

تنظیم پارامترهای راه‌اندازی کرنل: با آزمایش بلوک‌ها و شبکه‌های مختلف، پیکربندی بهینه برای هر کرنل را پیدا کنید. پارامترهای راه‌اندازی بهینه به عواملی مانند تعداد ثبات‌های استفاده شده در هر رشته، استفاده از حافظه اشتراکی و ویژگی‌های معماری GPU بستگی دارد.
کاهش انتقال داده بین میزبان و دستگاه: مقدار داده منتقل شده بین میزبان (CPU) و دستگاه (GPU) را کاهش دهید با انجام هرچه بیشتر محاسبات در GPU. انتقال‌های کوچک را در انتقال‌های بزرگ‌تر قرار دهید تا هزینه هر انتقال کاهش یابد.
استفاده از عملیات غیرهمزمان: از عملیات غیرهمزمان مانند کپی‌های حافظه و راه‌اندازی کرنل‌های غیرهمزمان استفاده کنید تا محاسبه و ارتباطات را همپوشانی دهید. این امکان را به CPU می‌دهد تا در حالی که GPU در حال اجرا است، سایر وظایف را انجام دهد و در نتیجه عملکرد کلی برنامه بهبود یابد.

مثال: بهینه‌سازی الگوهای دسترسی به حافظه با استفاده از حافظه اشتراکی در CUDA

کد اصلی با دسترسی‌های ناکارآمد به حافظه جهانی:

__global__ void myKernel(float* data, int n) {
    // شناسه رشته
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < n) {
        float result = 0.0f;
        // محاسبه نتیجه با دسترسی به حافظه جهانی
        for (int i = 0; i < n; i++) {
            result += data[tid] * data[i];
        }
        data[tid] = result;
    }
}

کد بهینه‌شده با استفاده از حافظه اشتراکی:

__global__ void myKernel(float* data, int n) {
    // حافظه اشتراکی برای ذخیره داده‌ها
    __shared__ float sharedData[256];
    // شناسه رشته
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    // شناسه محلی در بلوک
    int localIdx = threadIdx.x;
 
    // کپی داده‌ها به حافظه اشتراکی
    if (tid < n) {
        sharedData[localIdx] = data[tid];
    }
    __syncthreads();
 
    if (tid < n) {
        float result = 0.0f;
        // محاسبه نتیجه با دسترسی به حافظه اشتراکی
        for (int i = 0; i < blockDim.x; i++) {
            result += sharedData[localIdx] * sharedData[i];
        }
        datاینجا ترجمه فارسی فایل مارک‌داون است:
 
a[tid] = result;
    }
}

در کد بهینه‌شده، داده‌های ورودی ابتدا در حافظه اشتراکی بارگذاری می‌شوند، که نسبت به حافظه جهانی دارای تأخیر بسیار کمتری است. محاسبات سپس با استفاده از حافظه اشتراکی انجام می‌شوند، که باعث کاهش تعداد دسترسی‌های به حافظه جهانی و بهبود عملکرد می‌شود.

نتیجه‌گیری

تحلیل و بهینه‌سازی عملکرد GPU برای توسعه برنامه‌های GPU کارآمد و با عملکرد بالا ضروری است. با درک شاخص‌های کلیدی عملکرد مانند پهنای باند، تأخیر و پهنای باند حافظه، توسعه‌دهندگان می‌توانند تصمیمات آگاهانه‌ای در مورد بهینه‌سازی کد خود اتخاذ کنند.

ابزارهای پروفایل‌گیری و بهینه‌سازی عملکرد نقش حیاتی در شناسایی گلوگاه‌های عملکرد و هدایت تلاش‌های بهینه‌سازی ایفا می‌کنند. این ابزارها بینش‌های ارزشمندی در مورد اجرای هسته، الگوهای دسترسی به حافظه، اشغال و استفاده از منابع ارائه می‌دهند، که به توسعه‌دهندگان امکان می‌دهد تلاش‌های بهینه‌سازی خود را بر مناطق بحرانی متمرکز کنند.

استراتژی‌های بهینه‌سازی رایج شامل حداکثرسازی موازی‌سازی، بهینه‌سازی الگوهای دسترسی به حافظه، کاهش واگرایی شاخه و غیره است.

برخی استراتژی‌های رایج برای بهینه‌سازی عملکرد GPU به شرح زیر است:

کاهش واگرایی شاخه: جریان کنترل واگرا در یک وارپ/موج می‌تواند به سریالی‌سازی و کاهش کارایی SIMD منجر شود. الگوریتم‌ها باید به گونه‌ای طراحی شوند که واگرایی شاخه را به حداقل برسانند. تکنیک‌هایی مانند پیش‌بینی شاخه، شاخه‌بندی وابسته به داده و برنامه‌نویسی سطح وارپ می‌توانند تأثیر واگرایی شاخه را کاهش دهند.
بهره‌گیری از سلسله‌مراتب حافظه: به طور مؤثر از سلسله‌مراتب حافظه GPU استفاده کنید با حداکثرسازی استفاده از ثبات‌ها و حافظه اشتراکی برای داده‌های دسترسی‌شده مکرر. از حافظه بافر و حافظه ثابت برای داده‌های فقط‌خواندنی که دارای محلی‌گرایی فضایی هستند یا به طور یکنواخت توسط رشته‌ها دسترسی می‌شوند، استفاده کنید.
همپوشانی محاسبات و انتقال حافظه: تأخیر انتقال حافظه را با همپوشانی محاسبات با انتقال حافظه با استفاده از جریان‌های CUDA یا صف‌های فرمان OpenCL مخفی کنید. این امکان را فراهم می‌کند کهHere is the Persian translation of the provided markdown file, with the code comments translated:
تنظیم پارامترهای راه‌اندازی کرنل: با آزمایش کردن اندازه‌های مختلف بلوک و شبکه، پیکربندی بهینه برای هر کرنل را پیدا کنید. پارامترهای بهینه راه‌اندازی به عواملی مانند تعداد ثبات‌های استفاده شده در هر رشته، استفاده از حافظه مشترک و ویژگی‌های معماری GPU بستگی دارد.
کاهش انتقال داده بین میزبان و دستگاه: میزان داده منتقل شده بین میزبان (CPU) و دستگاه (GPU) را کاهش دهید با انجام محاسبات بیشتر روی GPU. انتقال‌های کوچک را در انتقال‌های بزرگ‌تر ادغام کنید تا هزینه هر انتقال را کاهش دهید.
استفاده از عملیات‌های غیرهمزمان: از عملیات‌های غیرهمزمان مانند کپی‌های حافظه و راه‌اندازی کرنل‌های غیرهمزمان استفاده کنید تا محاسبه و ارتباطات را همپوشانی دهید. این به CPU اجازه می‌دهد تا در حالی که GPU در حال اجرا است، سایر وظایف را انجام دهد و در نتیجه عملکرد کلی برنامه بهبود می‌یابد.

مثال: بهینه‌سازی الگوهای دسترسی به حافظه با استفاده از حافظه مشترک در CUDA

کد اصلی با دسترسی‌های ناکارآمد به حافظه جهانی:

__global__ void myKernel(float* data, int n) {
    // شناسه رشته را محاسبه کن
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < n) {
        float result = 0.0f;
        // محاسبه نتیجه با استفاده از حافظه جهانی
        for (int i = 0; i < n; i++) {
            result += data[tid] * data[i];
        }
        data[tid] = result;
    }
}

کد بهینه‌شده با استفاده از حافظه مشترک:

__global__ void myKernel(float* data, int n) {
    // ایجاد یک آرایه حافظه مشترک
    __shared__ float sharedData[256];
    // شناسه رشته را محاسبه کن
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    // شناسه محلی رشته را محاسبه کن
    int localIdx = threadIdx.x;
 
    // بارگذاری داده‌ها در حافظه مشترک
    if (tid < n) {
        sharedData[localIdx] = data[tid];
    }
    __syncthreads();
 
    // محاسبه نتیجه با استفاده از حافظه مشترک
    if (tid < n) {
        float result = 0.0f;
        for (int i = 0; i < blockDim.x; i++) {
            result += sharedData[localIdx] * sharedData[i];
        }
        data[tid] = result;
    }
}

در کد بهینه‌شده، داده‌های ورودی ابتدا در حافظه مشترک بارگذاری می‌شوند که نسبت به حافظه جهانی دارای زمان دسترسی کمتری است.اینجا ترجمه فارسی فایل مارک‌داون است:

حافظه جهانی

محاسبه در ابتدا با استفاده از حافظه جهانی انجام می‌شود. سپس محاسبه با استفاده از حافظه اشتراکی انجام می‌شود که باعث کاهش تعداد دسترسی‌های به حافظه جهانی و بهبود عملکرد می‌شود.

// این تابع محاسبه را با استفاده از حافظه اشتراکی انجام می‌دهد
__global__ void kernel_function(int *input, int *output, int size) {
    // شناسه بلوک و شناسه رشته را به دست می‌آوریم
    int block_id = blockIdx.x;
    int thread_id = threadIdx.x;
 
    // محاسبه را در حافظه اشتراکی انجام می‌دهیم
    __shared__ int shared_memory[BLOCK_SIZE];
    shared_memory[thread_id] = input[block_id * BLOCK_SIZE + thread_id];
    __syncthreads();
 
    // محاسبه را انجام می‌دهیم
    int result = 0;
    for (int i = 0; i < BLOCK_SIZE; i++) {
        result += shared_memory[i];
    }
 
    // نتیجه را در خروجی ذخیره می‌کنیم
    output[block_id] = result;
}

Chapter 5 Gpu Memory System Design Chapter 7 Streaming Multiprocessor Design