Лекция2 - 2015

Представление целых чисел

•в традиционных позиционных системах счисления, арифметические операции не могут быть распараллелены (т.к. перенос распространяется последовательно)

•реализация, в которой одна нить работает с одним «длинным» числом не будет эффективной на CUDA (из-за особенностей доступа к памяти на этой архитектуре)

•следовательно, необходимо изменить способ представления чисел. Новый способ представления должен обеспечивать распараллеливание арифметических операций

•возможные варианты: Многомодульная система счисления и Знакоразрядная система счисления

21

Распараллеливание арифметических операций

Схема организации вычислений

Организация программы

уровень нити скрыт от пользователя

22

Многомодульная система счисления

A = (A0 ,A1 ,K,AN −1 ) AI = A(MOD MI ) M = M0M1 K,MN −1

Отображение положительных чисел:

Отображение отрицательных чисел:

M

M

0,1,K,

→ 0,1,K,

;

2

2

M

M

−

,K,−2,−1

→

,K,M − 2,M −1

2

2

Операция		Результат
A + B	CI	= (AI	+ BI )(MOD MI )
A − B	CI	= (AI	− BI )(MOD MI )
A B	CI	= (AI BI )(MOD MI )

Нет необходимости в обменах между нитями

23

Многомодульная арифметика: исходные коды операций

_device__ inline multi_mod_num operator+ (const multi_mod_num& x, const multi_mod_num& y){

multi_mod_num res; int tmp=x.r+y.r;

if(tmp>=modules[threadIdx.x]) tmp-=modules[threadIdx.x];

res.r=tmp; return res;

}

__device__ inline multi_mod_num operator* (const multi_mod_num& x, const multi_mod_num& y){ multi_mod_num res;

long long tmp=((long long)x.r*(long long)y.r)%(long long)modules[threadIdx.x]; res.r=tmp;

return res;

}

__device__ inline bool operator==(const multi_mod_num& x, const multi_mod_num& y){ __shared__ bool tmpr[MAX_THREADS];

tmpr[threadIdx.x]=(x.r==y.r); __syncthreads();

for (int s=1; s<blockDim.x; s<<=1){

if( (threadIdx.x&((s<<1)-1)) == 0 ) if(threadIdx.x+s<blockDim.x)

tmpr[threadIdx.x]=tmpr[threadIdx.x]&&tmpr[threadIdx.x+s]; __syncthreads();

}

return tmpr[0];

}

24

Параллельная редукция

Дано: массив A из N элементов

бинарная ассоциативная операция

Найти: A0 A1 K AN 1

−

В качестве могут быть использованы: +, min, max

Пример: сложение элементов массива

распараллеливание вычисления выражения: требуется log2(N) шагов

25

Параллельная редукция

Дано: массив A из N элементов

бинарная ассоциативная операция

Найти: A0 A1 K AN 1

−

В качестве могут быть использованы: +, min, max

Пример: работа нитей GPU

	SHARED	GPU THREADS
	MEM

Проверка равенства: проверяем равенство остатков от деления на соответствующие модули, затем проводим редукцию полученного массива относительно операции логическое И

__device__ bool operator==(const multi_mod& x, const multi_mod& y){ __shared__ bool tmpr[MAX_THREADS]; tmpr[threadIdx.x]=(x.r==y.r);

__syncthreads();

for (int s=1; s<blockDim.x; s<<=1){

if( (threadIdx.x&((s<<1)-1)) == 0 ) if(threadIdx.x+s<blockDim.x)

tmpr[threadIdx.x]=tmpr[threadIdx.x]&&tmpr[threadIdx.x+s]; __syncthreads();

									}
0	1	2	3	4	5	6	7		return tmpr[0];
								}	26

Сравнение чисел: многомодульная система счисления

проверка равенства

T[i]= (A[i]==B[i])

Редукция T относительно операции &

сравнение

1)C=A-B

2)Определение знака C

определение знака числа

1)Преобразование в систему счисления со смешанным основанием C->CM

2)Сравнение с (M/2)M: поразрядное вычитание и редукция полученного массива относительно операции Op(a,b)= {a, a!=0; b, a==0}

27

mi−1 (mod m j ),i < j

A = D0 + D1 (M1 ) + D2 (M1M2 ) + K+ D

A0 = D0

Многомодульная система счисления: определение знака числа

Система счисления со смешанным основанием для ρ = [R1 , R2 ,K, RN ]

S = D0 + D1 (R1 ) + D 2 (R1R2 ) +K+ D N−1 (R1R2 KRN−1 )

0 ≤ Di < Ri +1 I = 0,1,K, N −1

1) Перевести число s в смешанную систему счисления со смешанным основанием

•	используем	ρM = [M0 ,M1 ,K,MN −1 ]
•	тогда для	A = (A0 ,A1 ,K,A N −1 )

ρM = [M1 , M2 ,K, MN ]

N −1 (M1M2 KMN −1 )

• поэтому

•исключаем A0 из основания, и повторяем предыдущие

шаги для числа (A − D0 ) M0−1

•следующие значения могут быть просчитаны заранее:

, M2 в системе со смешанным основанием для ρM

2) Сравнить полученное число с M/2

28

Многомодульная система счисления: определение знака числа

Система счисления со смешанным основанием для ρ = [R1 , R2 ,K, RN ]

S = D0 + D1 (R1 ) + D 2 (R1R2 ) +K+ D N−1 (R1R2 KRN−1 )

0 ≤ Di < Ri +1 I = 0,1,K, N −1

Алгоритм перевода в систему счисления со смешанным основанием

Алг Convert_Parallel Вход: src,modules,N Выход: res

tmp=src;

Для i от 1 до N цикл

если threadIdx==i, то res[i]=tmp[i];

кесли

если threadIdx>i, то

tmp[Idx]=(tmp[Idx]-tmp[i]) (mod modules[Idx]); m1 = modules[i]-1 (mod modules[Idx]);

tmp[Idx] = (m1* tmp[Idx]) (mod modules[Idx]);

кесли

кц

29

Знакоразрядная система счисления

Определение:

позиционная система счисления с основанием Q, цифры которой могут принимать как положительные так и отрицательные значения

		dI {− q +1,K,0,K, q −1}
X = D0D1 KDN −1

Пример:

Знакоразрядная система с основанием 10. Положительные значения разрядов обозначаются цифрами от 0 до 9, отрицательные значения разрядов обозначим как 0,1,...,9 (имеют значения от -1 до -9 соответственно). Таким образом 16 = 1 10 + (−6) = 4

Свойства:

1)Число может иметь несколько представлений:1 = 19 = 199 = 199 9 = ...

2)Число 0 имеет единственное представление

3)Знак числа равен знаку старшего разряда

4)Чтобы изменить знак числа, достаточно изменить знаки его цифр

30