Changed FP mode for the kernels. Seems to be required by the compiler Added a trigo...

diff --git a/kernels/binomialOption_kernels.cl b/kernels/binomialOption_kernels.cl
index bcf6e7d..885ab66 100644 (file)
--- a/kernels/binomialOption_kernels.cl
+++ b/kernels/binomialOption_kernels.cl
@@ -72,4 +72,4 @@ binomial_options(
 
     // write result for this block to global mem
     if(tid == 0) output[bid] = callA[0];
-}
\ No newline at end of file
+}

diff --git a/kernels/mersenne_kernels.cl b/kernels/mersenne_kernels.cl
index 9aa2fc8..f2f3d0d 100644 (file)
--- a/kernels/mersenne_kernels.cl
+++ b/kernels/mersenne_kernels.cl
@@ -201,7 +201,7 @@ void gaussianRand(const __global uint4 *seedArray,
         temp1 = convert_float4(temp[i]) * one / intMax;
         temp2 = convert_float4(temp[i + 1]) * one / intMax;
         
-        // Applying Box Mullar Transformations.
+        // Asinf Box Mullar Transformations.
         r = sqrt((-two) * log(temp1));
         phi  = two * PI * temp2;
         gaussianRand[actualPos + i * 2 + 0] = r * cos(phi);

diff --git a/kernels/test_trigo.cl b/kernels/test_trigo.cl
new file mode 100644 (file)
index 0000000..eae4b9a
--- /dev/null
+++ b/kernels/test_trigo.cl
@@ -0,0 +1,7 @@
+__kernel void
+test_trigo(__global float *dst, __global const float *src)
+{
+  const int offset = get_global_id(0);
+  dst[offset] = sin(src[offset]);
+}
+

diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt
index 1380ad7..34b8f0e 100644 (file)
--- a/src/CMakeLists.txt
+++ b/src/CMakeLists.txt
@@ -47,6 +47,7 @@ ADD_EXECUTABLE(test_private_memory tests/test_private_memory.c)
 ADD_EXECUTABLE(test_constant_memory tests/test_constant_memory.c)
 ADD_EXECUTABLE(test_memory_leak tests/test_memory_leak.c)
 ADD_EXECUTABLE(test_perf_report tests/test_perf_report.c)
+ADD_EXECUTABLE(test_trigo tests/test_trigo.c)
 ADD_EXECUTABLE(mersenneTwister tests/mersenneTwister.c)
 ADD_EXECUTABLE(blackscholes tests/blackscholes.c)
 ADD_EXECUTABLE(matmul tests/matmul.c)
@@ -67,6 +68,7 @@ TARGET_LINK_LIBRARIES(test_private_memory cl_test m)
 TARGET_LINK_LIBRARIES(test_constant_memory cl_test m)
 TARGET_LINK_LIBRARIES(test_memory_leak cl_test m)
 TARGET_LINK_LIBRARIES(test_perf_report cl_test m)
+TARGET_LINK_LIBRARIES(test_trigo cl_test m)
 TARGET_LINK_LIBRARIES(mersenneTwister cl_test m)
 TARGET_LINK_LIBRARIES(blackscholes cl_test m)
 TARGET_LINK_LIBRARIES(matmul cl_test m)

diff --git a/src/cl_command_queue.c b/src/cl_command_queue.c
index f9be3cc..c856bc0 100644 (file)
--- a/src/cl_command_queue.c
+++ b/src/cl_command_queue.c
@@ -99,8 +99,6 @@ cl_command_queue_add_ref(cl_command_queue queue)
   atomic_inc(&queue->ref_n);
 }
 
-#define SURFACE_SZ 32
-
 LOCAL cl_int
 cl_command_queue_bind_surface(cl_command_queue queue,
                               cl_kernel k,

diff --git a/src/cl_kernel.c b/src/cl_kernel.c
index ac12d80..990cb0b 100644 (file)
--- a/src/cl_kernel.c
+++ b/src/cl_kernel.c
@@ -305,7 +305,7 @@ cl_kernel_allocate_inline_buffer(cl_kernel k,
   const size_t sz = init->sz;
   cl_int err = CL_SUCCESS;
 
-  FATAL_IF (init->offset % 64, "Bad alignment for inline buffer offset");
+  FATAL_IF (init->offset % SURFACE_SZ, "Bad alignment for inline buffer offset");
   FATAL_IF (k->const_bo != NULL, "inline buffer already declared");
   assert(k->program && k->program->ctx);
   bufmgr = cl_context_get_intel_bufmgr(k->program->ctx);
@@ -314,7 +314,7 @@ cl_kernel_allocate_inline_buffer(cl_kernel k,
                                              sz,
                                              64));
   drm_intel_bo_subdata(k->const_bo, 0, sz, &init->data);
-  k->const_bo_index = init->offset / 64;
+  k->const_bo_index = init->offset / SURFACE_SZ;
   *read += sz;
   *patch += sz;
 

diff --git a/src/cl_kernel.h b/src/cl_kernel.h
index 9c1cd5f..704d1c5 100644 (file)
--- a/src/cl_kernel.h
+++ b/src/cl_kernel.h
@@ -277,6 +277,9 @@ struct _cl_kernel {
   uint8_t ref_its_program;      /* True only for the user kernel (those created by clCreateKernel) */
 };
 
+/* Size of the surface state as encoded in the binary blob */
+#define SURFACE_SZ 32
+
 /* Allocate an empty kernel */
 extern cl_kernel cl_kernel_new(void);
 

diff --git a/src/intel/intel_gpgpu.c b/src/intel/intel_gpgpu.c
index 779e906..566051c 100644 (file)
--- a/src/intel/intel_gpgpu.c
+++ b/src/intel/intel_gpgpu.c
@@ -948,6 +948,7 @@ gpgpu_build_idrt(intel_gpgpu_t *state,
   for (i = 0; i < ker_n; i++) {
     memset(desc, 0, sizeof(*desc));
     desc->desc0.kernel_start_pointer = kernel[i].bo->offset >> 6; /* reloc */
+    desc->desc1.floating_point_mode = 1;
     /* desc->desc1 = 0; - no exception control, no SPF */
     /* desc->desc2 = 0; - no samplers */
     desc->desc2.sampler_state_pointer = state->sampler_state_b.bo->offset >> 5;