| // |
| // Copyright (c) 2017 The Khronos Group Inc. |
| // |
| // Licensed under the Apache License, Version 2.0 (the "License"); |
| // you may not use this file except in compliance with the License. |
| // You may obtain a copy of the License at |
| // |
| // http://www.apache.org/licenses/LICENSE-2.0 |
| // |
| // Unless required by applicable law or agreed to in writing, software |
| // distributed under the License is distributed on an "AS IS" BASIS, |
| // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. |
| // See the License for the specific language governing permissions and |
| // limitations under the License. |
| // |
| #include "harness/compat.h" |
| |
| #include <stdio.h> |
| #include <stdlib.h> |
| #include <string.h> |
| #include <sys/types.h> |
| #include <sys/stat.h> |
| #include "harness/rounding_mode.h" |
| |
| #include "procs.h" |
| |
| static const char *global_linear_id_2d_code = |
| "__kernel void test_global_linear_id_2d(global int *dst)\n" |
| "{\n" |
| " int tid_x = get_global_id(0);\n" |
| " int tid_y = get_global_id(1);\n" |
| "\n" |
| " int linear_id = tid_y * get_global_size(0) + tid_x;\n" |
| " int result = (linear_id == (int)get_global_linear_id()) ? 0x1 : 0x0;\n" |
| " dst[linear_id] = result;\n" |
| "}\n"; |
| |
| static const char *global_linear_id_1d_code = |
| "__kernel void test_global_linear_id_1d(global int *dst)\n" |
| "{\n" |
| " int tid_x = get_global_id(0);\n" |
| "\n" |
| " int result = (tid_x == (int)get_global_linear_id()) ? 0x1 : 0x0;\n" |
| " dst[tid_x] = result;\n" |
| "}\n"; |
| |
| |
| static int |
| verify_global_linear_id(int *result, int n) |
| { |
| int i; |
| for (i=0; i<n; i++) |
| { |
| if (result[i] == 0) |
| { |
| log_error("get_global_linear_id failed\n"); |
| return -1; |
| } |
| } |
| log_info("get_global_linear_id passed\n"); |
| return 0; |
| } |
| |
| |
| int |
| test_global_linear_id(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements) |
| { |
| cl_mem streams; |
| cl_program program[2]; |
| cl_kernel kernel[2]; |
| |
| int *output_ptr; |
| size_t threads[2]; |
| int err; |
| num_elements = (int)sqrt((float)num_elements); |
| int length = num_elements * num_elements; |
| |
| output_ptr = (int *)malloc(sizeof(int) * length); |
| |
| streams = clCreateBuffer(context, CL_MEM_READ_WRITE, length * sizeof(int), |
| NULL, &err); |
| test_error( err, "clCreateBuffer failed."); |
| |
| err = create_single_kernel_helper(context, &program[0], &kernel[0], 1, |
| &global_linear_id_1d_code, |
| "test_global_linear_id_1d"); |
| test_error( err, "create_single_kernel_helper failed"); |
| err = create_single_kernel_helper(context, &program[1], &kernel[1], 1, |
| &global_linear_id_2d_code, |
| "test_global_linear_id_2d"); |
| test_error( err, "create_single_kernel_helper failed"); |
| |
| err = clSetKernelArg(kernel[0], 0, sizeof streams, &streams); |
| test_error( err, "clSetKernelArgs failed."); |
| err = clSetKernelArg(kernel[1], 0, sizeof streams, &streams); |
| test_error( err, "clSetKernelArgs failed."); |
| |
| threads[0] = (size_t)num_elements; |
| threads[1] = (size_t)num_elements; |
| err = clEnqueueNDRangeKernel(queue, kernel[1], 2, NULL, threads, NULL, 0, NULL, NULL); |
| test_error( err, "clEnqueueNDRangeKernel failed."); |
| |
| err = clEnqueueReadBuffer(queue, streams, CL_TRUE, 0, length*sizeof(int), output_ptr, 0, NULL, NULL); |
| test_error( err, "clEnqueueReadBuffer failed."); |
| |
| err = verify_global_linear_id(output_ptr, length); |
| |
| threads[0] = (size_t)num_elements; |
| err = clEnqueueNDRangeKernel(queue, kernel[1], 1, NULL, threads, NULL, 0, NULL, NULL); |
| test_error( err, "clEnqueueNDRangeKernel failed."); |
| |
| err = clEnqueueReadBuffer(queue, streams, CL_TRUE, 0, num_elements*sizeof(int), output_ptr, 0, NULL, NULL); |
| test_error( err, "clEnqueueReadBuffer failed."); |
| |
| err = verify_global_linear_id(output_ptr, num_elements); |
| |
| // cleanup |
| clReleaseMemObject(streams); |
| clReleaseKernel(kernel[0]); |
| clReleaseKernel(kernel[1]); |
| clReleaseProgram(program[0]); |
| clReleaseProgram(program[1]); |
| free(output_ptr); |
| |
| return err; |
| } |