cl: add support for global offset

Signed-off-by: Jan Vesely <jan.vesely@rutgers.edu>
Reviewed-by: Serge Martin <edb+piglit@sigluy.net>

11 files changed, 56 insertions, 15 deletions

diff --git a/tests/cl/api/create-program-with-binary.c b/tests/cl/api/create-program-with-binary.c
index 4873b76fc..423833158 100644
--- a/tests/cl/api/create-program-with-binary.c
+++ b/tests/cl/api/create-program-with-binary.c
@@ -183,7 +183,7 @@ piglit_cl_test(const int argc,
 		size_t global_work_size = 1;
 		size_t local_work_size = 1;
 		cl_command_queue queue = ctx->command_queues[i];
-		if (!piglit_cl_enqueue_ND_range_kernel(queue, kernel, 1,
+		if (!piglit_cl_enqueue_ND_range_kernel(queue, kernel, 1, NULL,
 					&global_work_size, &local_work_size)) {
 			fprintf(stderr, "Failed to execute binary kernel.");
 			piglit_merge_result(&result, PIGLIT_FAIL);
diff --git a/tests/cl/custom/buffer-flags.c b/tests/cl/custom/buffer-flags.c
index 6bf988f0a..115fb9228 100644
--- a/tests/cl/custom/buffer-flags.c
+++ b/tests/cl/custom/buffer-flags.c
@@ -130,7 +130,7 @@ buffer_test(piglit_cl_context *ctx,
 
 	printf("Running the kernel...\n");
 	if (!piglit_cl_enqueue_ND_range_kernel(context->command_queues[0],
-					kernel, 1, &global, &local)) {
+					kernel, 1, NULL, &global, &local)) {
 		ret = PIGLIT_FAIL;
 		goto cleanup;
 	}
diff --git a/tests/cl/custom/flush-after-enqueue-kernel.c b/tests/cl/custom/flush-after-enqueue-kernel.c
index 00a05168e..8d3fc8ccd 100644
--- a/tests/cl/custom/flush-after-enqueue-kernel.c
+++ b/tests/cl/custom/flush-after-enqueue-kernel.c
@@ -92,7 +92,7 @@ piglit_cl_test(const int argc,
 	}
 
 	if (!piglit_cl_enqueue_ND_range_kernel(context->command_queues[0],
-					kernel, 3, global_size, local_size)) {
+					kernel, 3, NULL, global_size, local_size)) {
 		return PIGLIT_FAIL;
 	}
 
diff --git a/tests/cl/custom/r600-create-release-buffer-bug.c b/tests/cl/custom/r600-create-release-buffer-bug.c
index 7c3a214bf..535e57dbf 100644
--- a/tests/cl/custom/r600-create-release-buffer-bug.c
+++ b/tests/cl/custom/r600-create-release-buffer-bug.c
@@ -74,8 +74,8 @@ piglit_cl_test(const int argc,
 		return PIGLIT_FAIL;
 	}
 
-	if (!piglit_cl_enqueue_ND_range_kernel(queue, kernel, 1, &global_size,
-						&local_size)) {
+	if (!piglit_cl_enqueue_ND_range_kernel(queue, kernel, 1, NULL,
+	                                       &global_size, &local_size)) {
 		return PIGLIT_FAIL;
 	}
 
@@ -84,8 +84,8 @@ piglit_cl_test(const int argc,
 		return PIGLIT_FAIL;
 	}
 
-	if (!piglit_cl_enqueue_ND_range_kernel(queue, kernel, 1, &global_size,
-						&local_size)) {
+	if (!piglit_cl_enqueue_ND_range_kernel(queue, kernel, 1, NULL,
+	                                       &global_size, &local_size)) {
 		return PIGLIT_FAIL;
 	}
 
@@ -98,8 +98,8 @@ piglit_cl_test(const int argc,
 		return PIGLIT_FAIL;
 	}
 
-	if (!piglit_cl_enqueue_ND_range_kernel(queue, kernel, 1, &global_size,
-						&local_size)) {
+	if (!piglit_cl_enqueue_ND_range_kernel(queue, kernel, 1, NULL,
+	                                       &global_size, &local_size)) {
 		return PIGLIT_FAIL;
 	}
 
diff --git a/tests/cl/custom/run-simple-kernel.c b/tests/cl/custom/run-simple-kernel.c
index 48d772514..db529d2fe 100644
--- a/tests/cl/custom/run-simple-kernel.c
+++ b/tests/cl/custom/run-simple-kernel.c
@@ -61,7 +61,7 @@ piglit_cl_test(const int argc,
 	                       sizeof(cl_int), &data);
 	piglit_cl_set_kernel_buffer_arg(kernel, 0, &buffer);
 	piglit_cl_execute_ND_range_kernel(context->command_queues[0], kernel, 1,
-	                                  &global_size, &local_size);
+	                                  NULL, &global_size, &local_size);
 
 	/* Read the buffer and check the result */
 	piglit_cl_read_buffer(context->command_queues[0], buffer, 0,
diff --git a/tests/cl/custom/use-sub-buffer-in-kernel.c b/tests/cl/custom/use-sub-buffer-in-kernel.c
index 10cfe0117..74c085882 100644
--- a/tests/cl/custom/use-sub-buffer-in-kernel.c
+++ b/tests/cl/custom/use-sub-buffer-in-kernel.c
@@ -98,8 +98,8 @@ piglit_cl_test(const int argc,
 		return PIGLIT_FAIL;
 	}
 
-	if (!piglit_cl_enqueue_ND_range_kernel(queue, kernel, 1, &global_size,
-						&local_size)) {
+	if (!piglit_cl_enqueue_ND_range_kernel(queue, kernel, 1, NULL,
+	                                       &global_size, &local_size)) {
 		return PIGLIT_FAIL;
 	}
 	clFinish(queue);
diff --git a/tests/cl/program/bitcoin-phatk.c b/tests/cl/program/bitcoin-phatk.c
index 422a3ce8e..f57d41b27 100644
--- a/tests/cl/program/bitcoin-phatk.c
+++ b/tests/cl/program/bitcoin-phatk.c
@@ -514,7 +514,7 @@ piglit_cl_test(const int argc,
 	piglit_cl_set_kernel_buffer_arg(env->kernel, 23, &buffer);
 
 	piglit_cl_execute_ND_range_kernel(env->context->command_queues[0],
-					env->kernel, 1,
+					env->kernel, 1, NULL,
 					&global_size, &local_size);
 	piglit_cl_read_whole_buffer(env->context->command_queues[0], buffer, data);
 
diff --git a/tests/cl/program/max-work-item-sizes.c b/tests/cl/program/max-work-item-sizes.c
index 6fdd18e11..5f160a572 100644
--- a/tests/cl/program/max-work-item-sizes.c
+++ b/tests/cl/program/max-work-item-sizes.c
@@ -116,6 +116,7 @@ piglit_cl_test(const int argc,
 		piglit_cl_execute_ND_range_kernel(env->context->command_queues[0],
 		                                  env->kernel,
 		                                  i+1,
+		                                  NULL,
 		                                  global_size,
 		                                  global_size);
 
diff --git a/tests/cl/program/program-tester.c b/tests/cl/program/program-tester.c
index 0a3b01161..8e3cb4da2 100644
--- a/tests/cl/program/program-tester.c
+++ b/tests/cl/program/program-tester.c
@@ -279,7 +279,9 @@ bool     expect_test_fail = false;
 cl_uint  work_dimensions = 1;
 size_t   global_work_size[3] = {1, 1, 1};
 size_t   local_work_size[3] = {1, 1, 1};
+size_t   global_offset[3] = {0, 0, 0};
 bool     local_work_size_null = false;
+bool     global_offset_null = true;
 
 /* Helper functions */
 
@@ -384,6 +386,8 @@ struct test {
 	size_t global_work_size[3];
 	size_t local_work_size[3];
 	bool local_work_size_null;
+	size_t global_offset[3];
+	bool global_offset_null;
 
 	bool expect_test_fail;
 
@@ -407,6 +411,7 @@ struct test create_test()
 		//.global_work_size = global_work_size,
 		//.local_work_size = local_work_size,
 		.local_work_size_null = local_work_size_null,
+		.global_offset_null = global_offset_null,
 
 		.expect_test_fail = expect_test_fail,
 
@@ -419,6 +424,7 @@ struct test create_test()
 
 	memcpy(t.global_work_size, global_work_size, sizeof(global_work_size));
 	memcpy(t.local_work_size, local_work_size, sizeof(local_work_size));
+	memcpy(t.global_offset, global_offset, sizeof(global_offset));
 
 	return t;
 }
@@ -1736,6 +1742,19 @@ parse_config(const char* config_str,
 					} else {
 						local_work_size_null = true;
 					}
+				} else if(regex_match(key, "^global_offset$")) {
+					if(!regex_match(value, REGEX_FULL_MATCH(REGEX_NULL))) {
+						int i;
+						uint64_t* int_global_offset;
+						get_uint_array(value, &int_global_offset, 3);
+						for(i = 0; i < 3; i++) {
+							global_offset[i] = int_global_offset[i];
+						}
+						global_offset_null = false;
+						free(int_global_offset);
+					} else {
+						global_offset_null = true;
+					}
 				} else {
 					fprintf(stderr,
 					        "Invalid configuration, key '%s' does not belong to a [config] section: %s\n",
@@ -1778,6 +1797,19 @@ parse_config(const char* config_str,
 					} else {
 						test->local_work_size_null = true;
 					}
+				} else if(regex_match(key, "^global_offset$")) {
+					if(!regex_match(value, REGEX_FULL_MATCH(REGEX_NULL))) {
+						int i;
+						uint64_t* int_global_offset;
+						get_uint_array(value, &int_global_offset, 3);
+						for(i = 0; i < 3; i++) {
+							test->global_offset[i] = int_global_offset[i];
+						}
+						test->global_offset_null = false;
+						free(int_global_offset);
+					} else {
+						test->global_offset_null = true;
+					}
 				} else if(regex_match(key, "^arg_in$")) {
 					get_test_arg(value, test, true);
 				} else if(regex_match(key, "^arg_out$")) {
@@ -2396,6 +2428,7 @@ test_kernel(const struct piglit_cl_program_test_config* config,
 	if(!piglit_cl_execute_ND_range_kernel(env->context->command_queues[0],
 	                                      kernel,
 	                                      test.work_dimensions,
+	                                      test.global_offset_null ? NULL : test.global_offset,
 	                                      test.global_work_size,
 	                                      test.local_work_size_null ? NULL : test.local_work_size)) {
 		printf("Failed to enqueue the kernel\n");
diff --git a/tests/util/piglit-util-cl.c b/tests/util/piglit-util-cl.c
index 51d680825..7e7d9850b 100644
--- a/tests/util/piglit-util-cl.c
+++ b/tests/util/piglit-util-cl.c
@@ -1263,14 +1263,15 @@ piglit_cl_set_kernel_buffer_arg(cl_kernel kernel, cl_uint arg_index,
 bool
 piglit_cl_enqueue_ND_range_kernel(cl_command_queue command_queue,
                                   cl_kernel kernel, cl_uint work_dim,
+                                  const size_t* global_offset,
                                   const size_t* global_work_size,
                                   const size_t* local_work_size)
 {
 	cl_int errNo;
 
 	errNo = clEnqueueNDRangeKernel(command_queue, kernel, work_dim,
-	                               NULL, global_work_size, local_work_size,
-	                               0, NULL, NULL);
+	                               global_offset, global_work_size,
+	                               local_work_size, 0, NULL, NULL);
 	if(!piglit_cl_check_error(errNo, CL_SUCCESS)) {
 		fprintf(stderr,
 		        "Could not enqueue ND range kernel: %s\n",
@@ -1284,6 +1285,7 @@ piglit_cl_enqueue_ND_range_kernel(cl_command_queue command_queue,
 bool
 piglit_cl_execute_ND_range_kernel(cl_command_queue command_queue,
                                   cl_kernel kernel, cl_uint work_dim,
+                                  const size_t* global_offset,
                                   const size_t* global_work_size,
                                   const size_t* local_work_size)
 {
@@ -1292,6 +1294,7 @@ piglit_cl_execute_ND_range_kernel(cl_command_queue command_queue,
 	if(!piglit_cl_enqueue_ND_range_kernel(command_queue,
 	                                     kernel,
 	                                     work_dim,
+	                                     global_offset,
 	                                     global_work_size,
 	                                     local_work_size)) {
 		return false;
diff --git a/tests/util/piglit-util-cl.h b/tests/util/piglit-util-cl.h
index e4730ccba..9fb4f9fda 100644
--- a/tests/util/piglit-util-cl.h
+++ b/tests/util/piglit-util-cl.h
@@ -704,6 +704,7 @@ piglit_cl_set_kernel_buffer_arg(cl_kernel kernel,
  * @param command_queue     Command queue to enqueue operation on.
  * @param kernel            Kernel to be enqueued.
  * @param work_dim          Work dimensions.
+ * @param global_offset     Global offset.
  * @param global_work_size  Global work sizes.
  * @param local_work_size   Local work sizes.
  * @return                  \c true on succes, \c false otherwise.
@@ -712,6 +713,7 @@ bool
 piglit_cl_enqueue_ND_range_kernel(cl_command_queue command_queue,
                                   cl_kernel kernel,
                                   cl_uint work_dim,
+                                  const size_t* global_offset,
                                   const size_t* global_work_size,
                                   const size_t* local_work_size);
 
@@ -721,6 +723,7 @@ piglit_cl_enqueue_ND_range_kernel(cl_command_queue command_queue,
  * @param command_queue     Command queue to enqueue operation on.
  * @param kernel            Kernel to be enqueued.
  * @param work_dim          Work dimensions.
+ * @param global_offset     Global offset.
  * @param global_work_size  Global work sizes.
  * @param local_work_size   Local work sizes.
  * @return                  \c true on succes, \c false otherwise.
@@ -729,6 +732,7 @@ bool
 piglit_cl_execute_ND_range_kernel(cl_command_queue command_queue,
                                   cl_kernel kernel,
                                   cl_uint work_dim,
+                                  const size_t* global_offset,
                                   const size_t* global_work_size,
                                   const size_t* local_work_size);