Лекция2 - 2015

Сложение в знакоразрядной системе счисления

Сложение чисел за 4 такта: теоретически, за O(1)

1) Вычисление частичных сумм разрядов: pI = aI + bI

	0,−Q −1 < PI < Q −1
			≥ Q −1
2) Вычисление переносов. CI	=	1, PI
		−1, PI	≤ −Q −1

3)	Вычисление вспомогательной величины WI: WI = PI − CI Q
4)	Вычисление разрядов суммы	SI		WI , I = 0
			=	+ CI −1 , I ≠ 0
			WI

31

Сложение и сдвиг в знакоразрядной системе счисления

пример выполнения операции сложения

сдвиг влево

частичная нормализация

1 9 9 9 0 1 9 9

0 1 9 9 << 1 = 1 9 9 0 = 10

32

Знакоразрядная система счисления: исходные коды операций

__device__ irr_num operator+(const irr_num& x, const irr_num& y){

__shared__ int tmpr[MAX_THREADS_IRR];
int w=x.part_val+y.part_val;	//partial sum
irr_num res;
if(w>=0xFFFFFF){
res.part_val=w-0x1000000;
tmpr[threadIdx.x]=1;
}
else{
if(w<=-0xFFFFFF){
res.part_val=w+0x1000000;
tmpr[threadIdx.x]=-1;
}
else{
res.part_val=w;
tmpr[threadIdx.x]=0;
}
}
__syncthreads();
if(threadIdx.x)
res.part_val+=tmpr[threadIdx.x-1];
__syncthreads();
norm(res);
return res;
}
__device__ int sgn(const irr_num& x){
__shared__ int tmpr[MAX_THREADS_IRR];
tmpr[threadIdx.x]=(x.part_val>0)-(x.part_val<0);		//0 for 0, -1 for<0, 1 for>0
__syncthreads();

//reduction

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+s]) //searching ron non-zero element with maximum index tmpr[threadIdx.x]=tmpr[threadIdx.x+s];

__syncthreads();

}

return tmpr[0];

}

33

Умножение в знакоразрядной системе счисления

Алгоритм Mul_IrrNum (A, B, j, q) Вход: A – множимое,

B – множитель,

n – количество разрядов множимого и множителя j – индекс потока

q - основание системы счисления

начало

SM – значение сумматора SM = 0

Для всех разрядов множителя (от старшего к младшему) цикл: i – индекс разряда

PP – частичное произведение

PM – разряды частичного произведения без учёта переносов

CR – переносы, возникшие при вычислении частичного произведения tmp=a[j]*b[i]

CR[j]=tmp/q
PM[j]=tmp%q
CR<<=1
PP = PM+CR
SM<<=1
SM = SM+PP
кц
вернуть SM	34

конец

Сравнение чисел: знакоразрядная система счисления

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

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

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

3)Инверсия полученного результата

сравнение

1)C=A-B

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

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

Выделение старшего разряда C, т.е. редукция C относительно операции Op(a,b)= {a, a!=0; b, a==0}

35

Оценки быстродействия

		Система счисления
	Операция
		Многомодульная	Знакоразрядная
	+, -	O(1)	O(1)
	·	O(1)	O(N)
	==	O(log(N))	O(log(N))
	<,>	O(N)	O(log(N))

36

Зависимость скорости вычислений от
диапазона представимых чисел
Задача: Определение ориентации тройки точек на плоскости.
orientatio n( A,B,C ) = sgn((BX					− AX ) (C Y	− AY )− (BY		− AY ) (CX		− AX ))
	4500
	4000
	3500
	3000
ms
мс time,	2500
											MultiMod
время, Computing
	2000										IrrNum
	1500
	1000
	500
	0
	0	32	64	96	128	160	192	224	256	288	320

Number of threads (proportional to number length)

количество элементов (пропорционально длине числа)

37

Пример исходного кода

template <class Num> __device__ int Orientation(SPoint<Num> &a, SPoint<Num> &b, SPoint<Num> &c){ Num res=(b.y-a.y)*(c.x-a.x)-(b.x-a.x)*(c.y-a.y);

return sgn(res);

}

__device__ multi_mod_num MultiModFromData(const int *data){ multi_mod_num res(data[blockIdx.x+blockIdx.y*gridDim.x]); return res;

}

__global__ void calcMultiMod(int *xa, int *ya, int *xb, int *yb, int *xc, int *yc, int *results){ SPoint<multi_mod_num> a,b,c;

a.x=MultiModFromData(xa); a.y=MultiModFromData(ya); b.x=MultiModFromData(xb); b.y=MultiModFromData(yb); c.x=MultiModFromData(xc); c.y=MultiModFromData(yc); int res=Orientation(a,b,c);

if(threadIdx.x==0)

results[blockIdx.x+blockIdx.y*gridDim.x]=res;

}

#include <cuda.h> #include <stdio.h>

int main(int argc, char **argv){ unsigned n_threads=10;

unsigned n_triangles=8, max_module=2048, max_block_size=1;

//инициализация

//...

calcMultiMod <<<dim3(max_block_size,n_triangles/max_block_size), n_threads, 2*sizeof(int)*n_threads>>> (d_xa, d_ya, d_xb, d_yb, d_xc, d_yc, d_mm);

cudaThreadSynchronize();

//освобождение ресурсов

//...

return 0;

38

}

Литература

•Nvidia CUDA Programming Guide http://developer.nvidia.com/object/gpucomputing.html

•А. В. Боресков, А. А. Харламов «Основы работы с технологией CUDA» http://steps3d.narod.ru/tutorials/cuda-tutorial.html

•http://www.gpgpu.ru/

•Грегори Р., Кришнамурти Е. Безошибочные вычисления. Методы и приложения: Пер. с англ. – М.: Мир, 1988 – 208 с., ил.

39

Вопросы ???

40