added unary minus operator for matrix and float3; updated tests

2 files changed, 35 insertions, 3 deletions

diff --git a/tests/linalg_test.cu b/tests/linalg_test.cu
index bce5ea6..b61488f 100644
--- a/tests/linalg_test.cu
+++ b/tests/linalg_test.cu
@@ -99,16 +99,28 @@ __global__ void floatdivfloat3(float3 *a, float c, float3 *dest)
 	dest[idx] = c/a[idx];
 }
 
-__global__ void dot(float3 *a, float3 *b, float* dest)
+__global__ void dot(float3 *a, float3 *b, float *dest)
 {
 	int idx = blockIdx.x*blockDim.x + threadIdx.x;
 	dest[idx] = dot(a[idx],b[idx]);
 }
 
-__global__ void cross(float3 *a, float3 *b, float3* dest)
+__global__ void cross(float3 *a, float3 *b, float3 *dest)
 {
 	int idx = blockIdx.x*blockDim.x + threadIdx.x;
 	dest[idx] = cross(a[idx],b[idx]);
 }
 
+__global__ void norm(float3 *a, float *dest)
+{
+	int idx = blockIdx.x*blockDim.x + threadIdx.x;
+	dest[idx] = norm(a[idx]);
+}
+
+__global__ void minusfloat3(float3 *a, float3 *dest)
+{
+	int idx = blockIdx.x*blockDim.x + threadIdx.x;
+	dest[idx] = -a[idx];
+}
+
 } // extern "c"
diff --git a/tests/linalg_test.py b/tests/linalg_test.py
index f5e947e..44c4b52 100644
--- a/tests/linalg_test.py
+++ b/tests/linalg_test.py
@@ -30,8 +30,10 @@ float3divfloatequal = mod.get_function('float3divfloatequal')
 floatdivfloat3 = mod.get_function('floatdivfloat3')
 dot = mod.get_function('dot')
 cross = mod.get_function('cross')
+norm = mod.get_function('norm')
+minusfloat3 = mod.get_function('minusfloat3')
 
-size = {'block': (100,1,1), 'grid': (1,1)}
+size = {'block': (256,1,1), 'grid': (1,1)}
 
 a = np.empty(size['block'][0], dtype=float3)
 b = np.empty(size['block'][0], dtype=float3)
@@ -187,4 +189,22 @@ def testcross():
         if not np.allclose(wdest['x'], w[0]) or \
                 not np.allclose(wdest['y'], w[1]) or \
                 not np.allclose(wdest['z'], w[2]):
+            print w
+            print wdest
             assert False
+
+def testnorm():
+    dest = np.empty(a.size, dtype=np.float32)
+    norm(cuda.In(a), cuda.Out(dest), **size)
+
+    for i in range(len(dest)):
+        if not np.allclose(np.linalg.norm((a['x'][i],a['y'][i],a['z'][i])), dest[i]):
+            assert False
+
+def testminusfloat3():
+    dest = np.empty(a.size, dtype=float3)
+    minusfloat3(cuda.In(a), cuda.Out(dest), **size)
+    if not np.allclose(-a['x'], dest['x']) or \
+            not np.allclose(-a['y'], dest['y']) or \
+            not np.allclose(-a['z'], dest['z']):
+        assert False