| 67 | } |
| 68 | |
| 69 | static int test_barrier_common(cl_device_id device, cl_context context, |
| 70 | cl_command_queue queue, int num_elements, |
| 71 | std::string barrier_str) |
| 72 | { |
| 73 | clMemWrapper streams[3]; |
| 74 | clProgramWrapper program; |
| 75 | clKernelWrapper kernel; |
| 76 | |
| 77 | cl_int output; |
| 78 | int err; |
| 79 | |
| 80 | size_t max_threadgroup_size = 0; |
| 81 | std::string build_options = std::string("-DBARRIER=") + barrier_str; |
| 82 | err = create_single_kernel_helper(context, &program, &kernel, 1, |
| 83 | &barrier_kernel_code, "compute_sum", |
| 84 | build_options.c_str()); |
| 85 | test_error(err, "Failed to build kernel/program."); |
| 86 | |
| 87 | err = get_max_allowed_1d_work_group_size_on_device(device, kernel, |
| 88 | &max_threadgroup_size); |
| 89 | test_error(err, "get_max_allowed_1d_work_group_size_on_device failed."); |
| 90 | |
| 91 | // work group size must divide evenly into the global size |
| 92 | while (num_elements % max_threadgroup_size) max_threadgroup_size--; |
| 93 | |
| 94 | std::vector<cl_int> input(num_elements); |
| 95 | |
| 96 | streams[0] = clCreateBuffer(context, CL_MEM_READ_WRITE, |
| 97 | sizeof(cl_int) * num_elements, nullptr, &err); |
| 98 | test_error(err, "clCreateBuffer failed."); |
| 99 | streams[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, sizeof(cl_int), |
| 100 | nullptr, &err); |
| 101 | test_error(err, "clCreateBuffer failed."); |
| 102 | streams[2] = |
| 103 | clCreateBuffer(context, CL_MEM_READ_WRITE, |
| 104 | sizeof(cl_int) * max_threadgroup_size, nullptr, &err); |
| 105 | test_error(err, "clCreateBuffer failed."); |
| 106 | |
| 107 | generate_random_inputs(input); |
| 108 | |
| 109 | err = clEnqueueWriteBuffer(queue, streams[0], CL_TRUE, 0, |
| 110 | sizeof(cl_int) * num_elements, input.data(), 0, |
| 111 | nullptr, nullptr); |
| 112 | test_error(err, "clEnqueueWriteBuffer failed."); |
| 113 | |
| 114 | err = clSetKernelArg(kernel, 0, sizeof(streams[0]), &streams[0]); |
| 115 | err |= clSetKernelArg(kernel, 1, sizeof(num_elements), &num_elements); |
| 116 | err |= clSetKernelArg(kernel, 2, sizeof(streams[2]), &streams[2]); |
| 117 | err |= clSetKernelArg(kernel, 3, sizeof(streams[1]), &streams[1]); |
| 118 | test_error(err, "clSetKernelArg failed."); |
| 119 | |
| 120 | size_t global_threads[] = { max_threadgroup_size }; |
| 121 | size_t local_threads[] = { max_threadgroup_size }; |
| 122 | |
| 123 | err = clEnqueueNDRangeKernel(queue, kernel, 1, nullptr, global_threads, |
| 124 | local_threads, 0, nullptr, nullptr); |
| 125 | test_error(err, "clEnqueueNDRangeKernel failed."); |
| 126 |
no test coverage detected