Use compound literal to make instances of vector types, now possible …

…to use it to initialize globals.

include/double3.h

@@ -6,6 +6,8 @@
 #ifndef CL_CUDA_DOUBLE3_H_
 #define CL_CUDA_DOUBLE3_H_
 
+#define __make_double3( x, y, z ) (double3){ x, y, z }
+
 __device__ double3 double3_add ( double3 a, double3 b )
 {
   return make_double3 ( a.x + b.x, a.y + b.y, a.z + b.z );

include/double4.h

@@ -6,6 +6,8 @@
 #ifndef CL_CUDA_DOUBLE4_H_
 #define CL_CUDA_DOUBLE4_H_
 
+#define __make_double4( x, y, z, w ) (double4){ x, y, z, w }
+
 __device__ double4 double4_add ( double4 a, double4 b )
 {
   return make_double4 ( a.x + b.x, a.y + b.y, a.z + b.z, a.w + b.w );

include/float3.h

@@ -6,6 +6,8 @@
 #ifndef CL_CUDA_FLOAT3_H_
 #define CL_CUDA_FLOAT3_H_
 
+#define __make_float3( x, y, z ) (float3){ x, y, z }
+
 __device__ float3 float3_add ( float3 a, float3 b )
 {
   return make_float3 ( a.x + b.x, a.y + b.y, a.z + b.z );

include/float4.h

@@ -6,6 +6,8 @@
 #ifndef CL_CUDA_FLOAT4_H_
 #define CL_CUDA_FLOAT4_H_
 
+#define __make_float4( x, y, z, w ) (float4){ x, y, z, w }
+
 __device__ float4 float4_add ( float4 a, float4 b )
 {
   return make_float4 ( a.x + b.x, a.y + b.y, a.z + b.z, a.w + b.w );

src/lang/built-in.lisp

@@ -143,10 +143,10 @@
              ((double) double* nil "&")
              ((curand-state-xorwow) curand-state-xorwow* nil "&"))
     ;; built-in vector constructor
-    float3 (((float float float) float3 nil "make_float3"))
-    float4 (((float float float float) float4 nil "make_float4"))
-    double3 (((double double double) double3 nil "make_double3"))
-    double4 (((double double double double) double4 nil "make_double4"))
+    float3 (((float float float) float3 nil "__make_float3"))
+    float4 (((float float float float) float4 nil "__make_float4"))
+    double3 (((double double double) double3 nil "__make_double3"))
+    double4 (((double double double double) double4 nil "__make_double4"))
     ;; Synchronization functions
     syncthreads ((() void nil "__syncthreads"))
     ;; type casting intrinsics

t/api/defkernel.lisp

@@ -266,6 +266,7 @@
 ;;;
 ;;; test MOD
 ;;;
+
 (defkernel test-mod (void ((x int*)))
   (set (aref x 0) (mod (aref x 0) 5)))
 
@@ -280,6 +281,40 @@
         "basic case 23")))
 
 
+;;;
+;;; Initializers
+;;;
+
+(defglobal c (float3 3.0 2.0 1.0))
+
+(defkernel initializer (float3 ())
+  (let ((x 1.0))
+    (return (float3 x 2.0 3.0))))
+
+(defkernel use-initializer (void ((x float3*) (y float3*)))
+  (set (aref x 0) (initializer))
+  (set (aref y 0) c)
+  (return))
+
+(subtest "Initializers"
+
+  (with-cuda (0)
+    (with-memory-blocks ((x 'float3 1)
+                         (y 'float3 1))
+      (use-initializer x y :grid-dim (list 1 1 1)
+                           :block-dim (list 1 1 1))
+      (sync-memory-block x :device-to-host)
+      (sync-memory-block y :device-to-host)
+      (is (memory-block-aref x 0)
+          (make-float3 1.0 2.0 3.0)
+          :test #'float3-=
+          "Ok. - returning with initializer")
+      (is (memory-block-aref y 0)
+          (make-float3 3.0 2.0 1.0)
+          :test #'float3-=
+          "Ok. - initializing with initializer"))))
+
+
 ;;;
 ;;; test EXPAND-MACRO function
 ;;;