1 files changed, 36 insertions, 12 deletions

diff --git a/tests/linalg_test.cu b/tests/linalg_test.cu
index 13d2ed0..81043f1 100644
--- a/tests/linalg_test.cu
+++ b/tests/linalg_test.cu
@@ -3,78 +3,102 @@
 extern "C"
 {
 
-__global__ void add(float3 *a, float3 *b, float3 *dest)
+__global__ void float3add(float3 *a, float3 *b, float3 *dest)
 {
 	int idx = blockIdx.x*blockDim.x + threadIdx.x;
 	dest[idx] = a[idx] + b[idx];
 }
 
-__global__ void addequal(float3 *a, float3 *b)
+__global__ void float3addequal(float3 *a, float3 *b)
 {
 	int idx = blockIdx.x*blockDim.x + threadIdx.x;
 	a[idx] += b[idx];
 }
 
-__global__ void sub(float3 *a, float3 *b, float3 *dest)
+__global__ void float3sub(float3 *a, float3 *b, float3 *dest)
 {
 	int idx = blockIdx.x*blockDim.x + threadIdx.x;
 	dest[idx] = a[idx] - b[idx];
 }
 
-__global__ void subequal(float3 *a, float3 *b)
+__global__ void float3subequal(float3 *a, float3 *b)
 {
 	int idx = blockIdx.x*blockDim.x + threadIdx.x;
 	a[idx] -= b[idx];
 }
 
-__global__ void addfloat(float3 *a, float c, float3 *dest)
+__global__ void float3addfloat(float3 *a, float c, float3 *dest)
 {
 	int idx = blockIdx.x*blockDim.x + threadIdx.x;
 	dest[idx] = a[idx] + c;
 }
 
-__global__ void addfloatequal(float3 *a, float c)
+__global__ void float3addfloatequal(float3 *a, float c)
 {
 	int idx = blockIdx.x*blockDim.x + threadIdx.x;
 	a[idx] += c;
 }
 
-__global__ void subfloat(float3 *a, float c, float3 *dest)
+__global__ void floataddfloat3(float3 *a, float c, float3 *dest)
+{
+	int idx = blockIdx.x*blockDim.x + threadIdx.x;
+	dest[idx] = c + a[idx];
+}
+
+__global__ void float3subfloat(float3 *a, float c, float3 *dest)
 {
 	int idx = blockIdx.x*blockDim.x + threadIdx.x;
 	dest[idx] = a[idx] - c;
 }
 
-__global__ void subfloatequal(float3 *a, float c)
+__global__ void float3subfloatequal(float3 *a, float c)
 {
 	int idx = blockIdx.x*blockDim.x + threadIdx.x;
 	a[idx] -= c;
 }
 
-__global__ void mulfloat(float3 *a, float c, float3 *dest)
+__global__ void floatsubfloat3(float3 *a, float c, float3 *dest)
+{
+	int idx = blockIdx.x*blockDim.x + threadIdx.x;
+	dest[idx] = c - a[idx];
+}
+
+__global__ void float3mulfloat(float3 *a, float c, float3 *dest)
 {
 	int idx = blockIdx.x*blockDim.x + threadIdx.x;
 	dest[idx] = a[idx]*c;
 }
 
-__global__ void mulfloatequal(float3 *a, float c)
+__global__ void float3mulfloatequal(float3 *a, float c)
 {
 	int idx = blockIdx.x*blockDim.x + threadIdx.x;
 	a[idx] *= c;
 }
 
-__global__ void divfloat(float3 *a, float c, float3 *dest)
+__global__ void floatmulfloat3(float3 *a, float c, float3 *dest)
+{
+	int idx = blockIdx.x*blockDim.x + threadIdx.x;
+	dest[idx] = c*a[idx];
+}
+
+__global__ void float3divfloat(float3 *a, float c, float3 *dest)
 {
 	int idx = blockIdx.x*blockDim.x + threadIdx.x;
 	dest[idx] = a[idx]/c;
 }
 
-__global__ void divfloatequal(float3 *a, float c)
+__global__ void float3divfloatequal(float3 *a, float c)
 {
 	int idx = blockIdx.x*blockDim.x + threadIdx.x;
 	a[idx] /= c;
 }
 
+__global__ void floatdivfloat3(float3 *a, float c, float3 *dest)
+{
+	int idx = blockIdx.x*blockDim.x + threadIdx.x;
+	dest[idx] = c/a[idx];
+}
+
 __global__ void dot(float3 *a, float3 *b, float* dest)
 {
 	int idx = blockIdx.x*blockDim.x + threadIdx.x;