Luận văn Thạc sĩ: Nghiên cứu công nghệ xử lý GPU và ứng dụng

i

ĐẠI HỌC THÁI NGUYÊN

ĐẠI HỌC CÔNG NGHỆ THÔNG TIN VÀ TRUYỀN THÔNG

ĐINH TIẾN NGỌC

NGHIÊN CỨU CÔNG NGHỆ XỬ LÝ GPU VÀ ỨNG DỤNG

THÁI NGUYÊN 2017

ii

LỜI CAM ĐOAN

Tôi xin cam đoan :

Những nghiên cứu dưới đây trong luận văn của tôi hoàn toàn trung thực

không vi phạm bất kỳ quyền sở hữu trí tuệ nào. Nếu sai tôi xin chịu hoàn toàn

trách nhiệm.

TÁC GIẢ LUẬN VĂN

Đinh Tiến Ngọc

iii

LỜI CẢM ƠN

Lời đầu tiên tôi xin chân thành cảm ơn đếnTS. Lê Quang Minh người thầy đã

tận tình giúp đỡ, hướng dẫn tôi hoàn thành luận văn này.

Tôi cũng xin chân thành cảm ơn các thầy, cô giảng viên cao học người đã giúp

đỡ tôi nâng cao kiến thức giúp tôi có những kiến thức bổ trợ giúp hoàn thiện cho luận

văn này.

Tôi cũng xin chân thành cảm ơn người thân, bạn bè đã giúp đỡ và động viên tôi

trong suốt thời gian học tập cũng như trong thời gian thực hiện đề tài.

Xin chân thành cảm ơn!

Thái Nguyên, ngày tháng 5 năm 2017

TÁC GIẢ LUẬN VĂN

Đinh Tiến Ngọc

iv

DANH MỤC THUẬT NGỮ

TiếngAnh TiếngViệt

Bộ xử lý đồ họa GPU

Tính toán thông dụng trên GPU gpgpu

Application Program Interface : Định nghĩa một giao diện API

chuẩn để triệu gọi một tập các chức năng.

coproccessor bộ đồng xử lý

hạt nhân kernel

Kết cấu: cấu trúc của đối tượng, nó được xem như mô hình texture

thu nhỏ của đối tượng.

texturefetches Hàm đọc kết cấu

texturereference Tham chiếu kết cấu

Mỗi khối được tách thành các nhóm SIMD của các luồng. warp

Single Instruction Multiple Data: đơn lệnh đa dữ liệu SIMD

Dòng stream

streamingprocessor Bộ xử lý dòng

MIMD Multiple Instruction Multiple Data: đa lệnh đa dữ liệu

primarysurface Bề mặt chính

proccessor Bộ xử lý

Rasterization Sự quét mành trên màn hình

v

MỤC LỤC

LỜI CAM ĐOAN .............................................................................................................. i

LỜI CẢM ƠN ................................................................................................................ iii

DANH MỤC THUẬT NGỮ .......................................................................................... iv

MỤC LỤC ....................................................................................................................... v

DANH MỤC HÌNHVẼ ................................................................................................ vii

LỜI MỞ ĐẦU ............................................................................................................. viii

CHƯƠNG I : KHÁI QUÁT VỀ BỘ XỬ LÝ ĐỒ HỌA GPU VÀ XỬ LÝ SONG SONG

......................................................................................................................................... 1

1.1 Khái quát về xử lý song song................................................................................. 1

1.1.1 Khái quát về xử lý song song .......................................................................... 1

1.1.2 Khái quát về Hệ thống máy tính song song ..................................................... 3

1.1.3 Khái quát về lập trình song song ..................................................................... 7

1.1.4 Các nguyên tắc khi thiết kế giải thuật xử lý song song ................................... 9

1.2. Khái quát về công nghệ GPU và các ứng dụng .................................................. 10

1.2.1. Tổng quan về GPU ....................................................................................... 11

1.2.2. Nguồn gốc và quá trình phát triển GPU ....................................................... 11

1.2.3. Cấu trúc của bộ xử lý đồ họa GPU ............................................................... 15

1.2.4. Lập trình trên GPU ....................................................................................... 19

1.2.5. Các hỗ trợ phần mềm cho xử lý tính toán trên GPU .................................... 22

1.2.6. Các kỹ thuật tính toán trên GPU ................................................................... 26

1.2.7.Các giải thuật ứng dụng trên GPU ............................................................. 29

CHƯƠNG II: XỬ LÝ SONG SONG TRÊN THIẾT BỊ ĐỒ HỌA GPU VỚI CUDA . 31

2.1. Khái quát về CUDA ............................................................................................ 31

2.2.Cơ chế lập trình và cách thức hoạt động của CUDA ........................................... 33

2.2.1.Cơ chế lập trình .............................................................................................. 33

2.2.2.Cách thức hoạt động của CUDA ................................................................... 33

2.3. Tổng quan về lập trình với CUDA ...................................................................... 38

2.3.1. Là ngôn ngữ lập trình mở rộng của ngôn ngữ lập trình C ............................ 38

vi

2.3.2. Các phần mở rộng của CUDA ...................................................................... 38

2.3.3.Biến Built-in trong CUDA ............................................................................. 41

2.3.4. Biên dịch CUDA thông qua NVCC.............................................................. 42

2.3.5.Một số trường hợp cụ thể tính toán song song bằng CUDA ......................... 42

2.4. Các ứng dụng của CUDA trong các lĩnh vực ..................................................... 45

2.4.1. Ứng dụng của CUDA trong game ................................................................ 45

2.4.2. Ứng dụng của CUDA với video số ............................................................... 45

CHƯƠNG III: SỬ DỤNG GPU ĐỂ LÀM TĂNG TỐC ĐỘ TÍNH TOÁN CHO BÀI

TOÁN MÃ HÓA AES .................................................................................................. 48

3.1 Giới thiệu về AES ................................................................................................ 48

3.2 Thuật toán mã hóa ........................................................................................... 48

3.2.1 Công đoạn mã hóa ............................................................................................. 50

3.2.2 Công đoạn giải mã ............................................................................................ 54

3.3 Chương trình thuật toán song song mã hóa AES sử dụng GPU ....................... 62

3.3.1. Giao diện chương trình demo ....................................................................... 92

3.3.2. Kết quả chương trình và đánh giá hiệu suất tính toán ................................. 93

KẾT LUẬN ................................................................................................................... 88

TÀI LIỆU THAM KHẢO ............................................................................................. 89

vii

DANH MỤC HÌNHVẼ

Hình 1 : Kiến trúc Von Neumann ............................................................................ 1

Hình 2 : Máy tính song song có bộ nhớ chia sẻ ....................................................... 4

Hình 3 : Máy tính song song có bộ nhớ phân tán .................................................... 5

Hình 4 : Kiến trúc máy SISD ................................................................................... 5

Hình 5 : Kiến trúc máy SIMD ................................................................................. 6

Hình 6 : Kiến trúc máy MISD ................................................................................. 6

Hình 7 : Kiến trúc máy MIMD ................................................................................ 7

Hình 8 : Mô hình lập trình truyền thông hai tác vụ của hai máy tính .................... 8

Hình 9 : Mô hình lập trình song song dữ liệu .......................................................... 9

Hình10: Kiến trúc GPU của NVIDIA và AMD................................................. 19

Hình 11: Kiến trúc phần mềm CUDA ................................................................ 31

Hình 12: Thao tác cấp phát và thu hồi bộ nhớ ................................................... 32

Hình 13: Vùng nhớ dùng chung mang dữ liệu gần ALU hơn .......................... 33

Hình 14: Sơ đồ hoạt động truyền dữ liệu giữa Host và Device ............................. 34

Hình 15: Khối luồng .............................................................................................. 36

Hình 16: Mô hình bộ nhớ trên GPU ...................................................................... 37

Hình 17: Chiều của lưới và khối với chỉ số khối và luồng .................................... 42

Hình 18: Phương pháp đánh chỉ số luồng. ............................................................. 45

Hình 19 : Mã hóa và giải mã ................................................................................. 49

Hình 20: Biến đổi SubBytes() đối với mảng trạng thái ......................................... 51

Hình 21: Mô tả Hàm ShiftRows() ......................................................................... 51

Hình 22: Mô tả hàm MixColumns() ...................................................................... 52

Hình 23: Mô tả hàm AddRoundKey() ................................................................... 53

Hình 24: Mô tả hàm InvShiftRow()....................................................................... 55

viii

LỜI MỞ ĐẦU

Với sự phát triển như vũ bão của công nghệ, ngày nay công nghệ thông tin đã

trở thành một phần không thể thiếu trong cuộc sống. Không những thế nó còn là một

công cụ hữu hiệu trong các ngành khoa học, công nghệ cao,… đặc biệt là những

ngành có nhu cầu tính toán lớn. Tuy nhiên trong khi với nhu cầu tính toán ngày cành

tăng cao đó, ngành công nghệ thông tin lại vấp phải một vấn đề tối quan trọng đó là

năng lực xử lý của CPU có hạn. Các nhà phát triển phần cứng đã thực hiện gia tăng

mức độ xử lý cho CPU bằng cách gia tăng xung cho CPU. Tuy nhiên việc này cũng

chạm ngưỡng bởi gặp phải vấn đề về tản nhiệt cho CPU do nhiệt độ CPU quá cao.

Một hướng mới đã được các nhà nghiên cứu đưa ra đó là phát triển bộ xử lý đa

nhân với cơ chế xử lý song song.

Một bước phát triển trong hướng mới đó chính là bộ xử lý đồ họa – GPU

(Graphics Processing Unit - bộ xử lý đồ họa). Khi mới ra đời, GPU chỉ được sử dụng

với mục đích công việc phù hợp với khả năng là tăng tốc độ xử lý đồ họa, cũng như

trong ngành trò chơi là chủ yếu. Nhưng với sự phát triển dần của các trò chơi và các

phần mềm đồ họa, đã khiến GPU phát triển thêm và đến thế hệ GPUNV30 của

NVIDIA ra đời người ta đã bắt đầu phát triển những công việc khác cho GPU như

hỗ trợ tính toán dấu chấm động đơn, hỗ trợ tính toán lên cả ngàn lệnh. Và đặc biệt với

tiềm năng như vậy có thể nghĩ tới việc sử dụng GPU ngoài đồ họa. Cùng với ý tưởng

như vậy tôi đã liên tưởng đến việc áp dụng việc xử lý song song trên GPU thông qua

ngôn ngữ lập trình CUDA. Xuất phát từ ý tưởng trên tôi đã chọn đề tài: NGHIÊN

CỨU CÔNG NGHỆ XỬ LÝ GPU VÀ ỨNG DỤNG.

Luận văn gồm 3 chương chính:

Chương 1: Khái quát về bộ xử lý đồ họa GPU và xử lý song song, Chương

này giới thiệu tổng quan về xử lý song song và bộ xử lý đồ họa GPU

Chương 2: Xử lý song song trên thiết bị đồ họa GPU với CUDA. Chương này

nghiên cứu về ngôn ngữ lập trình CUDA và cách xử lý song song bằng CUDA trên

GPU.

Chương 3: Sử dụng GPU để làm tăng tốc độ tính toán cho bài toán mã hóa

ix

AES. Chương này tiến hành cài đặt thử chương trình song song, xử lý song song mã

hóa AES trên GPU bằng ngôn ngữ CUDA và đưa ra kết quả cùng kết luận về hiệu

suất của GPU.

1

CHƯƠNG I : KHÁI QUÁT VỀ BỘ XỬ LÝ ĐỒ HỌA GPU VÀ XỬ LÝ SONG SONG

1.1 Khái quát về xử lý song song

1.1.1 Khái quát về xử lý song song

Nguồn gốc ra đời của xử lý songsong

Một trong những nền tảng máy tính cơ bản đó là thiết kế máy tính của John Von

Neumann. Đó là thiết kế mà ở đó một lệnh được thực hiện trên một bộ xử lý.

Hình1: Kiến trúc Von Neumann

Khi cần tính toán với lượng câu lệnh và phép tính lớn thì thiết kế trên trở nên

lỗi thời. Người ta đã đưa ra các phương pháp nhằm giải quyết vấn đề trên. Trong đó

có việc tăng số lượng nhân xử lý hoặc kết nối nhiều máy tính thông qua mạng để

tăng tốc độ xử lý.

Khi tăng tốc xử lý các phép tính trên máy tính song song, việc sử dụng các thuật

toán tuần tự đã không còn thích hợp và không tận dụng hết khả năng tiềm tàng của

máy tính song song. Dẫn đến việc ra đời các giải thuật song song.

Lý do phải xử lý song song

Như đã nói ở trên máy tính song song với bộ xử lý nhiều nhân đã thay thế dần

máy tính đơn nhân, một bộ xử lý. Và với những thuật toán, câu lệnh, phép xử lý tuần

tự đã không còn phù hợp với máy tính song song. Do vậy xử lý song song đã ra đời

thay thế cho xử lý tuần tự nhằm đem lại hiệu năng tính toán cao hơn.

Bằng chứng đã thấy trong thực tế với nhiều bài toán xử lý với lượng dữ liệu lớn

2

yêu cầu tốc độ nhanh và độ chính xác như các bài toán về đồ họa, xử lý ảnh, xử lý

tín hiệu, mô phỏng giao thông, mô phỏng sự chuyển động của các phân tử, nguyên

tử, dự báo thời tiết, mô phỏng bản đồ gen……Xử lý song song đã chứng minh được

khả năng xử lý cũng như khả năng phát triển của nó sau này.

Các khái niệm trong xử lý songsong

 Định nghĩa xử lý songsong

Xử lý song song là quá trình xử lý thực hiện nhiều tiến trình cùng một lúc để

xử lý một bài toán, trên nhiều bộ xử lý.

 Làm rõ giữa xử lý song song và xử lý tuần tự

Xử lý tuần tự là tại mỗi thời điểm chỉ xử lý một phép toán còn xử lý song song

tại một thời điểm có thể thực hiện nhiều phép toán cùng một lúc trên nhiều bộ xử lý,

làm cho khả năng xử lý tăng lên đáng kể so với xử lý tuần tự. Bảng dưới đây cho

thấy sự khác nhau giữa xử lý tuần tự và xử lý song song.

Bảng 1.1: So sánh sự khác nhau giữa lập trình tuần tự và song song

Lập trình tính toán tuần tự Lập trình tính toán song song

- Chương trình chạy trên một bộ xử lý - Chương trình ứng dụng chạy trên nhiều

(single processor). bộ xử lý.

- Các câu lệnh được bộ xử lý(CPU) - Các câu lệnh được các bộ vi xử lý thực

thực hiện một cách lần lượt. hiện một cách song song, đồng thời.

- Mỗi câu lệnh thực thiện trên duy nhất - Mỗi câu lệnh có thể thao tác trên nhiều

một thành phần dữ liệu. thành phần dữ liệu khác nhau.

3

- Chỉ cần viết đúng câu lệnh và giải - Viết đúng câu lệnh và giải thuật. Ngoài ra

thuật là có thể chạy chương trình phải chỉ rõ phần nào của chương trình cần

phải chạy song song.

- Dùng với các bài toán dữ liệu nhỏ, độ - Được sử dụng đối với các bài toán có dữ

phức tạp không cao và yêu cầu thời gian liệu lớn, độ phức tạp cao và thời gian ngắn.

chấp nhận được.

 Mục đích của xử lý song song

Dựa trên việc tính toán song song trên nhiều bộ xử lý cho tốc độ xử lý cao,

xử lý song song thường dùng cho các bài toán có độ phức tạp lớn, yêu cầu khối lượng

tính toán lớn.

1.1.2 Khái quát về Hệ thống máy tính song song

Là một hệ thống máy tính với nhiều bộ nhân xử lý có khả năng xử lý song

song. Bao gồm cả máy tính với bộ xử lý mà trong đó gồm nhiều lõi, cũng được gọi

là máy tính song song.

Phân loại máy tính song song phổ biến nhất được biết tới là phân loại máy

tính song song của Michael Flynn vào năm 1966. Dựa vào các đặc điểm về số lượng

bộ xử lý, cấu trúc bộ nhớ,… Michael Flynn đã phân máy tính thành bốn loại dựa trên

sự biểu hiện của cặp khái niệm: Dòng lệnh (instruction stream) và dòng dữ liệu (data

stream), mỗi loại nằm trong một trong hai trạng thái đơn (single) hoặc đa (multiple).

Dựa theo phân loại của Flynn có 4 loại máy tính song song theo bảng dưới:

4

Bảng 1.2: Mô tả phân loại kiến trúc của Flynn

Dòng lệnh (instruction Dòng dữ liệu Loại kiến trúc

stream) (data stream)

Trạng thái đơn (single) Trạng thái đơn SISD

(single) Single Instruction Single

Trạng thái đơn (single) Trạng thái đa SIMD Data

(multiple) Single Instruction Multiple

Trạng thái đa (multiple) Trạng thái đơn MISD Data

(single) Multiple Instruction Single

Trạng thái đa (multiple) Trạng thái đa MIMD Data

(multiple) Multiple Instruction Multiple

Data

Dựa trên cấu trúc bộ nhớ Flynn đã chia ra làm các loại máy tính trên. Các bộ

xử lý với bộ nhớ chia sẻ có thể truy cập đến vùng nhớ chung. Sự thay đổi về nội dung

bộ nhớ sẽ được nhận biết bởi các bộ xử lý khác.

Hình 2 : Máy tính song song có bộ nhớ chia sẻ

Lại có thể chia ra tiếp 2 lớp nhỏ hơn trong loại máy tính này: Lớp máy tính

UMA (Uniform Memory Access – Truy cập bộ nhớ đồng nhất) khả năng truy cập bộ

nhớ đối với mỗi bộ xử lý có thời gian là như nhau. Lớp máy tính NUMA (Non-

Uniform Memory Access – Truy cập bộ nhớ không đồng nhất) khả năng truy cập bộ

nhớ đối với mỗi bộ xử lý có thời gian là không như nhau.

5

Máy tính song song với bộ nhớ phân tán. Mỗi bộ xử lý có một bộ nhớ và khả

năng truy cập độc lập. Sự thay đổi nội dung vùng nhớ của một bộ xử lý không làm

ảnh hưởng đến các bộ xử lý khác

Hình 3 : Máy tính song song có bộ nhớ phân tán

Mô hình cấu trúc đơn dòng lệnh đơn luồng dữ liệu (SISD)

Là máy tính chỉ có một bộ xử lý tại mỗi thời điểm chỉ thực hiện một lệnh, đọc,

ghi một mục dữ liệu. Chỉ có một thanh ghi gọi là bộ đệm và kết quả đầu ra theo tuần

tự các câu lệnh.

Hình 4 : Kiến trúc máy SISD

Mô hình cấu trúc đơn dòng lệnh đa luồng dữ liệu (SIMD)

Mỗi bộ xử lý thực hiện xử lý một luồng dữ liêu. Các bộ xử lý cùng thực hiện

một phép toán trên nhiều luồng dữ liệu khác nhau và có một thành phần để điều khiển

cho các bộ xử lý thực hiện xử lý các luồng câu lệnh.

6

Hình 5 : Kiến trúc máy SIMD

Mô hình cấu trúc đa dòng lệnh đơn luồng dữ liệu (MISD)

Là loại máy tính có thể thực hiện nhiều câu lệnh trên cùng một mục dữ liệu.

Hình 6 : Kiến trúc máy MISD

Mô hình cấu trúc đa dòng lệnh đa luồng dữ liệu (MIMD)

Là loại máy tính đa nhân, đa bộ xử lý có thể thực hiện nhiều câu lệnh trên

nhiều luồng khác nhau.Các bộ xử lý đều có bộ nhớ riêng biệt nhưng cũng có thể truy

cập vào bộ nhớ chung khi cần giúp tăng tốc độ xử lý. Mô hình kiến trúc này là mô

hình kiến trúc phức tạp nhất nhưng cũng là mô hình ưu việt nhất và cũng đã có nhiều

máy tính được xây dựng trên kiến trúc này, ví dụ: BBN Butterfly.

7

Hình 7 : Kiến trúc máy MIMD

1.1.3 Khái quát về lập trình song song

Là việc lập trình các câu lệnh các đoạn chương trình song song để chạy trên

hệ thống máy tính song song. Cũng có thể hiểu là việc song song hóa các thuật toán

tuần tự nhằm tăng tốc độ xử lý tính toán lên nhiều lần.

Trong đó việc lập trình song song là chia các chương trình, bài toán lớn thành

các bài toán con rồi chia các bài toán con thành các bài toán con nhỏ hơn…rồi chia

các bài toán con nhỏ đó cho các bộ xử lý giải quyết và đồng bộ về mặt thời gian xử

lý công việc để nhận được kết quả cuối cùng . Việc quan trong nhất ở đây là xử lý

các bài toán đồng thời. Do vậy khi muốn lập trình song song một thuật toán ta cần

xác định là có thể song song hóa thuật toán đó hay không.

Các kiểu song song hóa:

 Song song hóa mặc định: Việc phân chia công việc đến các bộ xử lý là tự

động.

 Song song hóa thủ công: Việc phân chia công việc đến các bộ xử lý do người

lập trình tự quyết định.

Trong đó cần để ý đến vấn đề cân bằng tải. Khi một bộ xử đang giải quyết quá

nhiều công việc cần chuyển các công việc đến các bộ xử lý khác đang giải quyết các

công việc ít hơn.

8

Một vấn đề quan trọng trong lập trình song song đó là việc kết nối giữa các

bộ nhớ. Có hai kỹ thuật kết nối cơ bản là: sử dụng bộ nhớ chia sẻ và truyền thông

điệp.

Mô hình lập trình song song bao gồm các hệ thống truyền thông và vào/ra

song song, các ứng dụng, ngôn ngữ, bộ biên dịch, thư viện. Việc phân chia công việc

cho máy tính song song hiệu quả cho tất cả các bài toán là không khả thi. Vấn đề là

người lập trình viên song song phải biết sử dụng hiệu quả các mô hình cho từng bài

toán hoăc kết hợp chúng với nhau.

Hiện nay có các mô hình lập trình song song: Truyền thông điệp (Message

Passing) và Song song dữ liệu (Data Parallel).

Truyền thông điệp trong xử lý song song

Là mô hình được sử dụng rộng rãi cho các hệ phân tán. Bao gồm các đặc trưng

sau:

 Trong quá trình tính toán mỗi luồng sử dụng một vùng nhớ cục bộ riêng.

 Các luồng có thể sử dụng chung tài nguyên.

 Việc trao đổi giữa các luồng được thực hiện bằng cách gửi các thông điệp.

 Mỗi luồng sẽ thực hiện việc điều khiển việc truyền dữ liệu. Ví dụ mỗi thao tác gửi

ở một luồng thì phải ứng với một thao tác nhận ở luồng khác.

Hình 8: Mô hình lập trình truyền thông hai tác vụ của hai máy tính

Song song dữ liệu trong xử lý song song

Ở mô hình này chủ yếu việc song song được thực hiện trên một tập dữ liệu.

Cấu trúc của tập dữ liệu này là mảng hoặc khối. Cùng một phép toán, Các tác vụ sẽ

thực hiện trên cùng một kiểu dữ liệu nhưng trên các tập dữ liệu khác nhau. Mô hình

9

này chủ yếu dành cho máy tính song song kiểu bộ xử lý mảng.

Hình 9 : Mô hình lập trình song song dữ liệu

1.1.4 Các nguyên tắc khi thiết kế giải thuật xử lý song song

Để xét đến việc xử lý song song cần quan tâm đến giải thuật và việc thực hiện

giải thuật song song đó trên loại máy tính nào.

Những bước cần thực hiện khi thiết kế giải thuật song song:

 Phân nhỏ bài toán thành bài toán con nhỏ hơn sao cho độc lập về mặt dữ liệu

và chức năng, và giải quyết các bài toán này đồng thời.

 Chỉ rõ việc chia sẻ dữ liệu và truy cập.

 Chia các tác vụ cho các tiến trình cho các nhân xử lý.

 Việc đồng bộ các tiến trình.

Các bước thiết kế giải thuật song song và chương trình có thể xử lý

song song

Khi bắt tay thiết kế một giải thuật song song điều quan trọng là ta phải biết

được cần song song trong công đoạn nào. Trước khi song song thuật toán ta cần biết

thuật toán đó ở dạng tuần tự, đồng thời phải hiểu về giải thuật và ngôn ngữ lập trình

cụ thể. Khi đó ta xét các bước sau để xem bài toán đó có song song hóa được không:

- Phân tích bài toán, xác định các thành phần của giải thuật có thể hoặc không

thể song song hóa được.

- Để ý đến các hạn chế của xử lý song song trong đó hạn chế lớn nhất là sự

10

phụ thuộc dữ liệu.

Các ví dụ về chương trình có thể song song hóa

Một bài toán ví dụ có thể song song hóa đó là cộng hai mảng số nguyên có cùng

phần tử. Công đoạn có thể song song hóa là cộng các phẩn tử có cùng thứ tự sẽ độc

lập với các phần tử khác.

A

1

1

1

1

1

+

B

4

4

4

4

4

=

C

5

5

5

5

5

Các ví dụ về chương trình không thể song song hóa

Tính chuỗi Fibonacci (1, 1, 2, 3, 5, 8, 13, 21,…) bằng cách sử dụng công

thức: F(k+2)= F(k+1) + F(k), với n > 1.

Bài toán này không song song hóa được vì nguyên nhân tính số hạng của dãy

Fibonacci theo công thức là phụ thuộc chứ không phải là độc lập. Trong đó việc tính

giá trị thứ k+2 phải sử dụng giá trị của cả hai giá trị k+1 và k.

Cách thiết kế giải thuật song song

Cách thiết kế khái quát được đề cập đến trong luận văn này là phân chia dữ liệu.

 Áp dụng khi liên quan đến tính toán trên nhiều cấu trúc dữ liệu. Các cấu trúc

dữ liệu này có thể phân chia nhỏ hơn và có thể tính toán trên phần cấu trúc dữ liệu

nhỏ hơn này.

Minh chứng cho lập luận trên là bài toán cộng hai mảng số nguyên. Trường hợp

ta có n bộ xử lý cùng làm việc để cộng hai mảng X[0…N-1] và Y[0…N-1] lưu vào

mảng kết quả Z[0...N-1], việc phân chia dữ liệu sẽ đặt N/n phần tử của mỗi mảng

vào từng quá trình và nó sẽ tính toán N/n phần tử tương ứng của mảng kết quả. Như

11

vậy, với n bộ xử lý càng nhiều thì thời gian chạy càng nhanh, ngược lại thì chạy càng

chậm.

1.2. Khái quát về công nghệ GPU và các ứng dụng

1.2.1. Tổng quan về GPU

Một phần không thể tách rời của hệ thống máy tính ngày nay đó là bộ xử lý

đồ họa (Graphic Proccessing Unit) gọi tắt là GPU.Từ khi ra đời cho đến nay GPU

đã cho ta thấy sự phát triển ấn tượng về hiệu suất. GPU hiện nay không chỉ là một

công cụ xử lý đồ họa mạnh mà còn là một bộ xử lý hỗ trợ lập trình song song ở mức

cao, giúp giải các bài toán số học cần khả năng xử lý số học phức tạp và băng thông

bộ nhớ tăng hơn đáng kể so với CPU cùng loại. Sự phát triển mạnh về hiệu suất của

GPU trong cả việc hỗ trợ lập trình và khả năng tính toán của nó đã tạo ra một hướng

nghiên cứu mới. Một nhóm các chuyên gia đã nghiên cứu thành công một lượng

lớn các vấn đề phức tạp đòi hỏi tính toán lớn vào GPU. Việc này đã góp phần ứng

dụng GPU vào giải quyết các bài toán hiệu năng cao của tính toán hiện đại. Tính

toán và lập trình trên GPU là một thay thế tiềm năng cho CPU trong hệ thống máy

tính hiện đại. Trong một tương lai không xa, GPU sẽ đảm nhận thay cho CPU những

công việc như xử lý hình ảnh, đồ họa, các tính toán phức tạp thay vì chỉ dừng lại ở

những ứng dụng trò chơi 3D.

1.2.2. Nguồn gốc và quá trình phát triển GPU

GPU là bộ xử lý của card đồ họa, dùng để tính toán các phép toán dấu phảy

động.

Lúc đầu GPU là bộ xử lý gắn trên card đồ họa phục vụ cho việc tính toán các

phép toán dấu phảy động.

Với các vi mạch siêu nhỏ, cùng các phép toán đăc biệt , GPU được sử dụng chủ

yếu trong các hoạt động cần đến xử lý đồ họa cao như trong các game đồ họa cao

hoặc các xử lý đồ họa 3D.

Việc xử lý một số phép toán đồ họa nguyên thủy khiến GPU chạy nhanh hơn

nhiều so với việc vẽ trực tiếp trên màn hình với CPU.

12

 Thập kỷ 70:

Công bố bộ điều khiển phần cứng kết hợp đồ họa và chế độ text, tính toán

vị trí và hiển thị (theo khuôn dạng phần cứng hỗ trợ) và những hiệu ứng khác trên

các máy tính ATARI 8-bit, bởi hãng sản xuất chip ANTIC và CTIA. Chíp ANTIC

là một bộ xử lý dành riêng cho ánh xạ (lập trình được) giữa dữ liệu đồ họa tới đầu

ra video và text. Nhà thiết kế chip ANTIC, sau đó đã thiết kế chip đồ họa cho

Commodore Amiga.

 Thập kỷ 80:

Chiếc máy tính đầu tiên có chứa các bộ blit (Block Image Transfer là sự

chuyển động của một bitmap lớn trong game 2D) trong phần cứng video, hệ thống

đồ họa 8514 của IBM là Commodore Amiga, là một trong những card video đầu

tiên trên máy tính có thể thực thi các phép toán 2D nguyên thủy .

Những tính năng của Amiga bây giờ được công nhận là bộ gia tốc đồ họa

đầy đủ, giảm tải tất cả các chức năng thế hệ video cho phần cứng, bao gồm tô màu

vùng,chuyển khối hình ảnh ,vẽ đường thẳng, và bộ đồng xử lý đồ họa với cùng với

tập các chỉ thị lệnh nguyên thủy của riêng nó.

 Thập kỷ 90:

S3 Graphics giới thiệu bộ gia tốc chip 2D đầu tiên, các 86C911 S3 vào năm

1991.Các 86C911 sinh ra một làn sóng công nghệ. Đến năm 1995, tất cả các nhà

sản xuất chip đồ họa máy tính lớn đã thêm các hỗ trợ tăng tốc 2D cho chip của họ.

Vào thời điểm này, bộ tăng tốc Windows với các chức năng cố định khá đắt tiền và

cao hơn bộ đồng xử lý đồ họa mục đích chung trong hiệu năng của Windows.

2D GUI tiếp tục tăng tốc phát triển trong suốt những năm 1990. Từ khả năng sản

xuất được phát triển đã tác động vào các mức độ tích hợp chip đồ họa. Cộng với việc

các giao diện lập trình ứng dụng (API) cho ra một lượng lớn tác vụ, ví dụ thư viện đồ

họa của Microsoft WinG cho Windows 3.x, và giao diện DirectDraw của họ cho tăng

tốc phần cứng của game 2D trong Windows 95 và sau đó.

13

Với sự hỗ trợ CPU-thời gian thực, đồ họa 3D đã d ầ n trở nên phổ biến trong

máy tính và giao diện các trò chơi, dẫn đến nhu cầu phát triển rộng rãi phần cứng

tăng tốc cho đồ họa 3D. M ộ t điển hình về loạt trên thị trường phần cứng đồ họa

3D có thể thấy trong các trò chơi video thế hệ console thứ năm như PlayStation và

Nintendo 64. Một trong các thử nghiệm không thành công là các chip đồ họa 3D giá

thành rẻ là ViRGE S3, ATI Rage, và Matrox Mystique. Những chip này về cơ bản

là bộ gia tốc 2D thế hệ trước bổ sung thêm các tính năng 3D. Nhiều thành phần

trong đó được thiết kế tương thích với thế hệ chip đời trước để dễ thực hiện và chi

phí tối thiểu. hiệu năng đồ họa 3D đã chấp nhận được với bảng mạch rời dành

riêng cho các chức năng tăng tốc 3D (thiếu chức năng 2D GUI) như 3dfx Voodoo

vào thời điểm ban đầu. Công nghệ sản xuất một lần nữa tiến triển, Chipset Verite

của Rendition được là sản phẩm đầu tiên được tích hợp video, bộ tăng tốc 2D GUI,

và chức năng 3D tất cả vào một con chip.

Trở thành một lực lượng chi phối trên máy tính, và là một động lực cho phát

triển phần cứng OpenGL xuất hiện vào đầu những năm 90 như là API đồ họa chuyên

nghiệp. Phần mềm của OpenGL được phổ biến trong thời gian này, dù sau đó do

tác động của OpenGL đã dẫn đến hỗ trợ phần cứng rộng rãi. Một sự lựa chọn nổi lên

giữa các tính năng có trong phần cứng và những tính năng đó có trong OpenGL.

Trong thời gian cuối những năm 90 DirectX phổ biến với các nhà phát triển

game Windows. Microsoft khẳng định nghiêm ngặt về việc cung cấp sự hỗ trợ một-

một của phần cứng Không giống như OpenGL. Việc đó đã làm DirectX ít phổ biến

như là API đồ họa độc lập ngay từ đầu trong khi đó các GPU có nhiều tính năng

đặc biệt, và hiện đã được ứng dụng OpenGL có thể được thừa hưởng, để lại DirectX

một thế hệ sau. Theo thời gian, Microsoft đã bắt đầu làm việc chặt chẽ hơn với các

nhà phát triển phần cứng, và bắt đầu nhắm mục tiêu các bản phát hành của DirectX

với những phần cứng đồ họa hỗ trợ. phiên bản API đầu tiên đang phát triển để đạt

được áp dụng rộng rãi trên thị trường chơi game là Direct3D 5,0 , và nó cạnh tranh

trực tiếp với nhiều phần cứng, như các thư viện đồ họa độc quyền, trong khi

OpenGL duy trì điều đó. Direct3D 7,0 hỗ trợ phần cứng tăng tốc biến đổi và ánh

14

sáng (T & L). Bộ tăng tốc 3D biến đổi từ lúc dầu chỉ là bộ quét đường thẳng đơn

giản về sau có thêm phần cứng quan trọng dùng cho các đường ống dẫn biến

đổi 3D. Sản phẩm đầu tiên trên thị trường với khả năng này là NVIDIA Geforce

256 (còn được gọi là NV10). Phần cứng biến đổi và ánh sáng, cả hai đều đã có

trong OpenGL, trong phần cứng những năm 90 và đặt tiền đề cho các phát triển

sau đó là các đơn vị đổ bóng điểm ảnh và đổ bóng vector mà với đặc tính linh hoạt

hơn và lập trình được.

 Từ những năm 2000 đến hiện nay:

GPU đã có thêm tính năng đổ bóng lập trình được cùng với sự ra đời của

API OpenGL và các tính năng tương tự trong DirectX. Mỗi điểm ảnh, mỗi vector

hình học bây giờ có thể được xử lý bởi một chương trình ngắn và trước khi nó

được chiếu lên màn hình. Lần đầu tiên hãng NVIDIA đã được sản xuất một

con chip có khả năng lập trình đổ bóng, GeForce 3 (tên mã NV20). Sự ra đời

của ATI Radeon 9.700 (còn gọi là R300), bộ tăng tốc Direct3D 9.0 lần đầu tiên

trên thế giới vào tháng 10 năm 2002, bộ đổ bóng điểm ảnh và vector có thể thực

hiện vòng lặp và các phép toán dấu phảy động dài, đã nhanh chóng trở nên linh

h o ạ t như CPU, và đòi hỏi cần có sự phát triển nhanh hơn cho các phép toán

mảng liên quan đến hình ảnh (image-array operations). Những thứ như lập bản đồ

bump, thêm vào các kết cấu (texture) thì đổ bóng điểm ảnh thường được sử dụng,

để làm cho một đối tượng trông bóng, căng mịn hoặc lồi lõm, hoặc thậm chí ảm

đạm, thô ráp. Khi khả năng xử lý của GPU tăng lên dẫn đến nhu cầu nguồn điện

cao hơn.Với GPU hiệu suất cao, khả năng tiêu thụ năng lượng nhiều hơn các CPU

hiện nay. Ngày nay, GPU song song đã bắt đầu thực hiện xâm nhập máy tính và cạnh

tranh với CPU, sự phổ biến của GPU đã len lỏi trong các lĩnh vực cần đến sự tính toán

cao như đại số tuyến tính, xử lý hình ảnh khoa học, tái tạo 3D và hỗ trợ lựa chọn

giá cổ phiếu , thăm dò dầu . Việc này đã làm tăng thêm động lực cho các nhà

sản xuất GPU từ "người dùng GPGPU" để giúp tiến thiết kế phần cứng tốt hơn.

15

1.2.3. Cấu trúc của bộ xử lý đồ họa GPU

Với năng lực của mình GPU là một bộ xử lý luôn dồi dào khả năng tính toán. Nhưng

điều hiển nhiên đó là nó phải trưng bày khả năng tính toán đó cho các lập trình

viên. Trong thời gian gần đây GPU đã phát triển từ một hàm cố định, bộ xử lý chuyên

dụng tới bộ xử lý lập trình song song, đầy đủ tính năng độc lập với việc bổ sung

thêm các chức năng cố định, và các chức năng chuyên biệt. Giờ đây các khía cạnh

về khả năng lập trình của bộ xử lý chiếm đã vị trí trung tâm. Chúng ta đi xem xét

từng khía cạnh.

Đường ống dẫn đồ họa (GraphicsPipeline)

Đầu vào của GPU là các kiểu hình học cơ bản, điển hình là tam giác, trong thế

giới không gian 3 chiều. Sau một quá trình những khối hình nguyên thủy đó được

làm bóng mờ (shade) và được vẽ lên màn hình, nơi chúng được lắp ráp để tạo ra

hình ảnh cuối cùng.

Các phép toán vector:

Hình học nguyên thủy (primary geometric) cấu thành từ các vector riêng lẻ.

Bằng cách tính toán tương tác với các luồng ánh sáng trong một bối cảnh cụ thể

mỗi vector phải được chuyển thành không gian trên màn hình và có bóng mờ. Những

bối cảnh tiêu biểu có thể có hàng chục đến hàng trăm ngàn vector, và mỗi vector

có thể được tính toán độc lập. Việc tính toán này rất phù hợp cho phần cứng song

song.

Các Thành phần nguyên tố:

Phần tử hỗ trợ phần cứng cơ bản trong GPU ngày nay đó chính là các vector

được lắp ráp vào các hình tam giác.

Sự quét mành:

Việc xác định những vị trí điểm ảnh nào trong không gian màn hình được bao

chứa bởi mỗi tam giác được gọi là Quét mành (rasterization). Mỗi một tam giác

16

tạo ra một phần tử cơ bản được gọi là "mảnh" ở vị trí điểm ảnh trong không gian

màn hình mà nó bao chứa. Một chú ý là do nhiều tam giác có thể chồng lên nhau tại

một vị trí điểm ảnh bất kỳ nên giá trị màu của mỗi điểm ảnh có thể được tính từ nhiều

mảnh.

Thao tác trên mảnh:

Lấy dữ liệu bổ sung từ bộ nhớ toàn cục trong các hình dạng của sự kết hợp

(sự kết hợp là hình ảnh được ánh xạ lên bề mặt), sử dụng thông tin màu sắc từ

vector, mỗi mảnh được làm bóng mờ để xác định màu sắc cuối cùng của nó. Cũng

giống trường hợp vector, mỗi mảnh được tính toán song song. Giai đoạn này đòi hỏi

nhiều tính toán nhất trong đường ống dẫn đồ họa.

Thành phần:

Với mỗi điểm ảnh các mảnh được lắp ráp thành hình ảnh cuối cùng với một

màu, và bằng cách giữ lại mảnh gần ống dẫn đồ họa nhất cho mỗi vị trí điểm ảnh.

Đã có các phép toán tại khung cảnh vector và mảnh đã được cấu hình nhưng

không thể lập trình được. Tính toán chủ yếu ở khung cảnh vector là tính toán màu

sắc ở mỗi vector và các độ sáng trong bối cảnh đó. Trong đường ống dẫn đồ họa,

các lập trình viên có thể kiểm soát được vị trí và màu sắc của các vector và ánh

sáng, nhưng không phải là mô hình chiếu sáng mà xác định tương tác giữa chúng.

Sự phát triển của cấu trúc GPU

Điều kiện tiên quyết cho các hiệu ứng phức tạp là biểu diễn hiệu quả các trường

hợp làm bóng mờ phức tạp hơn và các phép toán ánh sáng. Giai đoạn quan trọng

trên đã được thay thế bằng các phép toán trên mỗi mảnh và các hàm cố định chức

năng trên mỗi vector với chương trình chỉ định người sử dụng chạy trên từng vector

và từng mảnh. Qua thời gian các chương trình vector và chương trình mảnh đã có

ngày càng nhiều khả năng, với bộ chỉ thị (tập lệnh) đầy đủ tính năng, với giới hạn

lớn hơn về kích cỡ và tiêu thụ tài nguyên, và với các phép toán điều khiển luồng linh

hoạt hơn.

17

GPU hiện tại hỗ trợ mô hình bóng mờ thống nhất 4.0 (unified Shader Model

4.0) trên cả bóng mờ vector và mảnh:

 Hỗ trợ phần cứng đổ bóng mờ đến 65K chỉ thị tĩnh đồng thời chỉ thị động

không giới hạn.

 Lần đầu tiên các tập lệnh hỗ trợ cả số nguyên 32 bit và số dấu phảy động 32bit.

 Cho phép thao tác đọc trực tiếp và gián tiếp từ bộ nhớ toàn cục một cách tùy

ý .

 Hỗ trợ điều khiển luồng động trong các dạng vòng lặp và rẽ nhánh.

Tất cả các loại ứng dụng GPU đã tăng độ phức tạp chương trình vector và mảnh

Khi mô hình đổ bóng ra đời và ngày càng phát triển với hiệu suất cao hơn, kiến

trúc GPU ngày càng tập trung vào các bộ phận lập trình được của đường ống dẫn đồ

họa. Ngày nay GPU được khắc họa tốt hơn, như là công cụ lập trình được bao

quanh bởi các đơn vị hỗ trợ có chức năng cố định.

Cấu trúc của GPU hiện đại

GPU được xây dựng cho các nhu cầu ứng dụng khác nhau so với CPU, đó là

các yêu cầu tính toán lớn chạy song song, với trọng tâm là thông lượng hơn là độ

trễ. Do đó, các kiến trúc của GPU phát triển theo một hướng khác so với CPU.

Trong một đường ống dẫn của các tác vụ, hầu hết các giao diện lập trình đồ

họa phải xử lý một lượng lớn các dữ liệu đầu vào. Các nhiệm vụ nối tiếp nhau được

đưa vào đường ống dẫn đồ họa. Một cơ chế song song được đặt ra trong đường ống

dẫn đồ họa, cụ thể dữ liệu trong nhiều khung cảnh trong đường ống có thể được

tính toán cùng một thời điểm, đó là cơ chế song song dữ liệu.Lần lượt từng khung

cảnh, CPU có thể lấy một phần tử đơn và xử lý khung cảnh (stage) đầu tiên trong

đường ống. Tận dụng tất cả năng lực xử lý từng khung cảnh và chia đường

ống dẫn đồ họa theo thời gian.

GPU chia năng lực xử lý của bộ xử lý xét về mặt tài nguyên, theo khung cảnh.

Sao cho các đường ống được chia theo không gian.

18

Đây là một cơ chế xử lý tốt. Thứ nhất trong bất kỳ khung cảnh nào phần cứng

có thể khai thác cơ chế song song dữ liệu nên cơ chế song song nên công việc được

chạy bất kỳ lúc nào và được xử lý nhiều phần tử cùng một lúc.Với cơ chế như vậy

GPU có thể đáp ứng nhu cầu tính toán rất lớn của các đường ống dẫn đồ họa.

Thứ hai với mỗi công việc riệng biệt cần đến khả năng tính toán lớn phần cứng của

mỗi khung cảnh có thể được thay đổi với phần cứng chuyên dụng. Như trong giai

đoạn rasterization, khi cần tính thông tin bao phủ điểm ảnh của từng điểm ảnh

tam giác đầu vào, hiệu quả hơn khi thực hiện trên ứng dụng. Các chức năng đơn

giản được thay thế bằng thành phần lập trình được, nhưng với nhiệm vụ song song

không thay đổi.

Và cuối cùng ta có một đường ống dài feed-forward nhiều khung cảnh, mỗi khung

cảnh dùng cho mục đích khác nhau, và thích hợp cho việc song song hóa và phần cứng

song song. Thường với CPU mỗi phép toán thường mất 20 chu kỳ hoạt động. Với

GPU mỗi phép toán thường mất hàng ngàn chu kỳ hoạt động và độ trễ tương đối lâu.

Nhưng với cơ chế song song tác vụ và dữ liệu giữa các khung cảnh tạo ra khả năng

tính toán cao. Có một vấn đề đáng quan tâm là cân bằng tải. Hiệu suất của đường

ống GPU phụ thuộc vào khung cảnh chậm nhất của nó. Khi các chương trình mảnh

đơn giản, chương trình vector phức tạp, lúc đó hiệu suất chung phụ thuộc vào hiệu

suất của các chương trình vector. Ở thời điểm đầu sau khi lập trình các khung cảnh,

tập chỉ thị của các các chương trình mảnh và chương trình vector khá khác nhau, nên

chúng được tách riêng. Kiến trúc GPU xem xét lại đường ống song song tác vụ

nghiêm ngặt trong lợi thế của kiến trúc đổ bóng hợp nhất (unified shader), trong

đó tất cả đơn vị lập trình được trong đường ống chia sẻ một đơn vị phần cứng lập

trình được duy nhất khi cả hai chương trình vector và chương trình mảnh trở nên

đầy đủ tính năng, và tập chỉ thị lệnh hội tụ như nhau. các đơn vị lập trình bây giờ

phân chia thời gian của nó giữa công việc vector, công việc mảnh, và công việc hình

học (với DirectX có bộ đổ bóng 10 loại hình học khác nhau) trong khi phần lớn

các đường ống vẫn còn là song song tác vụ. Các đơn vị này có thể khai thác cả hai

cơ chế song song dữ liệu và song song tác vụ. Kiến trúc của GPU chuyển từ kiến

19

trúc song song tác vụ trong một đường ống nghiêm ngặt sang kiến trúc được phát

triển xung quanh một đơn vị lập trình được theo cơ chế song song dữ liệu thống

nhất khi các bộ phận lập trình được của đường ống chịu trách nhiệm tính toán ngày

càng nhiều trong các đường ống dẫn đồ họa. Sản phẩm GPU Xenos GPU của nó trong

Xbox 360 (2005) được AMD giới thiệu như các kiến trúc đổ bóng hợp nhất đầu tiên.

Ngày nay, cả GPU của AMD và NVIDIA đều có tính năng đổ bóng hợp nhất (unified

shaders). Lợi ích cho người sử dụng GPU là với chi phí cho phần cứng cao hơn ta có

cân bằng tải tốt hơn. Ngoài ra với tất cả nguồn lực lập trình được trong một đơn vị

phần cứng duy nhất, lập trình viên GPGPU bây giờ có thể tiếp cận đơn vị lập trình

được theo cách trực tiếp, hơn hẳn trước cách tiếp cận trước đây.

Hình10: Kiến trúc GPU của NVIDIA và AMD

1.2.4. Lập trình trênGPU

Chúng ta đi tìm hiểu mô hình lập trình của GPU. Lập trìnhtrênGPU

Các đơn vị lập trình của GPU tuân theo mô hình lập trình SPMD (single

program, multiple data): đơn chương trình, đa dữ liệu. Sao cho hiệu quả, bằng cách

sử dụng nhiều chương trình giống nhau GPU xử lý rất nhiều yếu tố (vector hoặc

20

mảnh) song song. Trong lập trình mô hình cơ sở, các yếu tố không thể giao tiếp với

nhau mỗi phần tử được độc lập với các phần tử khác. . Tất cả các chương trình GPU

được tổ chức theo cách: mỗi thành phần được xử lý song song bởi một đơn chương

trình và song song nhiều thành phần. Mỗi thành phần có thể hoạt động trên dữ liệu

dấu phảy động với một tập các chỉ thị lệnh vừa đủ dùng cho mục đích thông dụng

(general purpose) hay số nguyên 32-bit. Các thành phần co thể đọc dữ liệu từ bộ nhớ

chia sẻ toàn cục và cũng có thể ghi vào bộ nhớ này. Cách xử lý này đặc biệt thích

hợp với các chương trình làm việc với đường thẳng,cụ thể nhiều thành phần có thể

được xử lý trong các bước nối tiếp. Ta nghiên cứu mô hình SPMD tổng quát hơn.

Trên GPU mô hình này được hỗ trợ như sau.

Phần lớn tài nguyên dành cho việc tính toán là một trong những lợi ích của

GPU. Điều khiến đòi hỏi đáng kể phần cứng điều khiển là việc cho phép các con

đường thực thi khác nhau cho từng phần tử. GPU ngày nay hỗ trợ luồng điều khiển

riêng cho từng luồng, nhưng đặt ra việc bắt lỗi cho những luồng tạp nham. Các yếu

tố được nhóm lại với nhau thành những khối và các khối được xử lý song song.

Phần cứng tính cả hai bên của nhánh cho tất cả các phần tử trong khối nếu các yếu

tố phân nhánh ra các hướng khác nhau trong một khối. Với thế hệ GPU gần đây

kích cỡ của khối được giảm, ngày này đó là thứ tự của 16 phần tử.

Khi viết chương trình thi câu lệnh rẽ nhánh được phép nhưng không dư thừa tài

nguyên để tận dụng. Và người lập trình cần tận dụng tốt tài nguyên tính toán sao cho

hợp lý nhất.

Các xử lý tính toán trên GPU(GPGPU)

GPGPU là việc ánh xạ các bài toán tính toán mục đích thông thường lên GPU

sử dụng phần cứng đồ họa theo cách giống như bất cứ ứng dụng đồ họa chuẩn nào.

Do sự tương đương này, nó vừa cũng khó khăn hơn và dễ dàng hơn trong việc giải

thích quá trình hoạt động. Một mặt có điểm khác nhau giữa đồ họa và sử dụng cho

mục đích thông thường. Một mặt, các hoạt động thực tế là như nhau và rất dễ làm

theo. Chúng ta bắt đầu lần lượt đi tìm hiểu.

 Sử dụng GPU cho mục đích đồ họa:

21

Chúng ta bắt đầu với một đường ống dẫn GPU mà chúng ta đã mô tả và tập

- Khi lập trình cần xác định dạng hình học sẽ bao phủ khu vực trên màn hình.

trung vào các khía cạnh lập trình được.

- chương trình mảnh sẽ tạo bóng mờ của mỗi mảnh.

- Bằng cách kết hợp của phép toán toán học các chương trình mảnh sẽ tính giá

Sau đó quá trình quét mành sẽ tạo ra mỗi mảnh trên vị trí bao phủ đó.

- Kết quả sẽ là kết cấu tiền đề để đi qua các đường ống dẫn đồ họa.

trị của các mảnh và bộ nhớ toàn cục đọc từ bộ nhớ kết cấu toàn cục.

 Sử dụng GPU để tính toán cho các chương trình mục đích thông dụng:

Vẫn với đường ống dẫn đã chọn thực hiện tính toán general-purpose liên quan

đến các bước giống nhau, nhưng ký hiệu khác nhau.

Một ví dụ đặc biệt tiêu biểu là một mô phỏng tính chất lỏng được tính toán trên

lưới: tại mỗi bước, chúng ta tính toán trạng thái tiếp theo của chất lỏng cho mỗi

điểm lưới từ tình trạng hiện tại trên lưới của nó và trạng thái các điểm lân cận của nó

- Lập trình viên đưa ra một hình nguyên thủy bao gồm một miền tính toán

trên lưới.

quen thuộc. Các chương trình quét mành tạo ra một mảnh (fragment) ở mỗi vị trí

- Mỗi mảnh được làm bóng mờ bởi chương trình gerenal - purpose SPMD.

- Chương trình mảnh tính giá trị của mảnh bằng cách kết hợp các phép toán

điểm ảnh trong hình đó.

toán học và các truy cập từ bộ nhớ toàn cục. Mỗi điểm lưới có thể truy cập trạng

thái của các lân cận của nó ở bước tính toán trước đó trong khi tính toán giá trị hiện

- Bộ nhớ đệm chứa kết quả sau đó có thể được sử dụng như là một khởi đầu

tại của nó.

cho các chu kỳ tiếp theo

 Sử dụng GPU để tính toán cho chương trình mục đích thông dụng hiện

nay :

Một trong những khó khăn trong lập trình ứng dụng GPGPU đó là các tác vụ

general-purpose không liên quan gì tới đồ họa, nhưng các ứng dụng vẫn phải được

22

lập trình bằng cách sử dụng các API đồ họa. Các lập trình viên muốn truy cập vào

các đơn vị lập trình được trực tiếp trong khi chương trình đã được cấu trúc trong điều

kiện của đường ống đồ họa, với các đơn vị lập trình được chỉ có thể truy cập được

như một bước trung gian trong đường ống.

1) Xác định tên miền tính toán như một lưới cấu trúc của các luồng(thread).

2) Chương trình general-purpose SPMD tính giá trị của từng luồng.

3) Giá trị cho mỗi luồng được tính bằng cách thực hiện đồng thời kết hợp các

Hiện nay ứng dụng tính toán GPU được tổ chức theo cách sau:

4) Bộ nhớ đệm chứa kết quả tính toán của bộ nhớ toàn cục sẽ là đầu vào cho

phép toán toán học và cả truy cập đọc và ghi bộ nhớ toàn cục.

các tính toán tiếp theo.

Mô hình lập trình này mạnh vì nguyên nhân sau. Thứ nhất, bằng cách xác định

rõ ràng cơ chế song song trong chương trình nó cho phép các phần cứng khai thác

triệt để cơ chế song song dữ liệu của các ứng dụng. Thứ hai, bằng việc tạo ra sự cân

bằng vững chắc giữa sự hạn chế để đảm bảo hiệu năng tốt (mô hình SPMD, có các

hạn chế về phân nhánh cho hiệu quả, có hạn chế về dữ liệu giao tiếp giữa các

thành phần và giữa hạt nhân /chu kỳ, v.v..) và tính phổ biến (một thủ tục hoàn toàn

có thể lập trình tại mỗi phần tử) cho thấy điểm ấn tượng của nó. Cuối cùng nó đã

loại bỏ nhiều thách thức phức tạp của các lập trình viên GPGPU trước đây trong

việc đồng thời chọn giao diện đồ họa cho lập trình mục đích thông dụng nhờ khả

năng truy cập trực tiếp đến các đơn vị lập trình được.

Từ đó các chương trình thường được thể hiện bằng ngôn ngữ lập trình quen

thuộc hơn (chẳng hạn như ngôn ngữ lập trình của NVIDIA là CUDA), đơn giản

hơn, dễ dàng hơn để lập trình và gỡ lỗi. Kết quả ta có một mô hình lập trình cho

phép người dùng tận dụng đầy đủ các sức mạnh phần cứng của GPU

1.2.5. Các hỗ trợ phần mềm cho xử lý tính toán trên GPU

Trước đây phần lớn các chương trình GPGPU được thực hiện trực tiếp

thông qua các API đồ họa. Nhưng có một điều không phù hợp cơ bản giữa mô hình

lập trình truyền thống mà mọi người đang dùng và các mục tiêu của các API đồ họa

23

dù các nhà nghiên cứu đã thành công trong việc làm cho các ứng dụng làm việc

thông qua các API đồ họa. Vào thời gian đầu người ta sử dụng các hàm cố định, các

đơn vị đồ họa cụ thể (ví dụ như các bộ lọc kết cấu (texture filter), trộn (blending),

và các phép toán tạo mẫu tô đệm để thực hiện các thao tác GPGPU. Cách này vẫn

khó tiếp cận cho dù điều này tốt hơn với phần cứng là bộ xử lý các mảnh hoàn

toàn lập trình được với ngôn ngữ assembly mã giả. Lập trình đổ bóng cao cấp đã

được thực hiện có thể thông qua ngôn ngữ đổ bóng cấp cao ("high-level shading

language” - HLSL), Với DirectX 9. Nó được biểu diễn giống như giao diện lập

trình C cho lập trình đổ bóng. NVIDIA Cg giờ đây đã có thể cung cấp các tính năng

tương tự như HLSL, và hữu ích hơn là đã có thể biên dịch ra nhiều đích và cung cấp

ngôn ngữ lập trình cấp cao đầu tiên cho OpenGL. hiện nay ngôn ngữ đổ bóng

OpenGL là ngôn ngữ đổ bóng tiêu chuẩn cho OpenGL. Mặc dù vậy chúng vẫn là

ngôn ngữ đổ bóng và do vậy nên chúng vẫn khó tiếp cận với một lập trình viên thông

thường.

Một ngôn ngữ cấp cao hơn được thiết kế để tính toán một cách rõ ràng và trừu

tượng hóa tất cả các cơ chế đồ họa của GPU là những gì các nhà phát triển thực sự

muốn và có được. Với mục tiêu trừu tượng GPU như là bộ xử lý dòng (streaming

processor), BrookGPU và Sh là hai đầu dự án nghiên cứu đầu tiên về vấn đề đó. 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 điều nay hoàn toàn thuận lợi với tài nguyên

sẵn có trên GPU. Một chương trình dòng bao gồm một tập hạt nhân (kernel), các

dòng (stream), các tập được sắp xếp dữ liệu, và các hàm chức năng được thiết lập

với từng phần tử trong tập các dòng tạo ra một hay nhiều dòng đầu ra.

Brook tiếp cận theo hướng trừu tượng tính toán dòng đơn giản, để biểu diễn

dữ liệu như là các dòng và tính toán như là các hạt nhân. Không có khái niệm

về kết cấu vector, mảnh, hoặc trộn (blending) trong Brook. Hạt nhân là các tính

toán được viết trong một tập hợp con giới hạn của C, đặc biệt là không có con trỏ

và scatter (sự tán xạ- thao tác ghi bộ nhớ), với đầu vào, đầu ra định nghĩa trước, và

trùm các dòng được sử dụng trong hạt nhân như một phần của định nghĩa của nó.

24

Brook chứa các chức năng truy cập dòng như: rút gọn các dòng, lặp lại và thoát

khỏi vòng lặp, tập con các dòng để sử dụng như đầu vào và đầu ra và khả năng xác

định tên miền. Hạt nhân của người dùng được ánh xạ tới đoạn code đổ bóng cho

mảnh và đến các dòng liên quan tới kết cấu. Những hạt nhân được chạy cho mỗi

phần tử trong miền các dòng đầu ra.

Thông qua các lời gọi đọc / ghi rõ ràng được phiên dịch thao tác cập nhật kết

cấu và cập nhật vào bộ đệm phản hồi dữ liệu tải lên và tải về GPU. Sau đó, tính

toán được thực hiện bởi một biến đổi vào không gian 3 chiều vùng các điểm ảnh

trong miền đầu ra.

Thay vì sử dụng biên dịch offline, dự án Microsoft’s Accelerator (bộ gia tốc

của Microsoft) có mục tiêu tương tự như Brook tập trung vào khía cạnh tính toán,

nhưng bộ gia tốc dựa vào biên dịch tức thời (just-in-time) của các phép toán dữ liệu

song song cho bộ đổ bóng mảnh. Bộ gia tốc là ngôn ngữ dựa trên mảng (array-base

language) phát triển từ ngôn ngữ C #, và tất cả các tính toán được thực hiện thông

qua các phép toán trên các mảng, không giống như mô hình của Brook và Sh

được phần lớn các phần mở rộng từ C. Khác với Brook, nhưng giống Sh, mô hình

đánh giá độ trễ cho biên dịch tức thời hiệu suất cao hơn làm cho khả năng chuyên

biệt hơn và tối ưu code tạo ra để thực hiện trên GPU.

Đã có những thay đổi lớn trong môi trường phần mềm cho phép phát triển

các ứng dụng GPGPU dễ dàng hơn nhiều và tạo ra các hệ thống phát triển mạnh mẽ

hơn, chất lượng thương mại hơn trong thời gian vừa qua. RapidMind đã thương mại

hóa Sh và hiện nay đang đặt mục tiêu nhiều platform trong một GPU, CPU đa

lõi, các STI Cell Broadband Engine, và hệ thống mới tập trung nhiều hơn nữa vào

tính toán so với SH trong việc bao gồm nhiều phép toán đồ họa trung tâm.

RapidMind sử dụng ước lượng độ trễ và biên dịch online để chụp lại và

tối ưu hóa mã nguồn ứng dụng của người dùng cùng với các phép toán và mở

rộng kiểu của C ++ để tạo ra những hỗ trợ trực tiếp cho mảng giống như bộ gia

tốc của Microsoft. , Được thiết kế xoay quanh các phép toán trên mảng, PeakStream

là hệ thống mới, sáng tạo từ Brook. Cũng giống như bộ gia tốc và RapidMind,

25

PeakStream chỉ sử dụng trong biên dịch tức thời, nhưng linh hoạt hơn nhiều trong

việc vector hóa code của người dùng nhằm đạt hiệu suất cao nhất trên kiến trúc

SIMD. Khía cạnh mà là một vấn đề hóc búa trong phát triển GPGPU được

PeakStream cung cấp platform đầu tiên hỗ trợ profiling và gỡ lỗi. Các nỗ lực này

giúp cho các nhà cung cấp của bên thứ ba tạo các hệ thống với sự hỗ trợ từ các

nhà cung cấp GPU. Với lợi ích trên, Google đ ã mua PeakStream trong năm 2007.

Cả AMD và NVIDIA hiện nay đều sở hữu riêng hệ thống lập trình GPGPU.

AMD công bố hệ thống r i ên g của họ vào cuối năm 2006. CTM, hay "Close To

The Metal", cung cấp mức trừu tượng phần cứng ở cấp thấp (HAL) cho dòng

R5XX và dòng R6XX của GPU ATI. CTM-HAL cung cấp truy cập mức assembly

thô cho động cơ mảnh (bộ xử lý dòng - stream processor) cùng với bộ lắp ráp và

bộ đệm lệnh để điều khiển thực thi trên phần cứng. Với các giao diện này không tính

năng đồ họa cụ thể nào được xuất. Các phép tính toán được thực hiện bằng cách

ràng buộc bộ nhớ như là đầu vào và đầu ra các bộ vi xử lý dòng, định nghĩa một

miền các kết quả đầu ra mà trên đó để thực thi nhị phân, tải mã nhị phân ELF.

Hãng AMD cũng tiếp tục đưa ra tầng trừu tượng tính toán - Compute Abstraction

Layer (CAL). Tầng này đưa thêm các cấu trúc (construct) cấp cao hơn, giống như

thành phần tương tự trong hệt thống chạy của Brook, và hỗ trợ biên dịch GPU

ISA cho GLSL, HLSL, và mã giả Assembly như Pixel Shader 3.0. Hãng cũng

AMD hỗ trợ biên dịch các chương trình Brook trực tiếp đến phần cứng R6XX,

cung cấp một mức lập trình trừu tượng cao hơn so với CAL hoặc HAL đối với lập

trình cấp cao hơn. Cũng giống như Brook, CUDA cung cấp một cú pháp giống C

để thực hiện trên GPU và biên dịch offline.

CUDA khai thác hai cấp xử lý song song là song song dữ liệu và đa luồng,

không giống như Brook chỉ khai thác một hướng xử lý song song là song song

dữ liệu thông qua cơ chế dòng. 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; bộ

nhớ bo mạch, và bộ nhớ máy chủ, 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. Các hạt nhân trong CUDA cũng linh

26

hoạt hơn bằng việc cho phép sử dụng con trỏ (nhưng dữ liệu phải ở trên bo mạch),

việc lấy ra/lưu trữ thông thường vào bộ nhớ cho phép người sử dụng đồng bộ giữa

các luồng trong một khối luồng, và tán xạ (scatter) dữ liệu từ bên trong một hạt

nhân. Nhưng tất cả sự hữu ích này và kết quả tiềm năng đạt được đi kèm với việc

yêu cầu người sử dụng phải biết nhiều hơn các chi tiết ở cấp thấp của phần cứng,

như là sử dụng thanh ghi, các hành vi của các mẫu truy cập bộ nhớ và luồng và

lập lịch cho khối luồng.

Tất cả các hệ thống này cho phép người phát triển xây dựng các ứng dụng

lớn dễ dàng hơn. CUDA cung cấp điều chỉnh và tối ưu hóa thư viện Blas và FFT

để sử dụng như xây dựng khối cho các ứng dụng lớn. Truy cập hệ thống GPGPU

cụ thể như CUDA, hoặc truy cập cấp thấp vào phần cứng, như là cung cấp bởi CTM,

cho phép các người phát triển vượt qua một cách có hiệu quả các trình điều khiển

đồ họa và duy trì ổn định hiệu năng và tính đúng đắn. Việc tối ưu được thực hiện để

tối ưu hóa cho hiệu năng game có thể ảnh hướng tới tính ổn định và hiệu năng của

các ứng dụng GPGPU.

1.2.6. Các kỹ thuật tính toán trên GPU

Chúng ta khảo sát một số đặc tính tính toán quan trọng, ứng dụng tính toán

GPU và các thuật toán. Chúng ta tìm hiểu bốn phép toán song song dữ liệu tập trung

ở tính toán GPU: ánh xạ một chức năng vào nhiều yếu tố song song, thực hiện phép

toán tán xạ (scatter) / tập hợp (gather) bộ nhớ, tính toán rút gọn cho trước một

mảng song song, giảm một bộ sưu tập các yếu tố thành một yếu tố hoặc một giá

trị. Chúng ta nghiên cứu tính toán cơ bản cốt lõi ở một số chi tiết trước khi nghiên

cứu với một cách nhìn tổng quan mức cao về các vấn đề thuật toán mà các nhà

nghiên cứu đã nghiên cứu trên GPU: phương trình vi phân, đại số tuyến tính, quét,

tìm kiếm, sắp xếp, truy vấn dữ liệu. Các thuật toán cho phép nhiều ứng dụng khác

nhau, từ khai phá dữ liệu, cơ sở dữ liệu, đến các mô phỏng khoa học, như là chuyển

động nhiệt của chất lỏng chuyển động vật lý trong trò chơi và động lực học phân tử

và động lực học.

27

Tính toán nguyên thủy:

Thường với các lập trình viên ngày nay trưởng thành từ máy tính tuần tự hoặc

cụm máy tính kết nối lỏng lẻo nhưng các kiến trúc song song dữ liệu của GPU

đòi hỏi thuật ngữ lập trình quen thuộc từ lâu với người sử dụng siêu máy tính song

song. Chúng ta nghiên cứu cơ bản về bốn các yếu tố quan trọng: ánh xạ, rút gọn,

tán xạ / tập hợp (scatter/gather), và quét. Ta tìm hiểu những tính toán nguyên thủy

này trong bối cảnh cả dựa trên đồ họa và tính toán trực tiếp trên tính toán GPU để

nhấn mạnh sự đơn giản và tính linh hoạt của cách tiếp cận tính toán trực tiếp.

Tán xạ/tập hợp :

Đọc ra hoặc viết vào một vị trí được tính toán trong bộ nhớ. Tính toán GPU

dựa trên đồ họa cho phép tập hợp hiệu quả bằng cách lưu trữ dữ liệu như hình ảnh

kết cấu và đánh địa chỉ dữ liệu bằng cách tính toán tọa độ hình ảnh tương ứng và

thực hiện phép nạp kết cấu, sử dụng các hệ thống con về kết cấu. Tuy vậy, các hạn

chế về kết cấu làm cho khó phát triển thêm nữa: phép nạp kết cấu đơn chỉ có thể lấy

4 giá trị dấu phảy động 32bit, hạn chế bộ nhớ lưu trữ mỗi phần tử, hạn chế kích

thước kết cấu đòi hỏi các mảng chứa trên 4.096 phần tử thành nhiều dòng của một

kết cấu 2D, bổ sung thêm phép toán đánh địa chỉ. Phép tán xạ trong tính toán GPU

đòi hỏi phải tái liên kết dữ liệu để thực thi như là các vector, hoặc sử dụng phép

nạp kết cấu đỉnh hoặc render- to-vertex-buffer và dựa trên đồ họa khó khăn. Ngược

lại lớp trực tiếp tính toán cho phép đọc và ghi không giới hạn đến các địa điểm

tùy ý trong bộ nhớ. CUDA của NVIDIA cho phép người dùng truy cập vào bộ nhớ

bằng cách sử dụng các cấu trúc C chuẩn (mảng,con trỏ, biến); CTM của AMT cũng

gần linh hoạt được như vậy, nhưng sử dụng địa chỉ 2D.

Ánh xạ (Map): ): Sử dụng một phép toán trên mọi phần tử trong bộ sưu tập.

Tiêu biểu là vòng lặp for trong chương trình tuần tự (như là một luồng trên một

CPU đơn lõi). Một tác vụ song song có thể giảm tối đa thời gian cần thiết khi áp

dụng phép toán đó đến nhiều phần tử song song. Như là chương trình mảnh được

28

gọi từ bộ sưu tập điểm ảnh (một điểm ảnh cho mỗi phần tử), tính toán GPU dựa

trên đồ họa thực hiện phép ánh xạ. Đ i ể m ả n h đ ư ợ c đọc dữ liệu từ kết cấu tại

một ví trí tương ứng với vị trí của điểm ảnh trong hình ảnh đã biến đổi, thực thi

phép toán đó bởi từng chương trình mảnh sau đó lưu trữ các kết quả tại điểm ảnh

đầu ra. Giống như vậy CTM và CUDA sinh ra một chương trình luồng để thực

hiện phép toán đó trong nhiều luồng, với mỗi luồng nạp vào một phần tử, thực hiện

tính toán, và lưu trữ kết quả. Và vì vòng lặp hỗ trợ mỗi luồng có thể cũng lặp nhiều

lần trên nhiều phần tử.

Rút gọn (Reduce): Để rút gọn một tập hợp các phần tử thành một phần tử

duy nhất hoặc một giá trị duy nhất, liên tục áp dụng một phép toán kết hợp nhị

phân. Ví dụ bao gồm việc tìm kiếm tổng (tối thiểu, tối đa, phương sai, trung bình,

vv...) của một tập các giá trị. Một thực thi tuần tự sẽ lặp trên một mảng, tính tổng

từng phần tử bằng cách chạy phép cộng tât cả các phần tử hiện có, đó là trên CPU

truyền thống. Và thực hiện nhiều lần phép cộng song song trên một tập thu hẹp các

phần tử trong một rút gọn tổng theo cơ chế song song. Các xử lý tính toán trên

GPU dựa trên đồ họa thực hiện rút gọn dựa trên biến đổi (rendering) tập giảm dần

các điểm ảnh. Trong từng biến đổi từng vượt qua chương trình mảnh đọc nhiều

giá trị từ một kết cấu (thực thi khoảng 4 hoặc 8 lần đọc kết cấu), tính tổng đó, và

ghi giá trị đó vào điểm ảnh đầu ra trong kết cấu khác (nhỏ hơn 4 hoặc 8 lần), mà

sau đó sẽ bị ràng buộc như là đầu vào cho bộ đổ bóng mảnh tương tự và quá

trình lặp đi lặp lại cho đến khi đầu ra là một điểm ảnh đơn chứa kết quả cuối

cùng của quá trình rút gọn.Tương tự như vậy, nhưng CTM và CUDA cùng cho ra

cùng một quá trình trực tiếp hơn, ví dụ bằng cách tạo ra một tập các luồng, mỗi

luồng dọc 2 phần tử và ghi tổng của chúng vào một phần tử đơn. Một nửa số luồng

lặp lại quá trình trên, sau đó là nửa còn lại, cứ như vậy cho đến khi còn lại một

luồng sống sót sẽ ghi kết quả cuối cùng ra bộ nhớ.

Quét (Scan):

29

Quét lấy một mảng A các phần tử và trả về một mảng B có cùng chiều dài,

trong đó mỗi phần tử B [i] đại diện cho một phép rút gọn mảng con A[1...i] được

gọi là tổng tiền tố song song. Công cụ quét là một công cụ xây dựng khối dữ

cực kỳ hiệu quả cho thuật toán song song dữ liệu. CUDA thực hiện nhanh hơn so

với CPU bởi một một thừa số lên đến 20 và OpenGL bởi một thừa số lên đến 7. Kết

quả minh họa cho những lợi thế của tính toán trực tiếp hơn là tính toán GPU dựa

trên đồ họa.

1.2.7.Các giải thuật ứng dụng trên GPU

Khi phát triển phần lớn vào các phép toán cơ bản trên, các chuyên gia đã

biểu diễn nhiều thuật toán mức cao và các ứng dụng khai thác các lợi thế tính toán

của GPU.

Sắp xếp (Sort): GPU đã có những bước tiến đáng kể trong sắp xếp từ khi

các nhà nghiên cứu GPU đã nghiên cứu lại, áp dụng, và cải thiện các thuật toán

sắp xếp, đáng chú ý là sắp xếp bitonic merge. Thuật toán "sorting network"

này về bản chất là song song và mù, có nghĩa là được thực hiện bất kể đầu vào

và tương tự nhau. Nhờ sử dụng hệ thống thiết kế cẩn thận và sự kết hợp của cải

tiến nhiều thuật toán Govindaraju và các đồng nghiệp đã giành giải hiệu năng

"PennySort" trong cuộc thi "TeraSort" năm 2005.

Tìm kiếm và truy vấn cơ sở dữ liệu : Các nhà phát triển cũng đã triển khai

một số hình thức tìm kiếm trên GPU, như các thuật toán sắp xếp nhanh ở trên, tìm

kiếm láng giềng gần nhất, tìm kiếm nhị phân, cũng như các thao tác cơ sở dữ liệu

được xây dựng trên phần cứng đồ họa mục đích đặc biệt (gọi là bộ đệm độ sâu

stencil) .

Phương trình vi phân : Cố gắng đầu tiên sử dụng GPU cho tính toán phi đồ

họa tập trung vào giải quyết các tập lớn phương trình vi phân. một ứng dụng GPU

phổ biến cho phương trình vi phân thường (ODEs) là phép tìm đạo hàm, được sử

dụng rất nhiều trong các hiệu ứng trực quan cho các chò trơi trên máy tính và

trong mô phỏng khoa học (ví dụ, hệ thống thăm dò lưu lượng của Kr¨uger ). GPU

30

đã được sử dụng nhiều để giải quyết các vấn đề trong phương trình vi phân riêng

(PDEs) như phương trình Navier- Stokes cho dòng chảy tự do. ứng dụng đặc biệt

thành công mà GPU PDE đã giải quyết bao gồm phương trình thiết lập phân chia

âm thanh và các động lực chất lỏng.

Đại số tuyến tính : chương trình đại số tuyến tính là các khối tạo dựng cốt

lõi cho một rất lớn các thuật toán số học, bao gồm cả giải pháp PDE . Ứng dụng

của nó bao gồm chứa mô phỏng các hiệu ứng vật lý như: bức xạ, hiệu ứng quang

học như lĩnh vực độ sâu và chất lỏng, nhiệt, và tương tự, qua đó vấn đề đại số tuyến

tính trên GPU đã nhận được nhiều sự chú ý. Điển hình là sản phẩm của Kr ¨uger và

Westermann đã giải quyết một lớp rộng của các vấn đề đại số tuyến tính bằng

cách tập trung vào biểu diễn ma trận và vectơ trong tính toán trên GPU dựa trên đồ

họa (đóng gói các vector dày đặc (dense) và thưa thớt (sparse) vào các kết cấu, bộ

đệm vector, v.v..). Một sản phẩm đáng chú ý khác là giải pháp cho các hệ thống

tuyến tính dày đặc của Gallapo và đồng nghiệp, và các phân tích về phép nhân ma

trận dày đặc của Fatahalian và đồng nghiệp. Tác giả đã cho thấy có hiệu suất lớn

hơn, và thậm chí các triển khai ATLAS tối ưu hoá mức cao. Ứng dụng

của các tầng trực tiếp tính toán như CTM và CUDA vừa đơn giản hoá đồng thời

cải thiện hiệu suất của đại số tuyến tính trên GPU. Như NVIDIA cung cấp uBLAS,

một gói đại số tuyến tính dày đặc thực thi trong CUDA và sau đó là các quy ước

BLAS phổ biến. Các nhà nghiên cứu mong có mã nguồn thưa thớt để kiểm chứng

lợi ích tương tự hoặc lớn hơn từ tầng tính toán mới GPU và các thuật toán đại

số tuyến tính thưa thớt có nhiều biến đổi và phức tạp hơn so với loại dày đặc đang là

một lĩnh mở và hướng nghiên cứu tích cực.

31

CHƯƠNG II: XỬ LÝ SONG SONG TRÊN THIẾT BỊ ĐỒ HỌA GPU VỚI CUDA

2.1. Khái quát về CUDA

Là từ viết tắt của thuật ngữ Compute Unified Device Architecture, CUDA tạm

dịch là kiến trúc thiết bị hợp nhất cho tính toán. Bắt đầu xuất hiện từ tháng bảy

năm 2007 với vai trò ban đầu là một bộ công cụ phát triển phần mềm dựa trên

ngôn ngữ lập trình C. Hiện nay CUDA đang tiến hóa thành kiến trúc điện toán

GPU, hay còn gọi là GPGPU của NVIDIA. CUDA có mặt trên hầu hết các GPU

đời mới của NVIDIA, từ Quadro giành cho điện toán hình ảnh chuyên nghiệp,

dòng Tesla cho tính toán hiệu năng cao và dòng GeForce giành cho giải trí.

Bộ phần mềm CUDA có các lớp mô tả trong Hình 11, gồm: API lập trình, bộ

điều khiển (dirver) cho phần cứng, hai thư viện toán học mức cao hơn của các

hàm thường dùng: CUFFT và CUBLAS và môi trường thực thi. Phần cứng được

thiết kế để hỗ trợ bộ điều khiển hạng nhẹ và lớp môi trường thực thi. Kết quả là làm

cho GPU có tốc độ cao .

Hình 11: Kiến trúc phần mềm CUDA

32

Thư viện lập trình của CUDA bao gồm các hàm mở rộng của ngôn ngữ C. Hình

12 mô tả CUDA cung cấp cách đánh địa chỉ DRAM thường dùng cho việc lập

trình linh hoạt hơn, bao gồm cả thao tác thu hồi bộ nhớ và cấp phát bộ nhớ. Giống

như CPU nhìn từ góc độ lập trình, điều đó tương ứng với khả năng ghi và đọc dữ

liệu tại bất kỳ địa chỉ nào trong DRAM.

Hình 12: Thao tác cấp phát và thu hồi bộ nhớ

Ngôn ngữ CUDA có đặc điểm lưu dữ liệu đệm song song, bộ nhớ chia sẽ trên

bộ vi xử lý 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. Ứng dụng có thể đạt kết quả tốt bằng cách tối thiểu việc lấy/trả dữ liệu từ

DRAM, như mô tả trong hình. Qua đó giảm phụ thuộc băng thông truyền bộ nhớ

DRAM.

33

Hình 13: Vùng nhớ dùng chung mang dữ liệu gần ALU hơn

2.2.Cơ chế lập trình và cách thức hoạt động của CUDA

2.2.1. Cơ chế lập trình

Cần phải có các thư viện hỗ trợ để chương trình CUDA hoạt động được

trong môi trường windows hoặc linux. Do NVIDIA cung cấp Các thư viện này,

gồm có các phần sau: Bộ công cụ phát triển CUDA (gọi là CUDA Toolkit) và bộ

CUDA SDK Trình điều khiển thiết bị đồ họa cho GPU của NIVIDA.

2.2.2.Cách thức hoạt động của CUDA

Mong muốn chương trình chạy nhanh hơn nhờ khả năng xử lý song song nên

sử dụng CUDA. Vậy ta cần loại bỏ các ảnh hưởng làm một chương trình chạy chậm

đi. Ảnh hưởng chính đến tốc độ của chương trình là sự không thống nhất và tranh

chấp vùng nhớ trong quá trình đọc và lưu dữ liệu. Một chương trình CUDA hoạt

động theo mô hình SIMD (single instruction multiple data) điều này buộc trình

biên dịch phải chọn giải pháp an toàn trong truy cập dữ liệu. Việc này biến một

chương trình song song theo mô hình SIMD thành mô hình nối tiếp.

Kích thước dữ liệu phải bằng 4, 8, 16 bytes vì kích thước của kiểu dữ liệu rất

quan trọng trong việc truy cập dữ liệu một cách thống nhất (coalescing). Để hạn

chế việc truy cập thường xuyên vào bộ nhớ chung làm chậm chương trình (do việc

truy cập vào bộ nhớ chung mất rất nhiều thời gian hơn truy cập vào bộ nhớ chia sẻ)

khi số lệnh tính toán lớn thì nên sao chép dữ liệu từ bộ nhớ chung (global memory)

vào bộ nhớ chia sẻ (shared memory).

34

Một chương trình CUDA thường sử dụng hai hàm: Một hàm kernel dùng cho

việc xử lý dữ liệu, Một hàm dành cho việc truy cập dữ liệu.

Để hiểu cách hoạt động một chương trình CUDA, cần thừa nhận một số các

khái niệm sau:

Hình 14: Sơ đồ hoạt động truyền dữ liệu giữa Host và Device

 Host: Là cấu trúc phần cứng, phần mềm và những tác vụ và được xử lý bởi

CPU.

 Device: Là cấu trúc phần cứng, phần mềm và những tác vụ và được xử lý bởi

GPU.

Cách hoạt động được mô tả như sau:

 Bước đầu tiên là sao chép dữ liệu cần tính toán từ bộ nhớ Host sang bộ nhớ

Device trước khi muốn thực hiện trên Device bởi dữ liệu cần tính toán luôn ở trên bộ

nhớ của Host.

 Sau khi sao chép dữ liệu, device sẽ thực hiện việc tính toán trên dữ liệu đó.

 Khi tính toán xong, dữ liệu cần được sao chép lại từ bộ nhớ Device sang bộ

nhớ Host.

Đồng xử lý đa luồng mức cao

GPU được xem như là một thiết bị tính toán có khả năng thực hiện một số

lượng rất lớn các luồng song song trong lập trình CUDA. GPU hoạt động như một

bộ đồng xử lý với CPU chính. Có thể nói phần tính toán chuyên dụng của các ứng

dụng chạy trên host, dữ liệu song song được tách rời khỏi thiết bị.

35

Được gọi là bộ nhớ host và bộ nhớ thiết bị, cả hai Host và Device duy trì

DRAM riêng của nó. Sử dụng cơ chế truy cập bộ nhớ trực tiếp tốc độ cao thiết bị,

có thể sao chép dữ liệu giữa DRAM của Host và Device thông qua API đã tối ưu

hóa.

Cụ thể một phần của một ứng dụng được thực hiện nhiều lần, độc lập về mặt

dữ liệu, và nhóm thành một chức năng được thực hiện trên thiết bị như nhiều

luồng khác nhau. Thực hiện bằng việc một chức năng được biên dịch thành các tập

lệnh của thiết bị và tạo ra chương trình, gọi là nhân, được tải vào thiết bị.

Gom lô cácluồng

Lô các luồng thực hiện được nhân (kernel) tổ chức thành một lưới các khối

luồng được thể hiện trong các phần sau.

Khối luồng

Một khối luồng là một tập các luồng, có thể đồng thời xử lý với nhau bằng

cách thực thi đồng bộ để phối hợp truy cập bộ nhớ và dùng dữ liệu trong bộ nhớ

dùng chung. Nói chính xác hơn, nơi các luồng trong khối sẽ dừng cho đến khi tất cả

các luồng tới điểm đồng bộ, có thể xác định các điểm đồng bộ trong nhân.

Mỗi luồng được xác định bởi số hiệu của luồng trong khối là ID. Để hỗ trợ việc

định địa chỉ phức tạp dựa trên ID luồng, một ứng dụng cũng có thể xác định từng

luồng bằng cách sử dụng chỉ số hai hoặc ba thành phần để thay thế hoặc chỉ định

một khối như một mảng hai hoặc ba chiều có kích thước tùy ý. Đối với các khối

kích thước hai chiều (Dx, Dy), ID luồng của phần tử có chỉ số (x, y) là (x + y Dx)

và đối với khối kích thước ba chiều (Dx, Dy, Dz), ID luồng của phần tử (x, y, z) là (x

+ yDx + z Dx Dy) .

36

Lưới các khối luồng (Grid of Thread Blocks)

Hình 15: Khối luồng

Số lượng luồng tối đa trong một khối có giới hạn. Các khối cùng số chiều

và kích thước thực thi trên cùng nhân có thể nhóm với nhau thành lưới các khối nên

tổng số luồng chạy trên một nhân là lớn hơn nhiều. Vì các luồng trong các lô khác

nhau trong lưới không thể trao đổi và đồng bộ với nhau nên chi phí hợp tác giữa

các luồng giảm. Như mô tả, cho thấy các nhân chạy hiệu quả mà không phải dịch

lại trên các loại thiết bị với khả năng chạy song song khác nhau : Một thiết bị có thể

chạy chạy song song nếu nó có khả năng chạy song song nhiều hoặc chạy trên tất

cả khối của lưới một cách tuần tự nếu thiết bị đó có rất ít khả năng chạy song song

hoặc kết hợp cả hai.

37

Mỗi khối được xác định bởi số khối trong lưới đó là ID của nó. Để hỗ trợ

việc định địa chỉ phức tạp dựa trên khối ID, một ứng dụng có thể xác định một lưới

như một mảng hai chiều với kích thước cố định và định danh mỗi khối sử dụng chỉ

mục hai thành phần. Với khối hai chiều kích thước (Dx, Dy), ID của khối (x,y) là (x

+ y Dx).

Cấu trúc bộ nhớ

Hình 16: Mô hình bộ nhớ trên GPU

Như mô tả trong Hình 16 một luồng thực thi trên thiết bị chỉ truy cập vào

DRAM của thiết bị và bộ nhớ trên bộ vi xử lý qua các không gian nhớ :

 Đọc và ghi bộ nhớ toàn cục (Global Memory) của mỗi lưới.

 Chỉ đọc bộ nhớ hằng số (Constant Memory) của mỗi lưới.

 Chỉ đọc bộ nhớ kết cấu (Texture Memory) của mỗi lưới.

38

 Đọc và ghi bộ nhớ cục bộ (Local Memory) của mỗi luồng.

 Đọc và ghi trên các thanh ghi (Registers) của mỗi luồng.

 Đọc và ghi bộ nhớ dùng chung (Shared Memory) của mỗi khối.

Các vùng nhớ toàn cục, hằng số và kết cấu được tối ưu hóa cho các cách sử

dụng bộ nhớ khác nhau. Vùng nhớ kết cấu cũng đưa ra các cơ chế đánh địa chỉ khác,

cũng như lọc dữ liệu cho một số loại dữ liệu đặc biệt.

Các vùng nhớ toàn cục, hằng số và kết cấu có thể đọc hoặc ghi bởi Host và liên

tục giữa các lần thực thi nhân bởi cùng một ứng dụng.

2.3. Tổng quan về lập trình với CUDA

2.3.1. Là ngôn ngữ lập trình mở rộng của ngôn ngữ lập trìnhC

Để có thể dễ dàng viết chương trình cho việc xử lý bằng các thiết bị, mục tiêu

của giao diện lập trình CUDA là cung cấp cách tiếp cận khá đơn giản cho những

người sử dụng quen với ngôn ngữ lập trình C. Lập trình CUDA gồm có:

 Thư viện chạy được chiathành:

 Một thành phần chung (commom componet): Cung cấp xây dựng

trong kiểu vector và là một tập con thư viện chuẩn của C. Thành phần

chung hỗ trợ cho cả Host và các thiết bị thành phần.

 Các thiết bị thành phần (device componet): Được chạy trên các thiết

bị và cung cấp các hàm riêng của thiết bị đó.

 Thành phần chính (host componet): Chạy trên Host và cung cấp các

chức năng cho việc điều khiển và truy nhập một hoặc nhiều thiết bị

khác từ Host.

Nên chú ý rằng chỉ có hàm từ thư viện chuẩn của C là được hỗ trợ cho việc

chạy trên các thiết bị có các chức năng được cung cấp bởi thành phần chạy chung .

2.3.2. Các phần mở rộng của CUDA

Ngôn ngữ lập trình CUDA là mở rộng của ngôn ngữ lập trình C ở bốn khía cạnh

 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.

39

 Từ khóa phạm vi kiểu hàm cho phép xác định liệu một hàm thực hiện

trên host hay trên thiết bị và nó có thể được triệu gọi từ host hoặc từ thiết

bị.

 Bốn biến build-in để xác định chiều của lưới và khối, chỉ số khối và luồng.

 Một chỉ thị mới để xác định cách nhân được thực hiện trên thiết bị từ phía

host.

NVCC sẽ đưa ra lỗi hoặc thông điệp cảnh báo một số xung đột của các phần

hạn chế, nhưng một số xung đột có thể không được nhận ra

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:

 device

:

+ Tồn tại trong không gian bộ nhớ toàn cục (có bộ nhớ lớn, độ trễ cao).

+ Được cấp phát với cudaMalloc.

+ Có vòng đời (lifetime) của một ứng dụng .

+ Truy nhập được từ tất cả các luồng bên trong lưới

 shared

:

+ Tồn tại trong không gian bộ nhớ chia sẻ của một luồng (bộ nhớ

nhỏ,độ trễ thấp).

+ Được cấp phát khi thực hiện việc cấu hình, hay khi biên dịch chương

trình.

+ Có vòng đời của một khối.

+ Chỉ có thể truy cập từ tất cả các luồng bên trong một khối (các luồng

thuộc khối khác không thể truy cập).

Từ khóa phạm vi kiểu hàm

40

Dùng để khai báo một hàm có phạm vi hoạt động ở trên Host hay trên Device,

và được gọi từ Host hay từ Device:

 Từ khóa device

:

+ Khai báo_ device: định nghĩa một hàm chỉ xử lý trên thiết bị (Device)..

+ Chỉ được gọi từ thiết bị.

+ Ví dụ: device_void HamDevice(parameter,…) {…}

 Từ khóa global

:

+ Khai báo trên thiết bị. global định nghĩa một hàm như là một hạt nhân , xử

lý

+ Chỉ có thể triệu gọi được từ Host

+ Ví dụ : global void HamTaiNhan(parameter,…){…}

 Từ khóa host :

+ Khai báo Host là định nghĩa một hàm xử lý trên Host.

+ Chỉ có thể triệu gọi được từ Host.

Các hạn chế:

 Các hàm của device

là hàm đóng (inlined).

 Các hàm của device và global không hỗ trợ sự đệ quy.

 Các hàm của device và global không thể khai báo các biến static

trong thân hàm.

 Các hàm của device và global không thể có số biến của thay đổi.

 global và host không thể sử dụng đồng thời. global phải có

kiểu trả về là kiểu void.

 Lời gọi hàm global phải chỉ rõ cấu hình thực hiện nó.

 Gọi tới một hàm __global

là không đồng bộ, có nghĩa là hàm global

trả về trước khi thiết bị hoàn thành xong xử lý .

Thực hiện cấuhình

Các lời gọi tới hàm toàn cục (global) phải xác định cấu hình thực hiện cho

41

lời gọi. Để sử dụng thực hiện chức năng trên thiết bị cần cấu hình xử lý xác định

kích thước lưới và khối. Được xác định bằng cách chèn một biểu thức mẫu dạng

<<< Dg, Db, Ns >>> giữa tên hàm và danh sách tham số được để trong ngoặc đơn, ở

đây:

 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ớ tĩnh. Ns Là một đối số tùy chọn mặc định là 0, 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.

 Db là kiểu dim3 và xác định mục đích v à kích thước của mỗi khối, sao

cho Db.x*Db.y*Db.z bằng số lượng các luồng trên khối.

 Dg là kiểu dim3 và xác định mục đích và kích thước của lưới, sao cho

Dg.x*Dg.y bằng với số khối được đưa ra.

Một ví dụ cho việc khai báo hàm:

global void Func(int*parameter);

Phải gọi hàm từ Host giống như sau :

Func<<>>(parameter);

2.3.3.BiếnBuilt-in trong CUDA

Biến build-in để xác định chiều của lưới và khối, chỉ số khối và luồng :

 blockDim là biến kiểu dim3 và chứa kích thước của một khối.

 threadIdx là biến kiểu unit3 và chứa các chỉ số luồng trong khối

 gridDim là biến kiểu dim3 và chứa các kích thước của lưới.

 blockIdx là biến thuộc kiểu unit3 và chứa các chỉ số khối trong lưới.

42

Hình 17: Chiều của lưới và khối với chỉ số khối và luồng

2.3.4. Biên dịch CUDA thông qua NVCC

Bao gồm biên dịch mã thiết bị sang dạng nhị phân hoặc các đối tượng cubin và

luồng công việc cơ bản trong việc tách mã thiết bị từ mã Host. Mã Host sinh ra là

đầu ra có thể là mã C để được biên dịch bằng cách sử dụng một công cụ khác.

Ứng dụng tải đối tượng cubin vào thiết bị và khởi động mã thiết bị sử dụng

trình điều khiểu API của CUDA hoặc liên kết tới mã Host sinh ra, trong đó bao gồm

các đối tượng cubin được xem như mảng dữ liệu khởi tạo toàn cục và chứa một bản

dịch các cú pháp thực thi cấu hình thành mã cần thiết khởi động trong thời gian

chạy CUDA để nạp và khởi động mỗi lần biên dịch hạt nhân.

Là một trình điều khiển biên dịch bằng việc đơn giản hóa quá trình biên

dịch mã CUDA. NVCC cung cấp các dòng lệnh đơn giản và quen thuộc thực hiện

chúng bằng cách gọi tập hợp của các công cụ với công đoạn biên dịch khác nhau.

Frond end của trình biên dịch xử lý các tập tin nguồn CUDA theo cú pháp

quy định C++. Tuy nhiên, chỉ có các tập con C của C++ được hỗ trợ. Điều này có

nghĩa là những đặc tính đặc trưng của C++ như các lớp (classes), sự kế thừa hoặc

việc khai báo các biến trong khối cơ bản là không được hỗ trợ. Như một hệ quả

của việc sử dụng cú pháp C++, con trỏ void (ví dụ như trả lại malloc()) không thể được

gán tới những con trỏ non-void mà không có ép kiểu .

2.3.5.Một số trường hợp cụ thể tính toán song song bằngCUDA

 Cộng hai số nguyên:

Code tuầntự

void CongHaiSoNguyen(int *a,int *b, int *c)

{

*c=*a+*b;

}

43

void main()

{

int *a,*b,*c; CongHaiSoNguyen(a,b,c);

}

Code CUDA

global void KernelCongHaiSoNguyen(int *a,int *b,int*c)

{

*c=*a+*b;

}

void main()

{

int *a,*b,*c;

*a=1; *b=5;

int *deva,*devb,*devc; cudaMalloc((void**)&deva, sizeof(int)

);

cudaMalloc((void**)&devb, sizeof(int) );

cudaMalloc((void**)&devc, sizeof(int) );

cudaMemcpy(deva, a, sizeof(int),

cudaMemcpyHostToDevice); cudaMemcpy(devb, b,

sizeof(int), cudaMemcpyHostToDevice);

KernelCongHaiSoNguyen<<<1,1>>>(deva, devb, devc);

cudaMemcpy(c, devc, sizeof(int), cudaMemcpyDeviceToHost);

}

 Cộng hai mảng số nguyên:

Cộng hai mảng số nguyên a[n] và b[n]

global void KernelAdd(int *a, int *b, int*c)

{

44

c[blockIdx.x]= a[blockIdx.x] + b[blockIdx.x];

}

Trên thiết bị, mỗi block sẽ thực hiện song song:

 Block 0 thực hiện: c[0]= a[0] +b[0];

 Block 1 thực hiện: c[1]= a[1] +b[1];

 Block 2 thực hiện: c[2]= a[2] +b[2];

 Block 3 thực hiện: c[3]= a[3] +b[3];

 Block 4 thực hiện: c[4]= a[4] +b[4];

 Block 5 thực hiện: c[5]= a[5] +b[5];

 Block 6 thực hiện: c[6]= a[6] +b[6];

 Block n-1 thực hiện: c[n-1]= a[n-1] +b[n-1];

Cách giải quyết thứ hai ta dùng luồng để song song, cấu hình gọi hàm sẽ là

<<<1,n>>> (một block với nhiều luồng):

Code song song của chúng ta sẽ là :

global void KernelAdd(int *a, int *b, int*c)

{

c[threadIdx.x]= a[threadIdx.x] + b[threadIdx.x];

}

Như vậy chúng ta đã thấy việc song song dùng :

 Nhiều block với một luồng cho mỗi block.

 Một block với nhiều luồng.

Cách giải quyết thứ ba là kết hợp cả block và luồng : ta hãy xem cách đánh

chỉ số của một mảng với một phần tử của mảng cho mỗi thread (8thread/block) như

trong Hình.

45

Hình 18: Phương pháp đánh chỉ số luồng.

 Với M thread/block, một chỉ số duy nhất cho mỗi luồng sẽ là:

int index = threadIdx.x + blockIdx.x * M;

 Dùng biến built-in blockDim.x (tương ứng với số lượng luồng trong một

block) thay cho M ta được:

int index = threadIdx.x + blockIdx.x * blockDim.x;

 Vậy code song song của chúng ta sẽlà:

global void KernelAdd(int *a, int *b, int*c)

{

int index= threadIdx.x + blockIdx.x * blockDim.x;

c[index]= a[index] + b[index];

}

void main()

{

……;

KernelAdd<<

>>>(deva,devb,devc);

……;

}

2.4. Các ứng dụng của CUDA trong các lĩnh vực

2.4.1. Ứng dụng của CUDA trong game

Ứng dụng của công nghệ CUDA trong ngành công nghiệp giải trí với lĩnh vực

46

trò chơi là sự thành công lớn. Do NVIDIA cung cấp, hình ảnh trong trò chơi trở nên

như thật nhờ bộ công cụ PhysX SDK và khung hình làm việc có khả năng mở rộng

động trên nhiều nền tảng có liên quan với nó gọi là APEX. Giành riêng cho vật lý

trong trò chơi đây là những công cụ đầy sức mạnh trong bộ các engine AXE, hay

nói cách khác được thiết kế để xử lý các di chuyển sinh động và tương tác của đối

tượng trong từng cảnh của trò chơi.

Tính năng đồ họa hiệu ứng vật lý trong trò chơi khiến cho trò chơi trở nên

sống động, và thời gian thực trong trò chơi sẽ trở thành hiện thực với sự hỗ trợ của

PhysX và APEX chẳng bao lâu nữa chuyện hiển thị cảnh như phim.

Với hơn 150 tựa trò chơi mới trên thị trường, bộ công cụ PhysX SDK hiện

nay đã có trên hầu hết các nền tảng máy trò chơi thông dụng, từ XBOX 360 sang

PlayStation 3 sang Wii rồi đến NVIDIA GPU.

2.4.2. Ứng dụng của CUDA với video số

CUDA rất thành công trong việc xử lý video. Rất nhiều ứng dụng video số

hóa dựa trên CUDA, chẳng hạn như mở rộng độ phân giải DVD với SimHD của

ArcSoft, cải tiến chất lượng hình ảnh video với phần mềm vReveal của MotionDSP.

Một trong số các ứng dụng hay này là vReveal đến từ MotionDSP, là phần mềm

cải thiện chất lượng hình ảnh như: điều chỉnh độ tương phản và ổn định hóa (xóa

run) các video, làm rõ nét. Trước kia vReveal thường sử dụng đến các hệ thống

CPU đa bộ vi xử lý đắt tiền để hiển thị video một cách chậm chạp. Bây giờ với

CUDA GPU đã có thể thực hiện công việc trên theo thời gian thực hoặc nhanh hơn

khoảng năm lần so với CPU. Đối với lĩnh vực tình báo và điều tra pháp luật

MotionDSP còn cung cấp một cung cấp một phiên bản cao cấp hơn, gọi là Ikenna.

Ngày nay khả năng thu dữ liệu hình ảnh, video với chất lượng cao của những

thiết bị di động đã khiến con người thỏa mái hơn trong việc thưởng thức âm nhạc,

phim, hình chụp cá nhân ở mọi lúc, mọi nơi. Tuy nhiên, cần phải có những nỗ lực

của riêng mình khi sử dụng chúng. Ví dụ như phải tốn thời gian để chuyển đổi nhạc,

phim trong một chiếc iPod Touch sang máy để bàn của mình và ngược lại. Quá

trình đó thường mất nhiều thời gian và tài nguyên. Phần mềm Badaboom của

47

Elemental Technologies có thể giúp ích rất nhiều trong trường hợp đó. Đó là bộ

chuyển đổi media nhanh nhất và khi so sánh bộ chuyển định dạng cuariTunes,

Badaboom có thể nhanh hơn đến 20 lần hoặc tối thiểu cũng nhanh hơn hai đến ba

lần ngay khi sử dụng CPU nhanh nhất và đắt tiền nhất hiện nay.

48

CHƯƠNG III: SỬ DỤNG GPU ĐỂ LÀM TĂNG TỐC ĐỘ TÍNH TOÁN CHO BÀI TOÁN MÃ HÓA AES

3.1 Giới thiệu về AES

Là một hệ mã khóa bí mật có tên là Rijdael cho phép xử lý các khối dữ liệu input

có kích thước 128 bit sử dụng các khóa có độ dài 128, 192 hoặc 256 bit. Được Ủy

ban tiêu chuẩn của Hoa Kỳ đưa ra năm 2001, hệ mã Rijdael được thiết kế để có thể

làm việc với các khóa và các khối dữ liệu có độ dài lớn hơn tuy nhiên khi được chọn

là một chuẩn, nó được qui định chỉ làm việc với các khối dữ liệu 128 bit và các khóa

có độ dài 128, 192 hoặc 256bit.[5].

3.2 Thuật toán mã hóa

Được thực hiện bởi nhiều bước biến đổi tuần tự, đầu vào của các bước là kết quả

của bước trước đó. Một mảng trạng thái chứa kết quả trung gian của phép biến đổi

chính. Độ dài của khối dữ liệu đầu vào của AES là cố định với Nb=4, tùy vào độ dài

khoá (Nk=4,6,8) ban đầu ta có số lần lặp Nr cho mỗi quá trình được xác định theo

công thức Nr=max{Nb,Nk}+6

Công đoạn mã hóa và công đoạn giải mã AES sử dụng hàm lặp là hàm kết hợp

của bốn hàm biến đổi (với đơn vị xử lý là byte) sau:

- Thực hiện biến đổi thay thế byte thông qua sử dụng một bảng thế S-box

- Trong mảng trạng thái thực hiện dịch các hàng của với số lần dịch của mỗi

hàng là khác nhau.

- Kết hợp dữ liệu của mỗi cột trong mảng trạng thái

- Cộng một khóa RoundKey vào mảng trạng thái.

Đối với công đoạn giải mã AES được thực hiện bằng cách biến đổi ngược của các

biến đổi ở phép mã hóa AES hoặc bằng các biến đổi tương đương.[4] Các biến đổi

này được minh họa như trong hình vẽ:

49

(a) Mã hóa (b) Giải mã

Hình 19 : Mã hóa và giải mã

Tên hàm Giải thích

AddRoundKey() Được sử dụng trong công đoạn mã hóa và công đoạn

giải mã. Hàm này thực hiện phép XOR bit giữa trạng

thái trung gian (state) và một khóa vòng lặp (Round

Key). kích thước của trạng thái bằng Kích thước của

một Round Key.

MixColumns() Hàm này sử dụng trong công đoạn mã hóa, nhận tất

cả các cột của một trạng thái và trộn với dữ liệu của

nó để nhận được cột mới

ShiftRows() Hàm này sử dụng trong công đoạn mã hóa, xử lý các

trạng thái bằng cách dịch vòng ba hàng cuối của trạng

thái với số lần dịch khác nhau

50

SubBytes()

Hàm này sử dụng trong công đoạn mã hóa, xử lý một

trạng thái bằng cách sử dụng một bảng thế phi tuyến

các byte, thao tác trên mỗi byte một cách độc lập

InvMixColumns() Hàm này sử dụng trong công đoạn giải mã, là hàm

ngược của hàm MixColumns()

InvShiftRows() Hàm này sử dụng trong công đoạn giải mã, là hàm

ngược của hàm ShiftRows()

Inv SubBytes() Hàm này sử dụng trong công đoạn giải mã, là hàm

ngược của hàm SubBytes()

3.2.1 Công đoạn mã hóa

Khởi tạo, bản rõ được sao chép vào mảng trạng thái sử dụng các quy ước như

trên. Mảng trạng thái khởi tạo được biến đổi bằng cách thực hiện một hàm vòng Nr

lần (10, 12 hoặc 14 phụ thuộc vào độ dài của khóa) sau khi cộng với khóa RoundKey,

trong đó lần cuối cùng thực hiện khác với các lần trước đó. Sau lần lặp cuối cùng

trạng thái cuối sẽ được chuyển thành output của thuật toán. Bằng cách sử dụng một

dãy các khóa được biểu diễn như là mảng một chiều của các word 4 byte được sinh

ra từ thử tục sinh khóa hàm vòng được tham số hóa.

Tất cả các vòng đều thực hiện công việc giống nhau dựa trên 4 hàm theo thứ tự

SubBytes(), ShiftRows(), MixColumns(), và AddRoundKey() ngoại trừ vòng cuối

cùng không thực hiện hàm MixColumns().[4].

Hàm SubBytes()

Thực hiện phép thay thế các byte của mảng trạng thái bằng cách sử dụng một

bảng thế khả nghịch S-box, xây dựng bằng cách kết hợp hai biến đổi sau:

1. Nhân nghịch đảo trên trường hữu hạn GF(28), phần tử {00} được ánh xạ thành

chính nó

2. Áp dụng biến đổi Affine sau (trên GF(2)):

51

bi’= bi b(i+4)mod8  b(i+5)mod8 b(i+6)mod8 b(i+7)mod8 ci, trong đó 0 i 8 là bit

thứ i của byte b tương ứng và ci là bit thứ I của byte c với giá trị {63} hay {01100011}

Dưới đây là hình minh họa kết quả áp dụng hàm biến đổi SubBytes() đối với

mảng trạng thái

Hình 20: Biến đổi SubBytes() đối với mảng trạng thái

Hàm ShiftRows()

Thực hiện bằng cách các byte trong ba hàng cuối của mảng trạng thái sẽ được dịch

vòng với số lần dịch khác nhau. Hàng đầu tiên r=0 không bị dịch như sau:

S’rc = Sr,(c+shift(r,Nb))modNb (Nb=4) trong đó giá trị dịch shift(r,Nb) phụ thuộc

vào số hàng r như sau:

shift(1,4)=1, shift(2,4)=2, shift(3,4)=3

Các byte thấp nhất sẽ được chuyển lên đầu hàng, trong khi các byte khác sẽ tới các

vị trí thấp hơn trong các hàng. Có thể xem minh họa như hình dưới:

Hình 21: Mô tả Hàm ShiftRows()

52

Hàm MixColumns()

Thực hiện trên các cột của mảng trạng thái, coi mỗi cột của mảng trạng thái như là một đa thức gồm 4 hạng tử. Các cột sẽ được xem như là các đa thức trên GF(28) và được nhân trên modulo x4+1 với một đa thức cố định a(x):

a(x) = {03}x3 + {02}x2 + {01}x + {02}

Có thể biểu diễn bằng phép nhân ma trận:

s’(x) = a(x)  s(x)

=

Với mọi 0≤c

Mỗi cột sẽ có bốn byte sẽ được thay thế theo công thức sau:

s’0,c = ({02}●s0,c)({03}●s1,c) s2,c s3,c

s’1,c = s0,c({02}●s1,c) ({03}●s2,c) s3,c

s’2,c = s0,cs1,c ({02}●s2,c) ({03}●s3,c)

s’3,c = ({03}●s0,c)s1,c s2,c ({02}●s3,c)

Dưới đây là hình minh họa:

Hình 22: Mô tả hàm MixColumns()

53

Hàm AddRoundKey()

Thực hiện bằng cách một khóa vòng sẽ được cộng vào mảng trạng thái bằng phép

toán XOR bit. Sinh ra bởi thủ tục sinh khóa, mỗi khóa vòng gồm Nb word. Sau đó

các word này sẽ được cộng vào mỗi cột của mảng trạng thái:

[s’0,c, s’1,c, s’2,c, s’3,c] = [s0,c, s1,c, s2,c, s3,c]  [wround*Nb+c] 0≤c≤Nb = 4

Round là lần lặp tương ứng với quy ước 0≤round≤Nb, trong đó [wi] là các word

của khóa. Trước khi các vòng lặp của thuật toán được thực hiện thuật toán mã hóa

phép cộng khóa vòng khởi tạo xảy ra với round=0. Hàm này được thực hiện trong

thuật toán mã hóa khi 1≤round≤Nb.

Dưới đây là minh họa, trong đó l=round*Nb: 

Hình 23: Mô tả hàm AddRoundKey()

Thuật toán sinh khóa (Key Expansion)

Để sinh một dãy các khóa cho việc mã hóa, thuật toán này nhận một khóa mã

hóa K sau đó thực hiện một thủ tục sinh khóa. Thủ tục sử dụng một tập khởi tạo Nb

word và mỗi lần lặp trong số Nr sẽ cần tới Nb word của dữ liệu khóa và sẽ sinh tổng

số Nb*(Nr+1) word. Kết quả là một mảng tuyến tính các word 4 byte được ký hiệu

là [wi] trong đó 0≤i

Áp dụng bảng thế S-box lên input để nhận được một word output, SubWord()

là hàm nhận một input 4 byte. Hàm RotWord() nhận một word input [a0, a1, a2, a3]

thực hiện một hoán vị vòng và trả về [a1, a2, a3, a0]. Các phần tử của mảng hằng số

54

Rcon[i] chứa các giá trị nhận được bởi [xi-1, {00}, {00}, {00}] trong đó xi-1 là mũ hóa

của x (x được biểu diễn dưới dạng {02} trên GF(28) và i bắt đầu từ 1)

Nk word của khóa kết quả sẽ được điền bởi khóa mã hóa. Các word sau đó w[i]

sẽ bằng XOR với word đứng trước nó w[i-1] và w[i-Nk]. Trước khi thực hiện phép

XOR bit với các word ở vị trí chia hết cho Nk một biến đổi sẽ được thực hiện với

w[i-1], sau đó là phép XOR với một hằng số Rcon[i]. Biến đổi này gồm một phép

dịch vòng các byte của một word sau đó là áp dụng một bảng tra lên tất cả 4 byte của

word. So với thủ tục cho các khóa có độ dài 128 hoặc 192 bit thủ tục mở rộng khóa

đối với các khóa có độ dài 256 bit hơi khác. Trước khi thực hiện phép XOR bit nếu

Nk=8 và i-4 là bội số của Nk thì SubWord() sẽ được áp dụng cho w[i-1].

3.2.2 Công đoạn giải mã

Giống với công đoạn mã hóa nhưng các hàm trong công đoạn giải mã là các

hàm ngược của các hàm ở công đoạn mã hóa.[4].

Hàm InvShiftRow()

Là hàm ngược của hàm ShiftRows(), các byte của ba hàng cuối của mảng trạng

thái sẽ được dịch vòng với vị trí dịch khác nhau. Ba hàng cuối bị dịch đi Nb-

shift(r,Nb) byte trong đó giá trị shift(r,Nb) phụ thuộc vào số hàng, hàng đầu tiên

không bị dịch.

Hàm này thực hiện như sau:

s’r,(c+shift(r,Nb))modNb = sr,c 0

Dưới đây là hình ảnh minh họa:

55

Hình 24: Mô tả hàm InvShiftRow()

Hàm InvSubBytes() Là hàm ngược của hàm SubBytes(). Bằng cách thực hiện nhân nghịch đảo trên

GF(28), hàm sử dụng nghịch đảo của biến đổi Affine.

Hàm InvMixColumns()

Là hàm ngược của hàm MixColumns(). Coi mỗi cột như một đa thức 4 hạng tử

hàm này thực hiện làm việc trên các cột của mảng trạng thái. Các cột được xem như

là các đa thức trên GF(28) và được nhân theo modulo x4+1 với một đa thức cố định

là a-1(x)

a-1(x) = {0b}x3 + {0d}x2 + {09}x + {0e}

Được mô tả bằng phép nhân ma trận sau:

s’(x) = a-1(x) s(x)

=

trong đó 0≤c

Mỗi cột sẽ có bốn byte được thay theo công thức sau:

s’0,c = ({0e}●s0,c)({0b}●s1,c) ({0d}●s2,c) ({09}●s3,c)

56

s’1,c = ({09}●s0,c)({0e}●s1,c) ({0b}●s2,c) ({0d}●s3,c)

s’2,c = ({0d}●s0,c)({09}●s1,c) ({0e}●s2,c) ({0b}●s3,c)

s’3,c = ({0b}●s0,c)({0d}●s1,c) ({09}●s2,c) ({0e}●s3,c)

Hàm nghịch đảo của hàm AddRoundKey()

Hàm nghịch của hàm AddRoundKey() cũng chính là nó vì hàm này chỉ có

phép toán XOR bit

3.3 Chương trình thuật toán song song mã hóa AES sử dụng GPU

Chương trình song song sử dụng GPU:

File thuật toán

#include "aes.h"

// state - array holding the intermediate results during decryption.

typedef uint8_t state_t[4][4];

// The array that stores the round keys.

//__device__ static const uint8_t* RoundKey;

__device__ uintmax_t get_global_index(void)

{

return blockIdx.x * blockDim.x + threadIdx.x;

}

// prints string as hex

__device__ static void phex(uint8_t* str) {

unsigned char i;

for (i = 0; i < 16; ++i)

printf("%.2x", str[i]);

printf("\n");

}

__device__ static void print_state(state_t* state, char message[]) {

57

uintmax_t idx = get_global_index();

uint8_t i, j;

//for (i = 0; i < 4; i++)

printf("[thread %lld] state %s\n%.2x %.2x %.2x %.2x\n%.2x %.2x %.2x %.2x\n%.2x %.2x %.2x %.2x\n%.2x %.2x %.2x %.2x\n", idx, message,

(*state)[0][0], (*state)[0][1], (*state)[0][2], (*state)[0][3],

(*state)[1][0], (*state)[1][1], (*state)[1][2], (*state)[1][3],

(*state)[2][0], (*state)[2][1], (*state)[2][2], (*state)[2][3],

(*state)[3][0], (*state)[3][1], (*state)[3][2], (*state)[3][3]);

}

//

//__device__ static void printKey() {

// printf("RoundKey:\n");

// unsigned char i, j;

// for (j = 0; j < ROUNDS + 1; ++j) {

//

for (i = 0; i < KEYLENGTH; ++i)

//

printf("%.2x", RoundKey[(j*KEYLENGTH) + i]);

//

printf("\n");

// }

//}

// Lookup-tables

__device__ __constant__ uint8_t d_sbox[256] = {

//0 1 2 3 4 5 6 7 8 9 A B C D E F

0x63, 0x7c, 0x77, 0x7b, 0xf2, 0x6b, 0x6f, 0xc5, 0x30, 0x01, 0x67, 0x2b, 0xfe, 0xd7, 0xab, 0x76,

0xca, 0x82, 0xc9, 0x7d, 0xfa, 0x59, 0x47, 0xf0, 0xad, 0xd4, 0xa2, 0xaf, 0x9c, 0xa4, 0x72, 0xc0,

0xb7, 0xfd, 0x93, 0x26, 0x36, 0x3f, 0xf7, 0xcc, 0x34, 0xa5, 0xe5, 0xf1, 0x71, 0xd8, 0x31, 0x15,

0x04, 0xc7, 0x23, 0xc3, 0x18, 0x96, 0x05, 0x9a, 0x07, 0x12, 0x80, 0xe2, 0xeb, 0x27, 0xb2, 0x75,

0x09, 0x83, 0x2c, 0x1a, 0x1b, 0x6e, 0x5a, 0xa0, 0x52, 0x3b, 0xd6, 0xb3, 0x29, 0xe3, 0x2f, 0x84,

58

0x53, 0xd1, 0x00, 0xed, 0x20, 0xfc, 0xb1, 0x5b, 0x6a, 0xcb, 0xbe, 0x39, 0x4a, 0x4c, 0x58, 0xcf,

0xd0, 0xef, 0xaa, 0xfb, 0x43, 0x4d, 0x33, 0x85, 0x45, 0xf9, 0x02, 0x7f, 0x50, 0x3c, 0x9f, 0xa8,

0x51, 0xa3, 0x40, 0x8f, 0x92, 0x9d, 0x38, 0xf5, 0xbc, 0xb6, 0xda, 0x21, 0x10, 0xff, 0xf3, 0xd2,

0xcd, 0x0c, 0x13, 0xec, 0x5f, 0x97, 0x44, 0x17, 0xc4, 0xa7, 0x7e, 0x3d, 0x64, 0x5d, 0x19, 0x73,

0x60, 0x81, 0x4f, 0xdc, 0x22, 0x2a, 0x90, 0x88, 0x46, 0xee, 0xb8, 0x14, 0xde, 0x5e, 0x0b, 0xdb,

0xe0, 0x32, 0x3a, 0x0a, 0x49, 0x06, 0x24, 0x5c, 0xc2, 0xd3, 0xac, 0x62, 0x91, 0x95, 0xe4, 0x79,

0xe7, 0xc8, 0x37, 0x6d, 0x8d, 0xd5, 0x4e, 0xa9, 0x6c, 0x56, 0xf4, 0xea, 0x65, 0x7a, 0xae, 0x08,

0xba, 0x78, 0x25, 0x2e, 0x1c, 0xa6, 0xb4, 0xc6, 0xe8, 0xdd, 0x74, 0x1f, 0x4b, 0xbd, 0x8b, 0x8a,

0x70, 0x3e, 0xb5, 0x66, 0x48, 0x03, 0xf6, 0x0e, 0x61, 0x35, 0x57, 0xb9, 0x86, 0xc1, 0x1d, 0x9e,

0xe1, 0xf8, 0x98, 0x11, 0x69, 0xd9, 0x8e, 0x94, 0x9b, 0x1e, 0x87, 0xe9, 0xce, 0x55, 0x28, 0xdf,

0x8c, 0xa1, 0x89, 0x0d, 0xbf, 0xe6, 0x42, 0x68, 0x41, 0x99, 0x2d, 0x0f, 0xb0, 0x54, 0xbb, 0x16 };

// XOR the round key on state.

__device__ void AddRoundKey(state_t* state, uint8_t* roundKey, uint8_t round) {

//uintmax_t idx = get_global_index();

%lld]

roundKey:

//printf("[Thread %.2x%.2x%.2x%.2x%.2x%.2x%.2x%.2x%.2x%.2x%.2x%.2x%.2x%.2x%.2x%.2x\n", idx,

roundKey[round*BLOCKSIZE

roundKey[round*BLOCKSIZE

0],

+

+

1],

// roundKey[round*BLOCKSIZE + 2], roundKey[round*BLOCKSIZE + 3],

roundKey[round*BLOCKSIZE

roundKey[round*BLOCKSIZE

4],

+

+

5],

// roundKey[round*BLOCKSIZE + 6], roundKey[round*BLOCKSIZE + 7],

roundKey[round*BLOCKSIZE

roundKey[round*BLOCKSIZE

8],

+

+

9],

// roundKey[round*BLOCKSIZE + 10], roundKey[round*BLOCKSIZE + 11],

roundKey[round*BLOCKSIZE

roundKey[round*BLOCKSIZE

12],

+

+

13],

// roundKey[round*BLOCKSIZE + 14], roundKey[round*BLOCKSIZE + 15]);

uint8_t i, j;

for (i = 0; i<4; ++i) {

for (j = 0; j < 4; ++j) {

//printf("[Thread %lld] (*state)[%d][%d] before: %.2x\n", idx, i,

j, (*state)[i][j]);

59

(*state)[i][j] ^= roundKey[round * LANESIZE * 4 + i * LANESIZE +

j];

//printf("[Thread %lld] (*state)[%d][%d] after: %.2x\n", idx, i, j,

(*state)[i][j]);

}

}

}

// The SubBytes Function Substitutes the values in the

// state matrix with values in an S-box.

__device__ void SubBytes(state_t* state, uint8_t* s_sbox)

{

uint8_t i, j;

for (i = 0; i < 4; ++i)

{

for (j = 0; j < 4; ++j)

{

(*state)[j][i] = s_sbox[(*state)[j][i]];

}

}

}

// The ShiftRows() function shifts the rows in the state to the left.

// Each row is shifted with different offset.

// Offset = Row number. So the first row is not shifted.

__device__ void ShiftRows(state_t* state)

{

uint8_t temp;

// Rotate first row 1 columns to left

temp = (*state)[0][1];

(*state)[0][1] = (*state)[1][1];

(*state)[1][1] = (*state)[2][1];

(*state)[2][1] = (*state)[3][1];

(*state)[3][1] = temp;

// Rotate second row 2 columns to left

60

temp = (*state)[0][2];

(*state)[0][2] = (*state)[2][2];

(*state)[2][2] = temp;

temp = (*state)[1][2];

(*state)[1][2] = (*state)[3][2];

(*state)[3][2] = temp;

// Rotate third row 3 columns to left

temp = (*state)[0][3];

(*state)[0][3] = (*state)[3][3];

(*state)[3][3] = (*state)[2][3];

(*state)[2][3] = (*state)[1][3];

(*state)[1][3] = temp;

}

__device__ uint8_t xtime(uint8_t x)

{

return ((x << 1) ^ (((x >> 7) & 1) * 0x1b));

}

// MixColumns function mixes the columns of the state matrix

__device__ void MixColumns(state_t* state)

{

uint8_t i;

uint8_t Tmp, Tm, t;

for (i = 0; i < 4; ++i)

{

t = (*state)[i][0];

Tmp = (*state)[i][0] ^ (*state)[i][1] ^ (*state)[i][2] ^ (*state)[i][3];

Tm = (*state)[i][0] ^ (*state)[i][1]; Tm = xtime(Tm); (*state)[i][0] ^=

Tm ^ Tmp;

Tm = (*state)[i][1] ^ (*state)[i][2]; Tm = xtime(Tm); (*state)[i][1] ^=

Tm ^ Tmp;

Tm = (*state)[i][2] ^ (*state)[i][3]; Tm = xtime(Tm); (*state)[i][2] ^=

Tm ^ Tmp;

Tm = (*state)[i][3] ^ t; Tm = xtime(Tm); (*state)[i][3] ^= Tm ^

Tmp;

61

}

}

// Cipher is the main function that encrypts the PlainText.

__device__ void Cipher(state_t* state, uint8_t* roundKey, uint8_t* s_sbox)

{

uint8_t round = 0;

// Add the First round key to the state before starting the rounds.

AddRoundKey(state, roundKey, round);

//print_state(state, "after first round key added");

// There will be ROUNDS rounds.

// The first ROUNDS-1 rounds are identical.

// These ROUNDS-1 rounds are executed in the loop below.

for (round = 1; round < ROUNDS; ++round)

{

SubBytes(state, s_sbox);

ShiftRows(state);

MixColumns(state);

AddRoundKey(state, roundKey, round);

//print_state(state, "after round key added");

}

// The last round is given below.

// The MixColumns function is not here in the last round.

SubBytes(state, s_sbox);

ShiftRows(state);

AddRoundKey(state, roundKey, ROUNDS);

//print_state(state, "after last round key added");

}

__device__ void AES128_ECB_encrypt(uint8_t* ciphertext_block, uint8_t* roundKey, uint8_t* s_sbox) {

62

state_t* state = (state_t*)ciphertext_block;

//print_state(state, "after init");

// The next function call encrypts the PlainText with the Key using AES algorithm.

Cipher(state, roundKey, s_sbox);

}

__global__ void cuda_encrypt_block(uint8_t* d_ciphertext, uint8_t* d_plaintext, uint8_t* d_roundKey, uintmax_t plaintext_blocks) {

uintmax_t idx = blockIdx.x * blockDim.x + threadIdx.x;

__shared__ uint8_t s_roundKey[BLOCKSIZE * (ROUNDS + 1)];

//__shared__ uint8_t s_ciphertext[BLOCKSIZE * THREADS_PER_BLOCK];

__shared__ uint8_t s_sbox[256];

uintmax_t offset = idx*BLOCKSIZE;

uintmax_t block_offset = (idx % THREADS_PER_BLOCK) * BLOCKSIZE;

// if there are enough THREADS_PER_BLOCK, the round key allocation to shared memory is performed by (ROUNDS + 1) threads in parallel

if (THREADS_PER_BLOCK >= (ROUNDS + 1) && (idx % THREADS_PER_BLOCK) < (ROUNDS + 1)) {

memcpy(s_roundKey + block_offset, d_roundKey + block_offset, BLOCKSIZE);

}

// if not, this is done only by the first thread in a block

else if ((idx % THREADS_PER_BLOCK) == 0) {

memcpy(s_roundKey, d_roundKey, BLOCKSIZE*(ROUNDS + 1));

}

// first thread in a block copies sbox from constant to shared memory

if ((idx % THREADS_PER_BLOCK) == 0) {

memcpy(s_sbox, d_sbox, sizeof(uint8_t) * 256);

}

__syncthreads();

if (idx < plaintext_blocks) {

//memcpy(s_ciphertext + block_offset, d_plaintext + offset, BLOCKSIZE);

63

memcpy(d_ciphertext + offset, d_plaintext + offset, BLOCKSIZE);

// each plaintext block is encrypted by an individual thread

AES128_ECB_encrypt(d_ciphertext + block_offset, s_roundKey, s_sbox);

//memcpy(d_ciphertext

+

offset,

s_ciphertext

+

block_offset,

sizeof(uint8_t)*BLOCKSIZE);

}

}

File chạy

#define DEBUG 0

#include "aes.h"

#include

static double encrypt_file(char* outfile, char* infile, uint8_t* key);

static void __host__ phex(uint8_t* str);

uint8_t key[16] = { (uint8_t)0x2b, (uint8_t)0x7e, (uint8_t)0x15, (uint8_t)0x16,

(uint8_t)0x28, (uint8_t)0xae, (uint8_t)0xd2, (uint8_t)0xa6,

(uint8_t)0xab, (uint8_t)0xf7, (uint8_t)0x15, (uint8_t)0x88,

(uint8_t)0x09, (uint8_t)0xcf, (uint8_t)0x4f, (uint8_t)0x3c };

// The array that stores the round keys.

uint8_t h_roundKey[176];

boolean silent = 0;

void print_usage() {

printf("Usage: aes_parallel.exe

[--silent]\n");

return;

}

int main(int argc, char *argv[]) {

64

if (argc < 3 || argc > 4) {

print_usage();

return 1;

}

double cpu_time_used;

if (argc == 4)

if (!strcmp(argv[3], "--silent"))

silent = 1;

cpu_time_used = encrypt_file(argv[1], argv[2], key);

printf("Execution time: %6.9f seconds\n", cpu_time_used);

printf("Press enter to continue...\n");

getchar();

return 0;

}

double encrypt_file(char* infile, char* outfile, uint8_t* key) {

FILE *fp_in;

FILE *fp_out;

#if defined(DEBUG) && DEBUG

uint8_t i;

#endif

fp_in = fopen(infile, "rb");

if (fp_in == NULL && !silent) {

fprintf(stderr, "Can't open input file %s!\n", infile);

exit(1);

}

fp_out = fopen(outfile, "wb+");

if (fp_out == NULL && !silent) {

fprintf(stderr, "Can't open output file %s!\n", outfile);

65

exit(1);

}

KeyExpansion(key);

#if defined(DEBUG) && DEBUG

printf("Round Keys:\n");

for (i = 0; i < ROUNDS + 1; i++) {

phex(h_roundKey + (i * BLOCKSIZE));

}

#endif

// determine size of file, read file into h_plaintext and determine number of plaintext blocks

fseek(fp_in, 0, SEEK_END);

uintmax_t plaintext_size = ftell(fp_in);

rewind(fp_in);

uint8_t* h_plaintext = (uint8_t*)malloc(plaintext_size);

uintmax_t bytes_read = fread(h_plaintext, sizeof(uint8_t), plaintext_size, fp_in);

assert(bytes_read == plaintext_size);

uintmax_t plaintext_blocks = (bytes_read + BLOCKSIZE - 1) / BLOCKSIZE;

uint8_t* h_ciphertext = (uint8_t*)malloc(plaintext_blocks*BLOCKSIZE);

if (!silent) {

printf("File size: %llu bytes\n", plaintext_size);

printf("Number of plaintext blocks: %llu (blocksize: %d bytes)\n",

plaintext_blocks, BLOCKSIZE);

}

#if defined(DEBUG) && DEBUG

printf("Plaintext:\n");

for (i = 0; i < plaintext_blocks; i++) {

phex(h_plaintext + (i * BLOCKSIZE));

}

#endif

66

cudaError_t cudaStatus;

uintmax_t threads_per_block = THREADS_PER_BLOCK;

uintmax_t number_of_blocks = (plaintext_blocks + threads_per_block - 1) / threads_per_block;

uintmax_t shared_memory_size = BLOCKSIZE * THREADS_PER_BLOCK + BLOCKSIZE * (ROUNDS + 1) + 256;

if (!silent) {

printf("Launching kernel with configuration:\n");

printf("Threads per block: %lld\n", threads_per_block);

printf("Number of blocks: %lld\n", number_of_blocks);

printf("Shared memory size (per block): %lld\n", shared_memory_size);

}

// measure time

double cpu_time_used;

LARGE_INTEGER frequency;

LARGE_INTEGER start, end;

QueryPerformanceFrequency(&frequency);

// start timer

QueryPerformanceCounter(&start);

// copy h_plaintext and h_roundKey into global device memory

uint8_t* d_plaintext;

cudaStatus = cudaMalloc((void**)&d_plaintext, sizeof(uint8_t) * (plaintext_blocks * BLOCKSIZE)); // TODO if last block is smaller than BLOCKSIZE, the block maybe needs to be initialized with zero bits, test if this has to be done

if (cudaStatus != cudaSuccess && !silent) {

fprintf(stderr, "cudaMalloc failed!");

goto Error;

}

// make sure the last block is padded with zero bytes by initializing the full array with zero bytes

cudaStatus = cudaMemset(d_plaintext, 0, sizeof(uint8_t) * (plaintext_blocks * BLOCKSIZE));

67

if (cudaStatus != cudaSuccess && !silent) {

fprintf(stderr, "cudaMemset failed!");

goto Error;

}

cudaMemcpy(d_plaintext,

=

h_plaintext,

cudaStatus sizeof(uint8_t)*plaintext_size, cudaMemcpyHostToDevice);

if (cudaStatus != cudaSuccess && !silent) {

fprintf(stderr, "cudaMemcpy failed!");

goto Error;

}

uint8_t* d_roundKey;

cudaMalloc((void**)&d_roundKey, sizeof(uint8_t)*BLOCKSIZE*(ROUNDS+1));

if (cudaStatus != cudaSuccess && !silent) {

fprintf(stderr, "cudaMalloc failed!");

goto Error;

}

cudaMemcpy(d_roundKey, h_roundKey, sizeof(uint8_t)*BLOCKSIZE*(ROUNDS + 1), cudaMemcpyHostToDevice);

if (cudaStatus != cudaSuccess && !silent) {

fprintf(stderr, "cudaMemcpy failed!");

goto Error;

}

// allocate space for the ciphertext on the device

uint8_t* d_ciphertext;

=

cudaMalloc((void**)&d_ciphertext,

sizeof(uint8_t)

*

cudaStatus (plaintext_blocks * BLOCKSIZE));

if (cudaStatus != cudaSuccess && !silent) {

fprintf(stderr, "cudaMalloc failed!");

goto Error;

}

// reset last error

cudaGetLastError();

68

cuda_encrypt_block<<>>(d_ciphertext, d_plaintext, d_roundKey, plaintext_blocks);

cudaStatus = cudaGetLastError();

if (cudaStatus != cudaSuccess && !silent) {

"Kernel

launch

failed:

%s\n",

fprintf(stderr, cudaGetErrorString(cudaStatus));

goto Error;

}

cudaStatus = cudaDeviceSynchronize();

if (cudaStatus != cudaSuccess && !silent) {

"cudaDeviceSynchronize

failed:

%s\n",

fprintf(stderr, cudaGetErrorString(cudaStatus));

goto Error;

}

// Copy ciphertext array from device memory to host memory.

cudaMemcpy(h_ciphertext,

d_ciphertext,

=

sizeof(uint8_t)

*

cudaStatus (plaintext_blocks * BLOCKSIZE), cudaMemcpyDeviceToHost);

if (cudaStatus != cudaSuccess && !silent) {

fprintf(stderr, "cudaMemcpy failed!");

goto Error;

}

// stop timer

QueryPerformanceCounter(&end);

=

((double)(end.QuadPart

-

start.QuadPart))

/

cpu_time_used ((double)frequency.QuadPart);

#if defined(DEBUG) && DEBUG

printf("Ciphertext after kernel returned:\n");

for (i = 0; i < plaintext_blocks; i++) {

phex(h_ciphertext + (i * BLOCKSIZE));

}

#endif

69

// write ciphertext to output file

fwrite(h_ciphertext, sizeof(uint8_t), BLOCKSIZE * plaintext_blocks, fp_out);

if (!silent)

printf("\nEncryption

of

%llu

plaintext

blocks

successful!\n",

plaintext_blocks);

return cpu_time_used;

Error:

free(h_plaintext);

free(h_ciphertext);

free(h_roundKey);

cudaFree(d_plaintext);

cudaFree(d_ciphertext);

cudaFree(d_roundKey);

fclose(fp_in);

fclose(fp_out);

exit(1);

}

// This function produces (ROUNDS+1) round keys. The round keys are used in each round to decrypt the states.

void KeyExpansion(uint8_t* key) {

uint32_t i, j, k;

uint8_t tempa[4]; // Used for the column/row operations

// The first round key is the key

for (i = 0; i < KEYWORDS; ++i)

{

h_roundKey[(i * 4) + 0] = key[(i * 4) + 0];

h_roundKey[(i * 4) + 1] = key[(i * 4) + 1];

h_roundKey[(i * 4) + 2] = key[(i * 4) + 2];

70

h_roundKey[(i * 4) + 3] = key[(i * 4) + 3];

}

// All other round keys are found from the previous round keys.

for (; (i < (LANESIZE * (ROUNDS + 1))); ++i)

{

for (j = 0; j < 4; ++j)

{

tempa[j] = h_roundKey[(i - 1) * 4 + j];

}

if (i % KEYWORDS == 0)

{

// This function rotates the 4 bytes in a word to the left once.

// [a0,a1,a2,a3] becomes [a1,a2,a3,a0]

// Function RotWord()

{

k = tempa[0];

tempa[0] = tempa[1];

tempa[1] = tempa[2];

tempa[2] = tempa[3];

tempa[3] = k;

}

// SubWord() is a function that takes a four-byte input word and

// applies the S-box to each of the four bytes to produce an output

word.

// Function Subword()

{

tempa[0] = sbox[tempa[0]];

tempa[1] = sbox[tempa[1]];

tempa[2] = sbox[tempa[2]];

tempa[3] = sbox[tempa[3]];

}

71

tempa[0] = tempa[0] ^ Rcon[i / KEYWORDS];

}

else if (KEYWORDS > 6 && i % KEYWORDS == 4)

{

// Function Subword()

{

tempa[0] = sbox[tempa[0]];

tempa[1] = sbox[tempa[1]];

tempa[2] = sbox[tempa[2]];

tempa[3] = sbox[tempa[3]];

}

}

h_roundKey[i * 4 + 0] = h_roundKey[(i - KEYWORDS) * 4 + 0] ^ tempa[0];

h_roundKey[i * 4 + 1] = h_roundKey[(i - KEYWORDS) * 4 + 1] ^ tempa[1];

h_roundKey[i * 4 + 2] = h_roundKey[(i - KEYWORDS) * 4 + 2] ^ tempa[2];

h_roundKey[i * 4 + 3] = h_roundKey[(i - KEYWORDS) * 4 + 3] ^ tempa[3];

}

}

// prints string as hex

static void phex(uint8_t* str) {

unsigned char i;

for (i = 0; i < 16; ++i)

printf("%.2x", str[i]);

printf("\n");

}

Chương trình sử dụng CPU:

File thuật toán

#include

#include // CBC mode, for memset

#include "aes.h"

// state - array holding the intermediate results during decryption.

72

typedef uint8_t state_t[4][4];

static state_t* state;

// The array that stores the round keys.

static const uint8_t* RoundKey;

// prints string as hex

static void phex(uint8_t* str) {

unsigned char i;

for (i = 0; i < 16; ++i)

printf("%.2x", str[i]);

printf("\n");

}

static void print_state() {

uint8_t i, j;

printf("state:\n");

for (i = 0; i < 4; i++) {

for (j = 0; j < 4; j++)

printf("%.2x", (*state)[i][j]);

printf("\n");

}

}

static void printKey() {

printf("RoundKey:\n");

unsigned char i, j;

for (j = 0; j < ROUNDS + 1; ++j) {

for (i = 0; i < KEYLENGTH; ++i)

printf("%.2x", RoundKey[(j*KEYLENGTH)+i]);

printf("\n");

}

}

// The round constant word array, Rcon[i], contains the values given by

73

// x to th power (i-1) being powers of x (x is denoted as {02}) in the field GF(2^8)

// Note that i starts at 1, not 0).

static const uint8_t Rcon[255] = {

0x8d, 0x01, 0x02, 0x04, 0x08, 0x10, 0x20, 0x40, 0x80, 0x1b, 0x36, 0x6c, 0xd8, 0xab, 0x4d, 0x9a,

0x2f, 0x5e, 0xbc, 0x63, 0xc6, 0x97, 0x35, 0x6a, 0xd4, 0xb3, 0x7d, 0xfa, 0xef, 0xc5, 0x91, 0x39,

0x72, 0xe4, 0xd3, 0xbd, 0x61, 0xc2, 0x9f, 0x25, 0x4a, 0x94, 0x33, 0x66, 0xcc, 0x83, 0x1d, 0x3a,

0x74, 0xe8, 0xcb, 0x8d, 0x01, 0x02, 0x04, 0x08, 0x10, 0x20, 0x40, 0x80, 0x1b, 0x36, 0x6c, 0xd8,

0xab, 0x4d, 0x9a, 0x2f, 0x5e, 0xbc, 0x63, 0xc6, 0x97, 0x35, 0x6a, 0xd4, 0xb3, 0x7d, 0xfa, 0xef,

0xc5, 0x91, 0x39, 0x72, 0xe4, 0xd3, 0xbd, 0x61, 0xc2, 0x9f, 0x25, 0x4a, 0x94, 0x33, 0x66, 0xcc,

0x83, 0x1d, 0x3a, 0x74, 0xe8, 0xcb, 0x8d, 0x01, 0x02, 0x04, 0x08, 0x10, 0x20, 0x40, 0x80, 0x1b,

0x36, 0x6c, 0xd8, 0xab, 0x4d, 0x9a, 0x2f, 0x5e, 0xbc, 0x63, 0xc6, 0x97, 0x35, 0x6a, 0xd4, 0xb3,

0x7d, 0xfa, 0xef, 0xc5, 0x91, 0x39, 0x72, 0xe4, 0xd3, 0xbd, 0x61, 0xc2, 0x9f, 0x25, 0x4a, 0x94,

0x33, 0x66, 0xcc, 0x83, 0x1d, 0x3a, 0x74, 0xe8, 0xcb, 0x8d, 0x01, 0x02, 0x04, 0x08, 0x10, 0x20,

0x40, 0x80, 0x1b, 0x36, 0x6c, 0xd8, 0xab, 0x4d, 0x9a, 0x2f, 0x5e, 0xbc, 0x63, 0xc6, 0x97, 0x35,

0x6a, 0xd4, 0xb3, 0x7d, 0xfa, 0xef, 0xc5, 0x91, 0x39, 0x72, 0xe4, 0xd3, 0xbd, 0x61, 0xc2, 0x9f,

0x25, 0x4a, 0x94, 0x33, 0x66, 0xcc, 0x83, 0x1d, 0x3a, 0x74, 0xe8, 0xcb, 0x8d, 0x01, 0x02, 0x04,

0x08, 0x10, 0x20, 0x40, 0x80, 0x1b, 0x36, 0x6c, 0xd8, 0xab, 0x4d, 0x9a, 0x2f, 0x5e, 0xbc, 0x63,

0xc6, 0x97, 0x35, 0x6a, 0xd4, 0xb3, 0x7d, 0xfa, 0xef, 0xc5, 0x91, 0x39, 0x72, 0xe4, 0xd3, 0xbd,

0x61, 0xc2, 0x9f, 0x25, 0x4a, 0x94, 0x33, 0x66, 0xcc, 0x83, 0x1d, 0x3a, 0x74, 0xe8, 0xcb };

// Lookup-tables

static const uint8_t sbox[256] = {

//0 1 2 3 4 5 6 7 8 9 A B C D E F

0x63, 0x7c, 0x77, 0x7b, 0xf2, 0x6b, 0x6f, 0xc5, 0x30, 0x01, 0x67, 0x2b, 0xfe, 0xd7, 0xab, 0x76,

74

0xca, 0x82, 0xc9, 0x7d, 0xfa, 0x59, 0x47, 0xf0, 0xad, 0xd4, 0xa2, 0xaf, 0x9c, 0xa4, 0x72, 0xc0,

0xb7, 0xfd, 0x93, 0x26, 0x36, 0x3f, 0xf7, 0xcc, 0x34, 0xa5, 0xe5, 0xf1, 0x71, 0xd8, 0x31, 0x15,

0x04, 0xc7, 0x23, 0xc3, 0x18, 0x96, 0x05, 0x9a, 0x07, 0x12, 0x80, 0xe2, 0xeb, 0x27, 0xb2, 0x75,

0x09, 0x83, 0x2c, 0x1a, 0x1b, 0x6e, 0x5a, 0xa0, 0x52, 0x3b, 0xd6, 0xb3, 0x29, 0xe3, 0x2f, 0x84,

0x53, 0xd1, 0x00, 0xed, 0x20, 0xfc, 0xb1, 0x5b, 0x6a, 0xcb, 0xbe, 0x39, 0x4a, 0x4c, 0x58, 0xcf,

0xd0, 0xef, 0xaa, 0xfb, 0x43, 0x4d, 0x33, 0x85, 0x45, 0xf9, 0x02, 0x7f, 0x50, 0x3c, 0x9f, 0xa8,

0x51, 0xa3, 0x40, 0x8f, 0x92, 0x9d, 0x38, 0xf5, 0xbc, 0xb6, 0xda, 0x21, 0x10, 0xff, 0xf3, 0xd2,

0xcd, 0x0c, 0x13, 0xec, 0x5f, 0x97, 0x44, 0x17, 0xc4, 0xa7, 0x7e, 0x3d, 0x64, 0x5d, 0x19, 0x73,

0x60, 0x81, 0x4f, 0xdc, 0x22, 0x2a, 0x90, 0x88, 0x46, 0xee, 0xb8, 0x14, 0xde, 0x5e, 0x0b, 0xdb,

0xe0, 0x32, 0x3a, 0x0a, 0x49, 0x06, 0x24, 0x5c, 0xc2, 0xd3, 0xac, 0x62, 0x91, 0x95, 0xe4, 0x79,

0xe7, 0xc8, 0x37, 0x6d, 0x8d, 0xd5, 0x4e, 0xa9, 0x6c, 0x56, 0xf4, 0xea, 0x65, 0x7a, 0xae, 0x08,

0xba, 0x78, 0x25, 0x2e, 0x1c, 0xa6, 0xb4, 0xc6, 0xe8, 0xdd, 0x74, 0x1f, 0x4b, 0xbd, 0x8b, 0x8a,

0x70, 0x3e, 0xb5, 0x66, 0x48, 0x03, 0xf6, 0x0e, 0x61, 0x35, 0x57, 0xb9, 0x86, 0xc1, 0x1d, 0x9e,

0xe1, 0xf8, 0x98, 0x11, 0x69, 0xd9, 0x8e, 0x94, 0x9b, 0x1e, 0x87, 0xe9, 0xce, 0x55, 0x28, 0xdf,

0x8c, 0xa1, 0x89, 0x0d, 0xbf, 0xe6, 0x42, 0x68, 0x41, 0x99, 0x2d, 0x0f, 0xb0, 0x54, 0xbb, 0x16 };

static uint8_t getSBoxValue(uint8_t num) {

return sbox[num];

}

// This function produces LANESIZE * (ROUNDS+1) round keys. The round keys are used in each round to decrypt the states.

void KeyExpansion(uint8_t* roundKey, uint8_t* key)

{

uint32_t i, j, k;

uint8_t tempa[4]; // Used for the column/row operations

// The first round key is the key

75

for (i = 0; i < KEYWORDS; ++i)

{

roundKey[(i * 4) + 0] = key[(i * 4) + 0];

roundKey[(i * 4) + 1] = key[(i * 4) + 1];

roundKey[(i * 4) + 2] = key[(i * 4) + 2];

roundKey[(i * 4) + 3] = key[(i * 4) + 3];

}

// All other round keys are found from the previous round keys.

for (; (i < (LANESIZE * (ROUNDS + 1))); ++i)

{

for (j = 0; j < 4; ++j)

{

tempa[j] = roundKey[(i - 1) * 4 + j];

}

if (i % KEYWORDS == 0)

{

// This function rotates the 4 bytes in a word to the left once.

// [a0,a1,a2,a3] becomes [a1,a2,a3,a0]

// Function RotWord()

{

k = tempa[0];

tempa[0] = tempa[1];

tempa[1] = tempa[2];

tempa[2] = tempa[3];

tempa[3] = k;

}

// SubWord() is a function that takes a four-byte input word and

// applies the S-box to each of the four bytes to produce an output

word.

// Function Subword()

{

tempa[0] = getSBoxValue(tempa[0]);

76

tempa[1] = getSBoxValue(tempa[1]);

tempa[2] = getSBoxValue(tempa[2]);

tempa[3] = getSBoxValue(tempa[3]);

}

tempa[0] = tempa[0] ^ Rcon[i / KEYWORDS];

}

else if (KEYWORDS > 6 && i % KEYWORDS == 4)

{

// Function Subword()

{

tempa[0] = getSBoxValue(tempa[0]);

tempa[1] = getSBoxValue(tempa[1]);

tempa[2] = getSBoxValue(tempa[2]);

tempa[3] = getSBoxValue(tempa[3]);

}

}

roundKey[i * 4 + 0] = roundKey[(i - KEYWORDS) * 4 + 0] ^ tempa[0];

roundKey[i * 4 + 1] = roundKey[(i - KEYWORDS) * 4 + 1] ^ tempa[1];

roundKey[i * 4 + 2] = roundKey[(i - KEYWORDS) * 4 + 2] ^ tempa[2];

roundKey[i * 4 + 3] = roundKey[(i - KEYWORDS) * 4 + 3] ^ tempa[3];

}

}

// XOR the round key on state.

static void AddRoundKey(uint8_t round) {

uint8_t i, j;

for (i = 0; i<4; ++i) {

for (j = 0; j < 4; ++j) {

(*state)[i][j] ^= RoundKey[round * LANESIZE * 4 + i * LANESIZE +

j];

}

}

}

77

// The SubBytes Function Substitutes the values in the

// state matrix with values in an S-box.

static void SubBytes(void)

{

uint8_t i, j;

for (i = 0; i < 4; ++i)

{

for (j = 0; j < 4; ++j)

{

(*state)[j][i] = getSBoxValue((*state)[j][i]);

}

}

}

// The ShiftRows() function shifts the rows in the state to the left.

// Each row is shifted with different offset.

// Offset = Row number. So the first row is not shifted.

static void ShiftRows(void)

{

uint8_t temp;

// Rotate first row 1 columns to left

temp = (*state)[0][1];

(*state)[0][1] = (*state)[1][1];

(*state)[1][1] = (*state)[2][1];

(*state)[2][1] = (*state)[3][1];

(*state)[3][1] = temp;

// Rotate second row 2 columns to left

temp = (*state)[0][2];

(*state)[0][2] = (*state)[2][2];

(*state)[2][2] = temp;

temp = (*state)[1][2];

(*state)[1][2] = (*state)[3][2];

(*state)[3][2] = temp;

78

// Rotate third row 3 columns to left

temp = (*state)[0][3];

(*state)[0][3] = (*state)[3][3];

(*state)[3][3] = (*state)[2][3];

(*state)[2][3] = (*state)[1][3];

(*state)[1][3] = temp;

}

static uint8_t xtime(uint8_t x)

{

return ((x << 1) ^ (((x >> 7) & 1) * 0x1b));

}

// MixColumns function mixes the columns of the state matrix

static void MixColumns(void)

{

uint8_t i;

uint8_t Tmp, Tm, t;

for (i = 0; i < 4; ++i)

{

t = (*state)[i][0];

Tmp = (*state)[i][0] ^ (*state)[i][1] ^ (*state)[i][2] ^ (*state)[i][3];

Tm = (*state)[i][0] ^ (*state)[i][1]; Tm = xtime(Tm); (*state)[i][0] ^=

Tm ^ Tmp;

Tm = (*state)[i][1] ^ (*state)[i][2]; Tm = xtime(Tm); (*state)[i][1] ^=

Tm ^ Tmp;

Tm = (*state)[i][2] ^ (*state)[i][3]; Tm = xtime(Tm); (*state)[i][2] ^=

Tm ^ Tmp;

Tm = (*state)[i][3] ^ t; Tm = xtime(Tm); (*state)[i][3] ^= Tm ^

Tmp;

}

}

static uint8_t Multiply(uint8_t x, uint8_t y)

{

return (((y & 1) * x) ^

((y >> 1 & 1) * xtime(x)) ^

79

((y >> 2 & 1) * xtime(xtime(x))) ^

((y >> 3 & 1) * xtime(xtime(xtime(x)))) ^

((y >> 4 & 1) * xtime(xtime(xtime(xtime(x))))));

}

// Cipher is the main function that encrypts the PlainText.

static void Cipher(void)

{

uint8_t round = 0;

// Add the First round key to the state before starting the rounds.

AddRoundKey(0);

//print_state();

// There will be ROUNDS rounds.

// The first ROUNDS-1 rounds are identical.

// These ROUNDS-1 rounds are executed in the loop below.

for (round = 1; round < ROUNDS; ++round)

{

SubBytes();

ShiftRows();

MixColumns();

AddRoundKey(round);

}

// The last round is given below.

// The MixColumns function is not here in the last round.

SubBytes();

ShiftRows();

AddRoundKey(ROUNDS);

}

static void BlockCopy(uint8_t* output, uint8_t* input)

{

uint8_t i;

80

for (i = 0; i

{

output[i] = input[i];

}

}

void AES128_ECB_encrypt(uint8_t* input, const uint8_t* roundKey, uint8_t* output)

{

// Copy input to output, and work in-memory on output

BlockCopy(output, input);

state = (state_t*)output;

RoundKey = roundKey;

// The next function call encrypts the PlainText with the Key using AES algorithm.

Cipher();

}

File chạy

#define DEBUG 0

#define _CRT_SECURE_NO_WARNINGS

#define _CRT_SECURE_NO_DEPRECATE

#include

#include

#include

#include

#include

#include "aes.h"

static double encrypt_file(char* outfile, char* infile);

static size_t read_plaintext_block();

81

static void phex(uint8_t* str);

// The array that stores the round keys.

static uint8_t roundKey[176];

// The array that holds the plaintext for the current block.

uint8_t plaintext_block[BLOCKSIZE];

// The array that stores the ciphertext for the current block.

uint8_t ciphertext_block[BLOCKSIZE];

// 128bit key

uint8_t key[16] = { (uint8_t)0x2b, (uint8_t)0x7e, (uint8_t)0x15, (uint8_t)0x16,

(uint8_t)0x28,

(uint8_t)0xae,

(uint8_t)0xd2,

(uint8_t)0xa6,

(uint8_t)0xab,

(uint8_t)0xf7,

(uint8_t)0x15,

(uint8_t)0x88,

boolean silent = 0;

int main(int argc, char *argv[]) {

if (argc < 3 || argc > 4) {

printf("Usage: aes_serial.exe

[--silent]\n",

argv[0]);

return 1;

}

if (argc == 4)

if (!strcmp(argv[3], "--silent"))

silent = 1;

double cpu_time_used;

cpu_time_used = encrypt_file(argv[1], argv[2]);

printf("Execution time: %6.9f seconds\n", cpu_time_used);

return 0;

}

(uint8_t)0x09, (uint8_t)0xcf, (uint8_t)0x4f, (uint8_t)0x3c };

82

double encrypt_file(char* infile, char* outfile) {

FILE *fp_in;

FILE *fp_out;

fp_in = fopen(infile, "rb");

if (fp_in == NULL && !silent) {

fprintf(stderr, "Can't open input file %s!\n", infile);

exit(1);

}

fp_out = fopen(outfile, "wb+");

if (fp_out == NULL && !silent) {

fprintf(stderr, "Can't open output file %s!\n", outfile);

exit(1);

}

KeyExpansion(roundKey, key);

#if defined(DEBUG) && DEBUG

printf("Round Keys:\n");

uint8_t i;

for (i = 0; i < ROUNDS + 1; i++) {

phex(roundKey + (i * ROUNDS));

}

#endif

// determine size of file, read file into plaintext and determine number of plaintext blocks

fseek(fp_in, 0, SEEK_END);

uintmax_t plaintext_size = ftell(fp_in);

rewind(fp_in);

uint8_t* plaintext = (uint8_t*)malloc(plaintext_size);

uintmax_t bytes_read = fread(plaintext, sizeof(uint8_t), plaintext_size, fp_in);

assert(bytes_read == plaintext_size);

uintmax_t plaintext_blocks = (bytes_read + BLOCKSIZE - 1) / BLOCKSIZE;

uint8_t* ciphertext = (uint8_t*)malloc(plaintext_blocks*BLOCKSIZE);

83

if (!silent) {

printf("File size: %llu bytes\n", plaintext_size);

printf("Number of plaintext blocks: %llu (blocksize: %d bytes)\n",

plaintext_blocks, BLOCKSIZE);

}

#if defined(DEBUG) && DEBUG

printf("Plaintext:\n");

for (i = 0; i < plaintext_blocks; i++) {

phex(plaintext + (i * BLOCKSIZE));

}

#endif

// measure time

double cpu_time_used;

LARGE_INTEGER frequency;

LARGE_INTEGER start, end;

QueryPerformanceFrequency(&frequency);

// start timer

QueryPerformanceCounter(&start);

uintmax_t j;

for (j = 0; j < plaintext_blocks; j++) {

// encrypt plaintext block

AES128_ECB_encrypt(plaintext + j*BLOCKSIZE, roundKey, ciphertext_block);

// write ciphertext block to output file

memcpy(ciphertext

+

j*BLOCKSIZE,

ciphertext_block,

sizeof(uint8_t)*BLOCKSIZE);

}

// stop timer

QueryPerformanceCounter(&end);

84

=

((double)(end.QuadPart

-

start.QuadPart))

/

cpu_time_used ((double)frequency.QuadPart);

// write ciphertext to output file

fwrite(ciphertext, sizeof(uint8_t), BLOCKSIZE * plaintext_blocks, fp_out);

#if defined(DEBUG) && DEBUG

printf("Ciphertext:\n");

for (i = 0; i < plaintext_blocks; i++) {

phex(ciphertext + (i * BLOCKSIZE));

}

#endif

fclose(fp_in);

fclose(fp_out);

if (!silent)

printf("Encryption

of

%llu

plaintext

blocks

successful!\n",

plaintext_blocks);

return cpu_time_used;

}

// Reads one block of plaintext of size BLOCKSIZE bytes from the file pointed to by the pointer fp.

// If the last block does not match BLOCKSIZE bytes, the block is padded with zero bytes.

static size_t read_plaintext_block(FILE *fp) {

size_t current_blocksize = fread(plaintext_block, sizeof(uint8_t), BLOCKSIZE, fp);

#if defined(DEBUG) && DEBUG

if (feof(fp))

printf("End-of-File reached.\n");

if (ferror(fp))

printf("An error occurred while accessing the file.\n");

if (current_blocksize == 0) return 0;

//printf("current_blocksize: %d\n", current_blocksize);

85

#endif

if (current_blocksize == 0) return 0;

// pad last block with zeroes if it does not match BLOCKSIZE

if (current_blocksize < BLOCKSIZE) {

uint8_t i;

for (i = 0; current_blocksize + i < BLOCKSIZE; ++i) {

plaintext_block[current_blocksize + i] = '0';

}

}

return current_blocksize;

}

// prints string as hex

static void phex(uint8_t* str) {

unsigned char i;

for(i = 0; i < 16; ++i)

printf("%.2x", str[i]);

printf("\n");

}

86

3.3.1. Giao diện chương trình demo:

Chạy trên CPU:

Chạy trên GPU:

87

3.3.2. Kết quả chương trình và đánh giá hiệu suất tính toán:

Ta dễ thấy GPU có hiệu suất tính toán cao hơn hẳn so với CPU.

Kích cỡ file đầu vào Thời gian chạy trên CPU Thời gian chạy trên GPU

664911 bytes 0.941079536 seconds 0.265264969 seconds

1189270 bytes 1.064967493 seconds 0.188594010 seconds

4337261 bytes 2.624178762 seconds 0.317496639 seconds

8132164 bytes 5.940836913 seconds 0.444080242 seconds

25724901 bytes 14.598369270 seconds 0.981748010 seconds

Hướng phát triển

Với bài toán thử nghiệm trong luận văn cho thấy đối với các dữ liệu đầu vào

dung lượng lớn thì việc chạy trên GPU cho tốc độ tính toán nhanh hơn nhiều lần so

với tiến hành trên CPU. Ngày nay với sự phát triển của các ngành khoa học, và của

các dịch vụ internet thương mại nhu cầu tính toán là rất lớn. Việc phát triển, cài đặt

các thuật toán trên 1 hoặc nhiều GPU là một nhu cầu có thực và cấp thiết. Trong tương

lai tác giả sẽ tiếp tục nghiên cứu phát triển cài đặt các thuật toán, các phương pháp

mã hóa khác trên nền tảng GPU.

88

KẾT LUẬN

Luận văn đã nghiên cứu tổng quan về xử lý song song. Đồng thời luận văn cũng

nghiên cứu về bộ xử lý đồ họa GPU, cùng công cụ lập trình GPU phổ biến hiện nay

là CUDA. Đó là tiền đề để nghiên cứu cách xử lý song song trên bộ xử lý đồ họa

GPU sao cho hiệu quả nhất. Trong chương 3 luận văn cũng trình bày một cách cơ

bản về thuật toán mã hóa tương đối mới hiện nay đó là thuật toán mã hóa AES. Với

những kiến thức đã nghiên cứu và tổng hợp của luận văn. Tác giả đã đưa ra chương

trình thực nghiệm song song xử lý mã hóa AES với GPU thông qua công cụ lập trình

CUDA. Và so sánh với chương trình mã hóa AES bằng chương trình tuần tự xử lý

trên GPU. Trên kết quả nghiên cưu được thông qua bộ dữ liệu đầu vào từ nhỏ và lớn

dần đã cho thấy sự khác biệt và hiệu quả rõ rệt khi sử dụng GPU.

Với các kết quả dựa trên các nghiên cứu tổng quan về công nghệ GPU cũng như

về CUDA và kết quả trên chương trình thực nghiệm. Hy vọng sẽ giúp đưa ra các giải

pháp cho các nhà quản lý và các nhà phát triển phần cứng cũng như phần mềm để

cải thiện tốc độ tính toán sao cho đạt được kết quả tốt nhất.

89

TÀI LIỆU THAM KHẢO

Tài liệu tiếng việt

[1] Trương Văn Hiệu (2011), “Nghiên cứu các giải thuật song song trên hệ thống

xử lý đồ họa GPU đa lõi”, luận văn thạc sĩ, trường Đại học Đà Nẵng.

[2] Nguyễn Việt Đức – Nguyễn Nam Giang (2012), ”Xây dựng thuật toán song song

tìm đường đi ngắn nhất với CUDA”, luận văn thạc sỹ, trường Đại học Công nghệ

Hồ Chí Minh.

[3] Nguyễn Thị Thùy Linh (2009), “Tính toán hiệu năng cao với bộ xử lý đồ họa

GPU và ứng dụng”, luận văn thạc sĩ, trường Đại học Công nghệ Hà Nội.

[4] Nguyễn Thị Dung (2013), “Hệ mã hóa AES”, tiểu luận, trường đại học công

nghệ , Đại học Quốc gia Hà Nội.

[5] Lưu Thị Thúy Linh (2013), “Nghiên cứu một vài khía cạnh về độ an toàn của

AES”, luận văn thạc sĩ, Học viện Công nghệ bưu chính viễn thông.

Tài liệu tiếng anh

[6] Jason Sanders, Edward Kandrot, “CUDA by example”, an introduction to

General- Purpose GPU programming.

[7] Maciej Matyka, “GPGPU programming on example of CUDA”, Institute of

Theoretical Physics University of Wroclaw.

[8] NVIDIA, “High performance computing with CUDA”, Users Group Conference

San Diego, CA June 15, 2009.