gdb/testsuite/ChangeLog

2011-01-25  Ken Werner  <ken.werner@de.ibm.com>

        * gdb.opencl/convs_casts.cl: Move program scope variables into the
        OpenCL kernel function. Add a comment as marker. Add address space
        qualifiers for the remaining program scope variables.
        * gdb.opencl/datatypes.cl: Likewise.
        * gdb.opencl/operators.cl: Likewise.
        * gdb.opencl/vec_comps.cl: Likewise.
        * gdb.opencl/convs_casts.exp: Replace gdb_test_multiple by gdb_test.
        Add breakpoint at the marker comment.
        * gdb.opencl/datatypes.exp: Likewise.
        * gdb.opencl/operators.exp: Likewise.
        * gdb.opencl/vec_comps.exp: Likewise.

9 files changed, 229 insertions, 191 deletions

diff --git a/gdb/testsuite/ChangeLog b/gdb/testsuite/ChangeLog
index 6a0259e..f1d5911 100644
--- a/gdb/testsuite/ChangeLog
+++ b/gdb/testsuite/ChangeLog
@@ -1,3 +1,17 @@
+2011-01-25  Ken Werner  <ken.werner@de.ibm.com>
+
+	* gdb.opencl/convs_casts.cl: Move program scope variables into the
+	OpenCL kernel function. Add a comment as marker. Add address space
+	qualifiers for the remaining program scope variables.
+	* gdb.opencl/datatypes.cl: Likewise.
+	* gdb.opencl/operators.cl: Likewise.
+	* gdb.opencl/vec_comps.cl: Likewise.
+	* gdb.opencl/convs_casts.exp: Replace gdb_test_multiple by gdb_test.
+	Add breakpoint at the marker comment.
+	* gdb.opencl/datatypes.exp: Likewise.
+	* gdb.opencl/operators.exp: Likewise.
+	* gdb.opencl/vec_comps.exp: Likewise.
+
 2011-01-24  Pedro Alves  <pedro@codesourcery.com>
 
 	* gdb.base/printcmds.c (some_struct): New struct and instance.
diff --git a/gdb/testsuite/gdb.opencl/convs_casts.cl b/gdb/testsuite/gdb.opencl/convs_casts.cl
index 8ab0805..c62d614 100644
--- a/gdb/testsuite/gdb.opencl/convs_casts.cl
+++ b/gdb/testsuite/gdb.opencl/convs_casts.cl
@@ -17,39 +17,41 @@
 
    Contributed by Ken Werner <ken.werner@de.ibm.com>  */
 
-int opencl_version = __OPENCL_VERSION__;
+__constant int opencl_version = __OPENCL_VERSION__;
 
 #ifdef HAVE_cl_khr_fp64
 #pragma OPENCL EXTENSION cl_khr_fp64 : enable
-int have_cl_khr_fp64 = 1;
+__constant int have_cl_khr_fp64 = 1;
 #else
-int have_cl_khr_fp64 = 0;
+__constant int have_cl_khr_fp64 = 0;
 #endif
 
 #ifdef HAVE_cl_khr_fp16
 #pragma OPENCL EXTENSION cl_khr_fp16 : enable
-int have_cl_khr_fp16 = 1;
+__constant int have_cl_khr_fp16 = 1;
 #else
-int have_cl_khr_fp16 = 0;
+__constant int have_cl_khr_fp16 = 0;
 #endif
 
-char c = 123;
-uchar uc = 123;
-short s = 123;
-ushort us = 123;
-int i = 123;
-uint ui = 123;
-long l = 123;
-ulong ul = 123;
+__kernel void testkernel (__global int *data)
+{
+  char c = 123;
+  uchar uc = 123;
+  short s = 123;
+  ushort us = 123;
+  int i = 123;
+  uint ui = 123;
+  long l = 123;
+  ulong ul = 123;
 #ifdef cl_khr_fp16
-half h = 123.0;
+  half h = 123.0;
 #endif
-float f = 123.0;
+  float f = 123.0;
 #ifdef cl_khr_fp64
-double d = 123.0;
+  double d = 123.0;
 #endif
 
-__kernel void testkernel (__global int *data)
-{
+  /* marker! */
+
   data[get_global_id(0)] = 1;
 }
diff --git a/gdb/testsuite/gdb.opencl/convs_casts.exp b/gdb/testsuite/gdb.opencl/convs_casts.exp
index 86a44e2..0dbfcdc 100644
--- a/gdb/testsuite/gdb.opencl/convs_casts.exp
+++ b/gdb/testsuite/gdb.opencl/convs_casts.exp
@@ -40,15 +40,19 @@ if { [gdb_compile_opencl_hostapp "${clprogram}" "${testfile}" "" ] != "" } {
 clean_restart ${testfile}
 
 # Set breakpoint at the OpenCL kernel
-gdb_test_multiple "break testkernel" "set pending breakpoint" {
-     -re ".*Function \"testkernel\" not defined.*Make breakpoint pending.*y or \\\[n\\\]. $" {
-            gdb_test "y" "Breakpoint.*testkernel.*pending." "set pending breakpoint (without symbols)"
-     }
-}
+gdb_test "break testkernel" \
+    "" \
+    "Set pending breakpoint" \
+    ".*Function \"testkernel\" not defined.*Make breakpoint pending.*y or \\\[n\\\]. $" \
+    "y"
 
 gdb_run_cmd
 gdb_test "" ".*Breakpoint.*1.*testkernel.*" "run"
 
+# Continue to the marker
+gdb_breakpoint [gdb_get_line_number "marker" "${clprogram}"]
+gdb_continue_to_breakpoint "marker"
+
 # Retrieve some information about availability of OpenCL extensions
 set have_cl_khr_fp64 [get_integer_valueof "have_cl_khr_fp64" 0]
 set have_cl_khr_fp16 [get_integer_valueof "have_cl_khr_fp16" 0]
diff --git a/gdb/testsuite/gdb.opencl/datatypes.cl b/gdb/testsuite/gdb.opencl/datatypes.cl
index ab982317..e3b7787 100644
--- a/gdb/testsuite/gdb.opencl/datatypes.cl
+++ b/gdb/testsuite/gdb.opencl/datatypes.cl
@@ -17,129 +17,131 @@
 
    Contributed by Ken Werner <ken.werner@de.ibm.com>  */
 
-int opencl_version = __OPENCL_VERSION__;
+__constant int opencl_version = __OPENCL_VERSION__;
 
 #ifdef HAVE_cl_khr_fp64
 #pragma OPENCL EXTENSION cl_khr_fp64 : enable
-int have_cl_khr_fp64 = 1;
+__constant int have_cl_khr_fp64 = 1;
 #else
-int have_cl_khr_fp64 = 0;
+__constant int have_cl_khr_fp64 = 0;
 #endif
 
 #ifdef HAVE_cl_khr_fp16
 #pragma OPENCL EXTENSION cl_khr_fp16 : enable
-int have_cl_khr_fp16 = 1;
+__constant int have_cl_khr_fp16 = 1;
 #else
-int have_cl_khr_fp16 = 0;
+__constant int have_cl_khr_fp16 = 0;
 #endif
 
-bool b = 0;
+__kernel void testkernel (__global int *data)
+{
+  bool b = 0;
 
-char   c   = 1;
-char2  c2  = (char2) (1, 2);
+  char   c   = 1;
+  char2  c2  = (char2) (1, 2);
 #ifdef CL_VERSION_1_1
-char3  c3  = (char3) (1, 2, 3);
+  char3  c3  = (char3) (1, 2, 3);
 #endif
-char4  c4  = (char4) (1, 2, 3, 4);
-char8  c8  = (char8) (1, 2, 3, 4, 5, 6, 7, 8);
-char16 c16 = (char16)(1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16);
+  char4  c4  = (char4) (1, 2, 3, 4);
+  char8  c8  = (char8) (1, 2, 3, 4, 5, 6, 7, 8);
+  char16 c16 = (char16)(1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16);
 
-uchar   uc   = 1;
-uchar2  uc2  = (uchar2) (1, 2);
+  uchar   uc   = 1;
+  uchar2  uc2  = (uchar2) (1, 2);
 #ifdef CL_VERSION_1_1
-uchar3  uc3  = (uchar3) (1, 2, 3);
+  uchar3  uc3  = (uchar3) (1, 2, 3);
 #endif
-uchar4  uc4  = (uchar4) (1, 2, 3, 4);
-uchar8  uc8  = (uchar8) (1, 2, 3, 4, 5, 6, 7, 8);
-uchar16 uc16 = (uchar16)(1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16);
+  uchar4  uc4  = (uchar4) (1, 2, 3, 4);
+  uchar8  uc8  = (uchar8) (1, 2, 3, 4, 5, 6, 7, 8);
+  uchar16 uc16 = (uchar16)(1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16);
 
-short   s   = -1;
-short2  s2  = (short2) (-1, -2);
+  short   s   = -1;
+  short2  s2  = (short2) (-1, -2);
 #ifdef CL_VERSION_1_1
-short3  s3  = (short3) (-1, -2, -3);
+  short3  s3  = (short3) (-1, -2, -3);
 #endif
-short4  s4  = (short4) (-1, -2, -3, -4);
-short8  s8  = (short8) (-1, -2, -3, -4, -5, -6, -7, -8);
-short16 s16 = (short16)(-1, -2, -3, -4, -5, -6, -7, -8, -9, -10, -11, -12, -13, -14, -15, -16);
+  short4  s4  = (short4) (-1, -2, -3, -4);
+  short8  s8  = (short8) (-1, -2, -3, -4, -5, -6, -7, -8);
+  short16 s16 = (short16)(-1, -2, -3, -4, -5, -6, -7, -8, -9, -10, -11, -12, -13, -14, -15, -16);
 
-ushort   us   = 1;
-ushort2  us2  = (ushort2) (1, 2);
+  ushort   us   = 1;
+  ushort2  us2  = (ushort2) (1, 2);
 #ifdef CL_VERSION_1_1
-ushort3  us3  = (ushort3) (1, 2, 3);
+  ushort3  us3  = (ushort3) (1, 2, 3);
 #endif
-ushort4  us4  = (ushort4) (1, 2, 3, 4);
-ushort8  us8  = (ushort8) (1, 2, 3, 4, 5, 6, 7, 8);
-ushort16 us16 = (ushort16)(1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16);
+  ushort4  us4  = (ushort4) (1, 2, 3, 4);
+  ushort8  us8  = (ushort8) (1, 2, 3, 4, 5, 6, 7, 8);
+  ushort16 us16 = (ushort16)(1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16);
 
-int   i   = -1;
-int2  i2  = (int2) (-1, -2);
+  int   i   = -1;
+  int2  i2  = (int2) (-1, -2);
 #ifdef CL_VERSION_1_1
-int3  i3  = (int3) (-1, -2, -3);
+  int3  i3  = (int3) (-1, -2, -3);
 #endif
-int4  i4  = (int4) (-1, -2, -3, -4);
-int8  i8  = (int8) (-1, -2, -3, -4, -5, -6, -7, -8);
-int16 i16 = (int16)(-1, -2, -3, -4, -5, -6, -7, -8, -9, -10, -11, -12, -13, -14, -15, -16);
+  int4  i4  = (int4) (-1, -2, -3, -4);
+  int8  i8  = (int8) (-1, -2, -3, -4, -5, -6, -7, -8);
+  int16 i16 = (int16)(-1, -2, -3, -4, -5, -6, -7, -8, -9, -10, -11, -12, -13, -14, -15, -16);
 
-uint   ui   = 1;
-uint2  ui2  = (uint2) (1, 2);
+  uint   ui   = 1;
+  uint2  ui2  = (uint2) (1, 2);
 #ifdef CL_VERSION_1_1
-uint3  ui3  = (uint3) (1, 2, 3);
+  uint3  ui3  = (uint3) (1, 2, 3);
 #endif
-uint4  ui4  = (uint4) (1, 2, 3, 4);
-uint8  ui8  = (uint8) (1, 2, 3, 4, 5, 6, 7, 8);
-uint16 ui16 = (uint16)(1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16);
+  uint4  ui4  = (uint4) (1, 2, 3, 4);
+  uint8  ui8  = (uint8) (1, 2, 3, 4, 5, 6, 7, 8);
+  uint16 ui16 = (uint16)(1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16);
 
-long   l   = -1;
-long2  l2  = (long2) (-1, -2);
+  long   l   = -1;
+  long2  l2  = (long2) (-1, -2);
 #ifdef CL_VERSION_1_1
-long3  l3  = (long3) (-1, -2, -3);
+  long3  l3  = (long3) (-1, -2, -3);
 #endif
-long4  l4  = (long4) (-1, -2, -3, -4);
-long8  l8  = (long8) (-1, -2, -3, -4, -5, -6, -7, -8);
-long16 l16 = (long16)(-1, -2, -3, -4, -5, -6, -7, -8, -9, -10, -11, -12, -13, -14, -15, -16);
+  long4  l4  = (long4) (-1, -2, -3, -4);
+  long8  l8  = (long8) (-1, -2, -3, -4, -5, -6, -7, -8);
+  long16 l16 = (long16)(-1, -2, -3, -4, -5, -6, -7, -8, -9, -10, -11, -12, -13, -14, -15, -16);
 
-ulong   ul   = 1;
-ulong2  ul2  = (ulong2) (1, 2);
+  ulong   ul   = 1;
+  ulong2  ul2  = (ulong2) (1, 2);
 #ifdef CL_VERSION_1_1
-ulong3  ul3  = (ulong3) (1, 2, 3);
+  ulong3  ul3  = (ulong3) (1, 2, 3);
 #endif
-ulong4  ul4  = (ulong4) (1, 2, 3, 4);
-ulong8  ul8  = (ulong8) (1, 2, 3, 4, 5, 6, 7, 8);
-ulong16 ul16 = (ulong16)(1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16);
+  ulong4  ul4  = (ulong4) (1, 2, 3, 4);
+  ulong8  ul8  = (ulong8) (1, 2, 3, 4, 5, 6, 7, 8);
+  ulong16 ul16 = (ulong16)(1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15, 16);
 
-half *ph;
+  half *ph;
 #ifdef cl_khr_fp16
-half   h   = 1.0;
-half2  h2  = (half2) (1.0, 2.0);
+  half   h   = 1.0;
+  half2  h2  = (half2) (1.0, 2.0);
 #ifdef CL_VERSION_1_1
-half3  h3  = (half3) (1.0, 2.0, 3.0);
+  half3  h3  = (half3) (1.0, 2.0, 3.0);
 #endif
-half4  h4  = (half4) (1.0, 2.0, 3.0, 4.0);
-half8  h8  = (half8) (1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0);
-half16 h16 = (half16)(1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0, 11.0, 12.0, 13.0, 14.0, 15.0, 16.0);
+  half4  h4  = (half4) (1.0, 2.0, 3.0, 4.0);
+  half8  h8  = (half8) (1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0);
+  half16 h16 = (half16)(1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0, 11.0, 12.0, 13.0, 14.0, 15.0, 16.0);
 #endif
 
-float   f   = 1.0;
-float2  f2  = (float2) (1.0, 2.0);
+  float   f   = 1.0;
+  float2  f2  = (float2) (1.0, 2.0);
 #ifdef CL_VERSION_1_1
-float3  f3  = (float3) (1.0, 2.0, 3.0);
+  float3  f3  = (float3) (1.0, 2.0, 3.0);
 #endif
-float4  f4  = (float4) (1.0, 2.0, 3.0, 4.0);
-float8  f8  = (float8) (1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0);
-float16 f16 = (float16)(1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0, 11.0, 12.0, 13.0, 14.0, 15.0, 16.0);
+  float4  f4  = (float4) (1.0, 2.0, 3.0, 4.0);
+  float8  f8  = (float8) (1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0);
+  float16 f16 = (float16)(1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0, 11.0, 12.0, 13.0, 14.0, 15.0, 16.0);
 
 #ifdef cl_khr_fp64
-double   d   = 1.0;
-double2  d2  = (double2) (1.0, 2.0);
+  double   d   = 1.0;
+  double2  d2  = (double2) (1.0, 2.0);
 #ifdef CL_VERSION_1_1
-double3  d3  = (double3) (1.0, 2.0, 3.0);
+  double3  d3  = (double3) (1.0, 2.0, 3.0);
 #endif
-double4  d4  = (double4) (1.0, 2.0, 3.0, 4.0);
-double8  d8  = (double8) (1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0);
-double16 d16 = (double16)(1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0, 11.0, 12.0, 13.0, 14.0, 15.0, 16.0);
+  double4  d4  = (double4) (1.0, 2.0, 3.0, 4.0);
+  double8  d8  = (double8) (1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0);
+  double16 d16 = (double16)(1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0, 11.0, 12.0, 13.0, 14.0, 15.0, 16.0);
 #endif
 
-__kernel void testkernel (__global int *data)
-{
+  /* marker! */
+
   data[get_global_id(0)] = 1;
 }
diff --git a/gdb/testsuite/gdb.opencl/datatypes.exp b/gdb/testsuite/gdb.opencl/datatypes.exp
index a79e981..1fa80e4cc 100644
--- a/gdb/testsuite/gdb.opencl/datatypes.exp
+++ b/gdb/testsuite/gdb.opencl/datatypes.exp
@@ -206,15 +206,19 @@ gdb_reinitialize_dir $srcdir/$subdir
 gdb_load ${objdir}/${subdir}/${testfile}
 
 # Set breakpoint at the OpenCL kernel
-gdb_test_multiple "break testkernel" "set pending breakpoint" {
-     -re ".*Function \"testkernel\" not defined.*Make breakpoint pending.*y or \\\[n\\\]. $" {
-            gdb_test "y" "Breakpoint.*testkernel.*pending." "set pending breakpoint (without symbols)"
-     }
-}
+gdb_test "break testkernel" \
+    "" \
+    "Set pending breakpoint" \
+    ".*Function \"testkernel\" not defined.*Make breakpoint pending.*y or \\\[n\\\]. $" \
+    "y"
 
 gdb_run_cmd
 gdb_test "" ".*Breakpoint.*1.*testkernel.*" "run"
 
+# Continue to the marker
+gdb_breakpoint [gdb_get_line_number "marker" "${clprogram}"]
+gdb_continue_to_breakpoint "marker"
+
 # Check if the language was switched to opencl
 gdb_test "show language" "The current source language is \"auto; currently opencl\"\."
 
diff --git a/gdb/testsuite/gdb.opencl/operators.cl b/gdb/testsuite/gdb.opencl/operators.cl
index 2f12772..fd5110e 100644
--- a/gdb/testsuite/gdb.opencl/operators.cl
+++ b/gdb/testsuite/gdb.opencl/operators.cl
@@ -17,89 +17,91 @@
 
    Contributed by Ken Werner <ken.werner@de.ibm.com>  */
 
-int opencl_version = __OPENCL_VERSION__;
+__constant int opencl_version = __OPENCL_VERSION__;
 
 #ifdef HAVE_cl_khr_fp64
 #pragma OPENCL EXTENSION cl_khr_fp64 : enable
-int have_cl_khr_fp64 = 1;
+__constant int have_cl_khr_fp64 = 1;
 #else
-int have_cl_khr_fp64 = 0;
+__constant int have_cl_khr_fp64 = 0;
 #endif
 
 #ifdef HAVE_cl_khr_fp16
 #pragma OPENCL EXTENSION cl_khr_fp16 : enable
-int have_cl_khr_fp16 = 1;
+__constant int have_cl_khr_fp16 = 1;
 #else
-int have_cl_khr_fp16 = 0;
+__constant int have_cl_khr_fp16 = 0;
 #endif
 
-char ca = 2;
-char cb = 1;
-uchar uca = 2;
-uchar ucb = 1;
-char4 c4a = (char4) (2, 4, 8, 16);
-char4 c4b = (char4) (1, 2, 8, 4);
-uchar4 uc4a = (uchar4) (2, 4, 8, 16);
-uchar4 uc4b = (uchar4) (1, 2, 8, 4);
-
-short sa = 2;
-short sb = 1;
-ushort usa = 2;
-ushort usb = 1;
-short4 s4a = (short4) (2, 4, 8, 16);
-short4 s4b = (short4) (1, 2, 8, 4);
-ushort4 us4a = (ushort4) (2, 4, 8, 16);
-ushort4 us4b = (ushort4) (1, 2, 8, 4);
-
-int ia = 2;
-int ib = 1;
-uint uia = 2;
-uint uib = 1;
-int4 i4a = (int4) (2, 4, 8, 16);
-int4 i4b = (int4) (1, 2, 8, 4);
-uint4 ui4a = (uint4) (2, 4, 8, 16);
-uint4 ui4b = (uint4) (1, 2, 8, 4);
-
-long la = 2;
-long lb = 1;
-ulong ula = 2;
-ulong ulb = 1;
-long4 l4a = (long4) (2, 4, 8, 16);
-long4 l4b = (long4) (1, 2, 8, 4);
-ulong4 ul4a = (ulong4) (2, 4, 8, 16);
-ulong4 ul4b = (ulong4) (1, 2, 8, 4);
+__kernel void testkernel (__global int *data)
+{
+  char ca = 2;
+  char cb = 1;
+  uchar uca = 2;
+  uchar ucb = 1;
+  char4 c4a = (char4) (2, 4, 8, 16);
+  char4 c4b = (char4) (1, 2, 8, 4);
+  uchar4 uc4a = (uchar4) (2, 4, 8, 16);
+  uchar4 uc4b = (uchar4) (1, 2, 8, 4);
+
+  short sa = 2;
+  short sb = 1;
+  ushort usa = 2;
+  ushort usb = 1;
+  short4 s4a = (short4) (2, 4, 8, 16);
+  short4 s4b = (short4) (1, 2, 8, 4);
+  ushort4 us4a = (ushort4) (2, 4, 8, 16);
+  ushort4 us4b = (ushort4) (1, 2, 8, 4);
+
+  int ia = 2;
+  int ib = 1;
+  uint uia = 2;
+  uint uib = 1;
+  int4 i4a = (int4) (2, 4, 8, 16);
+  int4 i4b = (int4) (1, 2, 8, 4);
+  uint4 ui4a = (uint4) (2, 4, 8, 16);
+  uint4 ui4b = (uint4) (1, 2, 8, 4);
+
+  long la = 2;
+  long lb = 1;
+  ulong ula = 2;
+  ulong ulb = 1;
+  long4 l4a = (long4) (2, 4, 8, 16);
+  long4 l4b = (long4) (1, 2, 8, 4);
+  ulong4 ul4a = (ulong4) (2, 4, 8, 16);
+  ulong4 ul4b = (ulong4) (1, 2, 8, 4);
 
 #ifdef cl_khr_fp16
-half ha = 2;
-half hb = 1;
-half4 h4a = (half4) (2, 4, 8, 16);
-half4 h4b = (half4) (1, 2, 8, 4);
+  half ha = 2;
+  half hb = 1;
+  half4 h4a = (half4) (2, 4, 8, 16);
+  half4 h4b = (half4) (1, 2, 8, 4);
 #endif
 
-float fa = 2;
-float fb = 1;
-float4 f4a = (float4) (2, 4, 8, 16);
-float4 f4b = (float4) (1, 2, 8, 4);
+  float fa = 2;
+  float fb = 1;
+  float4 f4a = (float4) (2, 4, 8, 16);
+  float4 f4b = (float4) (1, 2, 8, 4);
 
 #ifdef cl_khr_fp64
-double da = 2;
-double db = 1;
-double4 d4a = (double4) (2, 4, 8, 16);
-double4 d4b = (double4) (1, 2, 8, 4);
+  double da = 2;
+  double db = 1;
+  double4 d4a = (double4) (2, 4, 8, 16);
+  double4 d4b = (double4) (1, 2, 8, 4);
 #endif
 
-uint4 ui4 = (uint4) (2, 4, 8, 16);
-int2 i2 = (int2) (1, 2);
-long2 l2 = (long2) (1, 2);
+  uint4 ui4 = (uint4) (2, 4, 8, 16);
+  int2 i2 = (int2) (1, 2);
+  long2 l2 = (long2) (1, 2);
 #ifdef cl_khr_fp16
-half2 h2 = (half2) (1, 2);
+  half2 h2 = (half2) (1, 2);
 #endif
-float2 f2 = (float2) (1, 2);
+  float2 f2 = (float2) (1, 2);
 #ifdef cl_khr_fp64
-double2 d2 = (double2) (1, 2);
+  double2 d2 = (double2) (1, 2);
 #endif
 
-__kernel void testkernel (__global int *data)
-{
+  /* marker! */
+
   data[get_global_id(0)] = 1;
 }
diff --git a/gdb/testsuite/gdb.opencl/operators.exp b/gdb/testsuite/gdb.opencl/operators.exp
index e72cb85..3e96719 100644
--- a/gdb/testsuite/gdb.opencl/operators.exp
+++ b/gdb/testsuite/gdb.opencl/operators.exp
@@ -40,15 +40,19 @@ if { [gdb_compile_opencl_hostapp "${clprogram}" "${testfile}" "" ] != "" } {
 clean_restart ${testfile}
 
 # Set breakpoint at the OpenCL kernel
-gdb_test_multiple "break testkernel" "set pending breakpoint" {
-     -re ".*Function \"testkernel\" not defined.*Make breakpoint pending.*y or \\\[n\\\]. $" {
-            gdb_test "y" "Breakpoint.*testkernel.*pending." "set pending breakpoint (without symbols)"
-     }
-}
+gdb_test "break testkernel" \
+    "" \
+    "Set pending breakpoint" \
+    ".*Function \"testkernel\" not defined.*Make breakpoint pending.*y or \\\[n\\\]. $" \
+    "y"
 
 gdb_run_cmd
 gdb_test "" ".*Breakpoint.*1.*testkernel.*" "run"
 
+# Continue to the marker
+gdb_breakpoint [gdb_get_line_number "marker" "${clprogram}"]
+gdb_continue_to_breakpoint "marker"
+
 # Retrieve some information about availability of OpenCL extensions
 set have_cl_khr_fp64 [get_integer_valueof "have_cl_khr_fp64" 0]
 set have_cl_khr_fp16 [get_integer_valueof "have_cl_khr_fp16" 0]
diff --git a/gdb/testsuite/gdb.opencl/vec_comps.cl b/gdb/testsuite/gdb.opencl/vec_comps.cl
index aca72fe..853249c 100644
--- a/gdb/testsuite/gdb.opencl/vec_comps.cl
+++ b/gdb/testsuite/gdb.opencl/vec_comps.cl
@@ -17,43 +17,45 @@
 
    Contributed by Ken Werner <ken.werner@de.ibm.com>  */
 
-int opencl_version = __OPENCL_VERSION__;
+__constant int opencl_version = __OPENCL_VERSION__;
 
 #ifdef HAVE_cl_khr_fp64
 #pragma OPENCL EXTENSION cl_khr_fp64 : enable
-int have_cl_khr_fp64 = 1;
+__constant int have_cl_khr_fp64 = 1;
 #else
-int have_cl_khr_fp64 = 0;
+__constant int have_cl_khr_fp64 = 0;
 #endif
 
 #ifdef HAVE_cl_khr_fp16
 #pragma OPENCL EXTENSION cl_khr_fp16 : enable
-int have_cl_khr_fp16 = 1;
+__constant int have_cl_khr_fp16 = 1;
 #else
-int have_cl_khr_fp16 = 0;
+__constant int have_cl_khr_fp16 = 0;
 #endif
 
+__kernel void testkernel (__global int *data)
+{
 #define CREATE_VEC(TYPE, NAME)\
   TYPE NAME =\
   (TYPE)  (0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12, 13, 14, 15);
 
-CREATE_VEC(char16, c16)
-CREATE_VEC(uchar16, uc16)
-CREATE_VEC(short16, s16)
-CREATE_VEC(ushort16, us16)
-CREATE_VEC(int16, i16)
-CREATE_VEC(uint16, ui16)
-CREATE_VEC(long16, l16)
-CREATE_VEC(ulong16, ul16)
+  CREATE_VEC(char16, c16)
+  CREATE_VEC(uchar16, uc16)
+  CREATE_VEC(short16, s16)
+  CREATE_VEC(ushort16, us16)
+  CREATE_VEC(int16, i16)
+  CREATE_VEC(uint16, ui16)
+  CREATE_VEC(long16, l16)
+  CREATE_VEC(ulong16, ul16)
 #ifdef cl_khr_fp16
-CREATE_VEC(half16, h16)
+  CREATE_VEC(half16, h16)
 #endif
-CREATE_VEC(float16, f16)
+  CREATE_VEC(float16, f16)
 #ifdef cl_khr_fp64
-CREATE_VEC(double16, d16)
+  CREATE_VEC(double16, d16)
 #endif
 
-__kernel void testkernel (__global int *data)
-{
+  /* marker! */
+
   data[get_global_id(0)] = 1;
 }
diff --git a/gdb/testsuite/gdb.opencl/vec_comps.exp b/gdb/testsuite/gdb.opencl/vec_comps.exp
index cb275bc..de537c2 100644
--- a/gdb/testsuite/gdb.opencl/vec_comps.exp
+++ b/gdb/testsuite/gdb.opencl/vec_comps.exp
@@ -40,15 +40,19 @@ if { [gdb_compile_opencl_hostapp "${clprogram}" "${testfile}" "" ] != "" } {
 clean_restart ${testfile}
 
 # Set breakpoint at the OpenCL kernel
-gdb_test_multiple "break testkernel" "set pending breakpoint" {
-     -re ".*Function \"testkernel\" not defined.*Make breakpoint pending.*y or \\\[n\\\]. $" {
-            gdb_test "y" "Breakpoint.*testkernel.*pending." "set pending breakpoint (without symbols)"
-     }
-}
+gdb_test "break testkernel" \
+    "" \
+    "Set pending breakpoint" \
+    ".*Function \"testkernel\" not defined.*Make breakpoint pending.*y or \\\[n\\\]. $" \
+    "y"
 
 gdb_run_cmd
 gdb_test "" ".*Breakpoint.*1.*testkernel.*" "run"
 
+# Continue to the marker
+gdb_breakpoint [gdb_get_line_number "marker" "${clprogram}"]
+gdb_continue_to_breakpoint "marker"
+
 # Check if the language was switched to opencl
 gdb_test "show language" "The current source language is \"auto; currently opencl\"\."