QUDA
v0.5.0
A library for QCD on GPUs
Main Page
Namespaces
Classes
Files
File List
File Members
All
Classes
Namespaces
Files
Functions
Variables
Typedefs
Enumerations
Enumerator
Friends
Macros
Pages
quda
include
double_single.h
Go to the documentation of this file.
1
#pragma once
2
3
__device__
inline
void
dsadd
(
volatile
float2 &c,
const
volatile
float2 &a,
const
volatile
float2 &b) {
4
float
t1 = a.x + b.x;
5
float
e = t1 - a.x;
6
float
t2 = ((b.x - e) + (a.x - (t1 - e))) + a.y + b.y;
7
// The result is t1 + t2, after normalization.
8
c.x = e = t1 + t2;
9
c.y = t2 - (e - t1);
10
}
11
12
struct
doublesingle
{
13
float2
a
;
14
__device__
inline
doublesingle
() :
a
(make_float2(0.0f,0.0f)) { ; }
15
__device__
inline
doublesingle
(
const
volatile
doublesingle
&b) :
a
(make_float2(b.
a
.
x
, b.
a
.y)) { ; }
16
__device__
inline
doublesingle
(
const
float
a
) : a(make_float2(a, 0.0)) { ; }
17
18
__device__
inline
void
operator+=
(
const
doublesingle
&b) {
dsadd
(this->
a
, this->
a
, b.
a
); }
19
__device__
inline
void
operator+=
(
const
volatile
doublesingle
&b) {
dsadd
(this->
a
, this->
a
, b.
a
); }
20
__device__
inline
void
operator+=
(
const
float
&b) {
21
float2 b2 = make_float2(b, 0.0);
22
dsadd
(this->
a
, this->
a
, b2); }
23
24
__device__
inline
doublesingle
&
operator=
(
const
doublesingle
&b)
25
{
a
.x = b.
a
.x;
a
.y = b.
a
.y;
return
*
this
; }
26
27
__device__
inline
doublesingle
&
operator=
(
const
float
&b)
28
{
a
.x = b;
a
.y = 0.0f;
return
*
this
; }
29
};
30
31
__device__
inline
volatile
doublesingle
operator+=
(
volatile
doublesingle
&a,
const
volatile
doublesingle
&b)
32
{
dsadd
(a.
a
, a.
a
, b.
a
);
return
a;}
33
34
__host__
double
operator+=
(
double
& a,
doublesingle
&b) { a += b.
a
.x; a += b.
a
.y;
return
a; }
35
36
struct
doublesingle2
{
37
doublesingle
x
;
38
doublesingle
y
;
39
__device__
inline
doublesingle2
&
operator=
(
const
double
&a)
40
{
x
= a;
y
= a;
return
*
this
; }
41
__device__
inline
void
operator+=
(
const
double2 &b) {
x
+= b.x;
y
+= b.y;}
42
__device__
inline
void
operator+=
(
const
doublesingle2
&b) {
x
+= b.
x
;
y
+= b.
y
;}
43
};
44
45
__host__ double2
operator+=
(double2& a,
doublesingle2
&b)
46
{ a.x += b.
x
.
a
.x; a.x += b.
x
.
a
.y; a.y += b.
y
.
a
.x; a.y += b.
y
.
a
.y;
return
a; }
47
48
struct
doublesingle3
{
49
doublesingle
x
;
50
doublesingle
y
;
51
doublesingle
z
;
52
__device__
inline
void
operator+=
(
const
double3 &b) {
x
+= b.x;
y
+= b.y;
z
+= b.z;}
53
__device__
inline
void
operator+=
(
const
doublesingle3
&b) {
x
+= b.
x
;
y
+= b.
y
;
z
+= b.
z
;}
54
};
55
56
__host__ double3
operator+=
(double3& a,
doublesingle3
&b)
57
{ a.x += b.
x
.
a
.x; a.x += b.
x
.
a
.y; a.y += b.
y
.
a
.x; a.y += b.
y
.
a
.y; a.z += b.
z
.
a
.x; a.z += b.
z
.
a
.y;
return
a; }
Generated on Wed Mar 20 2013 12:52:14 for QUDA by
1.8.2