Mô phỏng song song sử dụng khối xử lý đồ họa GPGPU.

- 36 2.1.4 Các loại bộ nhớ của GPU.
- 46 2.2.8 Các lệnh với bộ nhớ.
- Các thao tác thu hồi và cấp phát bộ nhớ với CUDA.
- Cấu trúc bộ nhớ thiết bị GPU.
- Các thiết bị đồ họa trở thành các bộ nhớ đệm cho các phần tử ảnh có khả năng truy nhập được, gọi là các điểm ảnh (pixel), có thể được ánh xạ trực tiếp bởi thiết bị vào trong màn hình quét mành bằng việc làm tươi trong một khoảng tần số nào đó.
- Một CPU phổ biến được cấu thành bởi một bộ nhớ cache lớn, một hoặc nhiều khối xử lý điều khiển, một hoặc nhiều nhân tính toán (ALU).
- Các CPU được nối với một bộ nhớ ngoài (RAM).
- Ngược lại, một GPU được xây dựng với nhiều nhân, mỗi nhân có một bộ nhớ cache nhỏ, vài bộ xử lý điều khiển cho các nhân này và cho các luồng ứng dụng.
- Một bộ nhớ lớn có thể truy nhập được nối ngay sát các GPU trên các thiết bị đồ họa.
- Mỗi SM được kết nối với bộ nhớ toàn cục của card.
- Để sử dụng tính toán, một khối bộ nhớ bất biến được tách khỏi bộ nhớ toàn cục.
- Mỗi SM được cấu thành bởi nhiều bộ xử lý dòng (Streaming Processors -SP), một bộ nhớ cache và một bộ xử lý điều khiển.
- Tồn tại một khối bộ nhớ nhỏ để sử dụng cho dữ liệu tính toán chung, nó có thể truy nhập đồng thời bởi tất cả các SP khác.
- Bộ nhớ này được xem như là bộ nhớ chia sẻ của CUDA.
- Một bộ quản lý thực thi luồng với bộ nhớ cache lệnh được cài đặt trên đỉnh của các SM, sử dụng để lập lịch và điều khiển các luồng, các bộ xử lý và phát triển ứng dụng.
- Bộ nhớ toàn cục được nạp dữ liệu hiện thời để phục vụ cho các chương trình tính toán đổ bóng.
- Các tham số phần cứng khác nhau có một số lượng bộ nhớ thiết bị toàn cục, các nhân, các tần số bộ nhớ và một số lượng SM được cài đặt sẵn.
- Mỗi CU có bộ nhớ cục bộ tương đương với bộ nhớ chia sẻ trong CUDA.
- Tất cả các CU đều được kết nối với bộ nhớ thiết bị toàn cục.
- Các chương trình mảnh tính giá trị của các mảnh bằng cách kết hợp của phép toán toán học và bộ nhớ toàn cục đọc từ bộ nhớ kết cấu toàn cục.
- 3) Các giá trị cho mỗi luồng được tính bằng cách kết hợp các phép toán toán học và cả truy cập "thu thập" (đọc) và "scatter" (ghi) bộ nhớ toàn cục.
- 4) Các vùng đệm chứa kết quả trong bộ nhớ toàn cục sau đó có thể được sử dụng như là một đầu vào của tính toán sau đó.
- Mô hình lập trình dòng tổ chức chương trình để thực hiện song song và cho phép giao tiếp hiệu quả và truyền dữ liệu đồng thời phù hợp với các nguồn lực xử lý song song và hệ thống bộ nhớ có sẵn trên GPU.
- CUDA cũng khai thác các nguồn tài nguyên phần cứng nhiều hơn Brook, làm lộ nhiều cấp độ của bộ nhớ hệ thống phân cấp.
- các thanh ghi theo từng luồng, bộ nhớ chia sẻ nhanh chóng giữa các luồng trong một khối, bộ nhớ bo mạch, và bộ nhớ máy chủ.
- Thuật toán đã phải dành sự quan tâm đặc biệt tới việc truyền dữ liệu để đảm bảo rằng các bộ nhớ L1 cache, L2 cache và bộ nhớ đệm của hệ thống được tận dụng một cách tối đa.
- Do cấu trúc bộ nhớ 30 cache đơn giản hơn nhiều nên cấu trúc thuật toán GPGPU không phức tạp như các thuật toán trên CPU.
- Kích thước của cây con được lựa chọn một cách cẩn thận để dữ liệu nút cây con vừa vặn với khối bộ nhớ chia sẻ trên GPU và bộ nhớ truy nhập tối thiểu cần để xử lý truy vấn trong cây con.
- CUDA cung cấp cách đánh địa chỉ DRAM thường dùng như mô tả trong hình 2.2 cho việc lập trình linh hoạt hơn, bao gồm cả thao tác cấp phát và thu hồi bộ nhớ.
- [9] CUDA có đặc tính lưu dữ liệu đệm song song và bộ nhớ chia sẻ on-chip với tốc độ đọc ghi rất cao, các luồng dùng bộ nhớ này để chia sẻ dữ liệu với nhau.
- Như mô tả trong hình 2.3, ứng dụng có thể đạt kết quả tốt với việc tối thiểu hóa việc lấy/trả dữ liệu từ DRAM, từ đó giảm phụ thuộc băng thông truyền bộ nhớ DRAM.
- Nhờ có xử lý đồ họa tự nhiên, các sản phẩm được phát hành mới đây cung cấp các bộ xử lý song song tốc độ cao với băng thông bộ nhớ cao và sức mạnh tính toán lên tới hàng teraflops trên giây.
- Các bộ nhớ chia sẻ: CUDA cho phép các luồng sử dụng 6 kiểu bộ nhớ khác nhau được thiết kế để đáp ứng cho các yêu cầu khác nhau.
- Như mô tả ở hình 2.5, các SM cũng có 2 bộ chức năng đặc biệt: một bộ dòng lệnh đa luồng và một bộ nhớ chia sẻ on-chip.
- [9] 2.1.4 Các loại bộ nhớ của GPU Việc lập trình song song GPU một cách có hiệu quả được dựa trên việc sử dụng bộ nhớ thiết bị một cách thích hợp.
- Vì vậy, chi tiết về bộ nhớ thiết bị, gồm có 6 bộ nhớ khác nhau với các đặc điểm khác nhau cần được trình bày ở đây.
- Các không gian bộ nhớ được minh họa ở hình 2.7 đều rất phù hợp với 38 các yêu cầu khác nhau.
- Mục này đưa ra thông tin quan tâm chi tiết đến các không gian bộ nhớ này và cách sử dụng chúng có hiệu quả.
- Tuy vậy, bộ nhớ này chỉ có thể truy nhập bởi một luồng đơn và có không gian hạn chế.
- Bộ nhớ cục bộ (Local Memory): được sử dụng như không gian thêm vào khi các thanh ghi không đáp ứng được các yêu cầu.
- Vì không gian bộ nhớ này không nằm trên chip (off-chip) và không phải là cache (không phải bộ nhớ truy nhập nhanh), nên nó bị tối ưu hóa băng thông.
- Bộ nhớ chia sẻ (Shared Memory): được thiết kế để thi hành công việc trên GPU.
- Bộ nhớ on-chip này nằm gần với bộ xử lý dòng.
- Bộ nhớ chia sẻ có thể truy nhập bởi các luồng bên trong một khối và nhanh hơn bộ nhớ toàn cục.
- Hiện tại, không gian bộ nhớ chia sẻ bị giới hạn bởi 16KB/block.
- Không gian bộ nhớ này có thể truy nhập bởi tất cả các luồng trong một khối.
- Mặc dù được thiết kế để truy nhập nhanh, nhưng 39 bộ nhớ chia sẻ cũng chỉ cung cấp một khả năng giới hạn.
- Bởi vậy, khi số luồng được chọn vượt quá giới hạn tối đa là 512 luồng / 1 block, kích thước bộ nhớ là một vấn đề cần phải tính đến khi thực thi trên GPGPU.
- Thêm vào đó để truy nhập nhanh, bộ nhớ chia sẻ cũng cho phép đồng bộ hóa, đỏi hỏi phải lưu giữ các luồng sang một bên trong quá trình đợi các luồng khác thực hiện xong.
- Việc thay thế bộ nhớ truy nhập toàn cục bởi bộ nhớ chia sẻ đã tiết kiệm nhiều băng thông cho bộ nhớ toàn cục.
- Do đó sẽ rất có ý nghĩa nếu cố gắng xây dựng lại các thuật toán để tận dụng tối đa những lợi ích của các bộ nhớ chia sẻ.
- Bộ nhớ toàn cục: (Global Memory): không gian bộ nhớ off-chip này được xem xét như là một không gian bộ nhớ linh hoạt nhất.
- Tuy nhiên bộ nhớ toàn cục tiềm ẩn sự không hiệu quả với việc truy nhập.
- Vì không gian bộ nhớ này không phải là cache (không truy nhập nhanh) nên việc truy nhập bộ nhớ này được yêu cầu hiệu suất tốt hơn.
- Nó cũng không yêu cầu không gian bộ nhớ quá lớn.
- Bộ nhớ cố định (Constant Memory): đây là không gian bộ nhớ chỉ đọc và tương tự như bộ nhớ toàn cục, nó cũng là off-chip.
- Tuy nhiên, bộ nhớ cố định được truy cập nhanh và cung cấp hiệu năng tốt hơn.
- Hiện tại, không gian bộ nhớ cố định được giới hạn ở 64KB.
- Bộ nhớ kết cấu (Texture Memory): trên thực tế nó là một bộ nhớ toàn cục chỉ đọc.
- So sánh với bộ nhớ toàn cục thì không gian bộ nhớ này cung cấp hiệu năng tốt hơn với cấu trúc cache (truy nhập nhanh).
- Bộ nhớ kết cấu cho phép biểu thị dữ liệu ở một trong các tổ hợp dưới đây: o Không gian: 1,2 hoặc 3 chiều.
- Việc truy nhập bộ nhớ chỉ đọc được sử dụng thông qua bộ nhớ kết cấu mỗi khi có thể để cung cấp hiệu năng tốt hơn.
- Các chức năng sao chép bộ nhớ với hậu tố Async.
- Các chức năng sao chép bộ nhớ host tới thiết bị và thiết bị tới host.
- Các chức năng thiết lập bộ nhớ.
- Cũng có thể sử dụng đồng thời các sao chép bộ nhớ thiết bị và host page-locked với các thi hành trong nhân.
- Chiến lược này cực tiểu hóa việc truyền dữ liệu giữa host và thiết bị và sử dụng các bộ nhớ thiết bị một cách thích hợp.
- Nếu có thể, bộ nhớ chia sẻ hoặc bộ nhớ kết cấu truy nhập nhanh hiệu năng cao cần được lựa chọn thay cho bộ nhớ toàn cục.
- Từ khóa phạm vi kiểu biến: cho phép đặc tả vị trí bộ nhớ trên thiết bị của một biến.
- Tham số của hàm toàn cục hiện đang được truyền qua Bộ nhớ chia sẻ với thiết bị và giới hạn độ lớn 256 byte.
- Biến chia sẻ lựa chọn sử dụng với các thiết bị khác, miêu tả một biến tồn tại trong một không gian bộ nhớ chia sẻ của một luồng, có thời gian sống của một khối và chỉ có thể truy nhập từ tất cả các chủ thể trong khối.
- Trình biên dịch không bị ràng buộc để tối ưu hóa những lần đọc, ghi vào bộ nhớ chia sẻ miễn là những câu lệnh trước đó được đáp ứng.
- Khi khai báo một biến trong bộ nhớ chia sẻ như một mảng mở rộng như: extern __shared__ float shared.
- Tất cả các biến nếu khai báo trong thời điểm này, bắt đầu tại cùng một địa chỉ trong bộ nhớ.
- thì trong bộ nhớ chia sẻ động được tạo có thể khai báo và khởi tạo các mảng theo cách sau: extern __shared__ char array.
- Ns là một kiểu size_t và xác định số byte trong bộ nhớ chia sẻ.
- Ns cho phép khai báo động trên mỗi khối cho lời gọi ngoài việc cấp phát bộ nhớ 46 tĩnh.
- Việc cấp phát bộ nhớ động sử dụng bởi bất kỳ biến khai báo như là một mảng mở rộng, Ns là một đối số tùy chọn mặc định là 0.
- 2.2.8 Các lệnh với bộ nhớ Các lệnh bộ nhớ bao gồm các lệnh đọc và ghi tới vùng nhớ chia sẻ hoặc vùng nhớ toàn cục.
- Bộ đa xử lý mất bốn chu kỳ đồng hồ để đưa ra một lệnh bộ nhớ cho warp.
- Khi truy cập bộ nhớ toàn cục, thêm vào đó sẽ mất độ trễ là 400 tới 600 chu kỳ đồng hồ.
- 48 Độ trễ bộ nhớ toàn cục nhiều đến mức có thể ẩn bởi bộ lập lịch luồng nếu có các lệnh số học không phụ thuộc có thể ban hành trong khi chờ truy cập bộ nhớ kết thúc.
- Bộ nhớ toàn cục Không gian nhớ toàn cục không được lưu vào bộ nhớ đệm.
- Vì thế điều quan trọng là truy xuất đúng cách để có được băng thông tối đa, đặc biệt là chi phí cho việc truy cập bộ nhớ thiết bị.
- Đầu tiên, thiết bị có thể đọc cùng lúc 32, 64 hoặc 128 bit cùng lúc từ bộ nhớ toàn cục vào thanh ghi với một câu lệnh.
- Việc xếp bộ nhớ của các biến như vậy được làm tự động cho các kiểu có sẵn như float2 hoặc float4.
- Với các kiểu cấu trúc, kích thước và xếp bộ nhớ có thể được thi hành bởi trình biên dịch sử dụng những chỉ thị cụ thể như __align__(8) hoặc __align__(16), ví dụ: struct __align__(8.
- Địa chỉ bộ nhớ toàn cục được truy xuất đồng thời bởi từng luồng trong suốt việc thi hành của một lệnh máy đọc hoặc ghi nên được xắp xếp để việc truy cập bộ nhớ có thể kết hợp thành việc truy xuất một vùng nhớ liên tục duy nhất.
- Với HalfWarpBaseAddress là kiểu con trỏ type* tuân theo cách dàn bộ nhớ.
- cudaMallocPitch() và các hàm sao chép bộ nhớ có liên quan được mô tả trong D.3 và E.6 cho phép nguời dùng phát triển các dòng lệnh không phụ thuộc vào phần cứng để cấp phát các mảng thỏa mãn các điều kiện đó.
- Đọc bộ nhớ thiết bị qua việc lấy kết cấu có thể là một lựa chọn nâng cao để đọc bộ nhớ thiết bị từ bộ nhớ toàn cục hoặc bộ nhớ hằng số.
- Bộ nhớ chia sẻ Vì bộ nhớ chia sẻ gắn trên bộ vi xử lý, nên không gian bộ nhớ chia sẻ nhanh hơn nhiều so với các không gian bộ nhớ cục bộ và bộ nhớ toàn cục.
- Trong thực tế, để tất cả các luồng của một warp, truy cập vào bộ nhớ chia sẻ sẽ nhanh như truy cập vào một thanh ghi miễn là không có bất kỳ sự xung đột dải nhớ (bank) giữa các luồng, như chi tiết dưới đây.
- Để có được băng thông bộ nhớ cao, bộ nhớ chia sẻ được chia thành các mô-đun bộ nhớ có kích thước bằng nhau gọi là các dải nhớ, mà có thể được truy cập cùng một lúc.
- Vì vậy, bất kỳ bộ nhớ đọc hoặc ghi yêu cầu thực hiện của các n địa chỉ nằm trong n dải nhớ riêng biệt thì có thể được phục vụ đồng thời, hiệu suất của một băng thông hiệu quả cao hơn n lần băng thông của một mô-đun đơn lẻ.
- 52 Tuy nhiên, nếu hai địa chỉ của một yêu cầu bộ nhớ rơi cùng vào một dải nhớ, đó là một xung đột dải nhớ và việc truy cập vào các dải nhớ phải được nối tiếp.
- Nếu số lượng yêu cầu bộ nhớ riêng biệt là n, yêu cầu vùng nhớ khởi tạo ban đầu sẽ gây ra xung đột dải nhớ theo n cách.
- Đối với các thiết bị khả năng tính toán 1.x, kích thước warp là 32 và số lượng của các dải nhớ nhớ là 16, một yêu cầu bộ nhớ chia sẻ cho một warp được chia thành một yêu cầu cho nửa đầu của warp và yêu cầu một cho nửa sau của warp.
- Thực hiện nhân ma trân với giải thuật song song bằng Cuda C cần thực hiện các bước [19] 56 • Khởi tạo ma trận A, B trên bộ nhớ Host Memory • Copy dữ liệu của ma trận A, B được sang Device Memory • GPU tạo ra số Block cần thiết để tính toán.
- Copy dữ liệu trên bộ nhớ thiết bị sang bộ nhớ chính Bài toán nhân ma trận được đánh giá trên hệ thống CPU Intel Core 2 Duo P8600@ 2.4GHz, 4GB Ram và GPU Geforce 9600M GS, 32 cores, cores speed 430MHz, Memory, 1024M

Xem thử không khả dụng, vui lòng xem tại trang nguồn
hoặc xem Tóm tắt