1 //
2 // Copyright (c) 2017 The Khronos Group Inc.
3 //
4 // Licensed under the Apache License, Version 2.0 (the "License");
5 // you may not use this file except in compliance with the License.
6 // You may obtain a copy of the License at
7 //
8 // http://www.apache.org/licenses/LICENSE-2.0
9 //
10 // Unless required by applicable law or agreed to in writing, software
11 // distributed under the License is distributed on an "AS IS" BASIS,
12 // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13 // See the License for the specific language governing permissions and
14 // limitations under the License.
15 //
16 #include "harness/compat.h"
17
18 #include <stdio.h>
19 #include <string.h>
20 #include <sys/types.h>
21 #include <sys/stat.h>
22
23 #include "procs.h"
24 #include "harness/errorHelpers.h"
25
26 const char* pipe_readwrite_errors_kernel_code = {
27 "__kernel void test_pipe_write_error(__global int *src, __write_only pipe int out_pipe, __global int *status)\n"
28 "{\n"
29 " int gid = get_global_id(0);\n"
30 " int reserve_idx;\n"
31 " reserve_id_t res_id;\n"
32 "\n"
33 " res_id = reserve_write_pipe(out_pipe, 1);\n"
34 " if(is_valid_reserve_id(res_id))\n"
35 " {\n"
36 " write_pipe(out_pipe, res_id, 0, &src[gid]);\n"
37 " commit_write_pipe(out_pipe, res_id);\n"
38 " }\n"
39 " else\n"
40 " {\n"
41 " *status = -1;\n"
42 " }\n"
43 "}\n"
44 "\n"
45 "__kernel void test_pipe_read_error(__read_only pipe int in_pipe, __global int *dst, __global int *status)\n"
46 "{\n"
47 " int gid = get_global_id(0);\n"
48 " int reserve_idx;\n"
49 " reserve_id_t res_id;\n"
50 "\n"
51 " res_id = reserve_read_pipe(in_pipe, 1);\n"
52 " if(is_valid_reserve_id(res_id))\n"
53 " {\n"
54 " read_pipe(in_pipe, res_id, 0, &dst[gid]);\n"
55 " commit_read_pipe(in_pipe, res_id);\n"
56 " }\n"
57 " else\n"
58 " {\n"
59 " *status = -1;\n"
60 " }\n"
61 "}\n"
62 };
63
64
test_pipe_readwrite_errors(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)65 int test_pipe_readwrite_errors(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
66 {
67 clMemWrapper pipe;
68 clMemWrapper buffers[3];
69 void *outptr;
70 cl_int *inptr;
71 clProgramWrapper program;
72 clKernelWrapper kernel[2];
73 size_t global_work_size[3];
74 cl_int err;
75 cl_int size;
76 cl_int i;
77 cl_int status = 0;
78 clEventWrapper producer_sync_event;
79 clEventWrapper consumer_sync_event;
80 BufferOwningPtr<cl_int> BufferInPtr;
81 BufferOwningPtr<cl_int> BufferOutPtr;
82 MTdataHolder d(gRandomSeed);
83 const char *kernelName[] = { "test_pipe_write_error",
84 "test_pipe_read_error" };
85
86 size_t min_alignment = get_min_alignment(context);
87
88 global_work_size[0] = num_elements;
89
90 size = num_elements * sizeof(cl_int);
91
92 inptr = (cl_int *)align_malloc(size, min_alignment);
93
94 for (i = 0; i < num_elements; i++)
95 {
96 inptr[i] = (int)genrand_int32(d);
97 }
98 BufferInPtr.reset(inptr, nullptr, 0, size, true);
99
100 buffers[0] =
101 clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, size, inptr, &err);
102 test_error_ret(err, " clCreateBuffer failed", -1);
103
104 outptr = align_malloc(size, min_alignment);
105 BufferOutPtr.reset(outptr, nullptr, 0, size, true);
106
107 buffers[1] = clCreateBuffer(context, CL_MEM_USE_HOST_PTR, size, outptr, &err);
108 test_error_ret(err, " clCreateBuffer failed", -1);
109
110 buffers[2] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR, sizeof(int), &status, &err);
111 test_error_ret(err, " clCreateBuffer failed", -1);
112
113 //Pipe created with max_packets less than global size
114 pipe = clCreatePipe(context, CL_MEM_HOST_NO_ACCESS, sizeof(int), num_elements - (num_elements/2), NULL, &err);
115 test_error_ret(err, " clCreatePipe failed", -1);
116
117 // Create producer kernel
118 err = create_single_kernel_helper(
119 context, &program, &kernel[0], 1,
120 (const char **)&pipe_readwrite_errors_kernel_code, kernelName[0]);
121 test_error_ret(err, " Error creating program", -1);
122
123 //Create consumer kernel
124 kernel[1] = clCreateKernel(program, kernelName[1], &err);
125 test_error_ret(err, " Error creating kernel", -1);
126
127 err = clSetKernelArg(kernel[0], 0, sizeof(cl_mem), (void*)&buffers[0]);
128 err |= clSetKernelArg(kernel[0], 1, sizeof(cl_mem), (void*)&pipe);
129 err |= clSetKernelArg(kernel[0], 2, sizeof(cl_mem), (void*)&buffers[2]);
130 err |= clSetKernelArg(kernel[1], 0, sizeof(cl_mem), (void*)&pipe);
131 err |= clSetKernelArg(kernel[1], 1, sizeof(cl_mem), (void*)&buffers[1]);
132 err |= clSetKernelArg(kernel[1], 2, sizeof(cl_mem), (void*)&buffers[2]);
133
134 test_error_ret(err, " clSetKernelArg failed", -1);
135
136 // Launch Consumer kernel for empty pipe
137 err = clEnqueueNDRangeKernel( queue, kernel[1], 1, NULL, global_work_size, NULL, 0, NULL, &consumer_sync_event );
138 test_error_ret(err, " clEnqueueNDRangeKernel failed", -1);
139
140 err = clEnqueueReadBuffer(queue, buffers[2], true, 0, sizeof(status), &status, 1, &consumer_sync_event, NULL);
141 test_error_ret(err, " clEnqueueReadBuffer failed", -1);
142
143 if(status == 0){
144 log_error("test_pipe_readwrite_errors failed\n");
145 return -1;
146 }
147 else{
148 status = 0;
149 }
150
151 // Launch Producer kernel
152 err = clEnqueueNDRangeKernel( queue, kernel[0], 1, NULL, global_work_size, NULL, 0, NULL, &producer_sync_event );
153 test_error_ret(err, " clEnqueueNDRangeKernel failed", -1);
154
155 err = clEnqueueReadBuffer(queue, buffers[2], true, 0, sizeof(status), &status, 1, &producer_sync_event, NULL);
156 test_error_ret(err, " clEnqueueReadBuffer failed", -1);
157
158 if (status == 0)
159 {
160 log_error("test_pipe_readwrite_errors failed\n");
161 return -1;
162 }
163 else{
164 status = 0;
165 }
166
167 // We will reuse this variable so release the previous referred event.
168 clReleaseEvent(consumer_sync_event);
169
170 // Launch Consumer kernel
171 err = clEnqueueNDRangeKernel( queue, kernel[1], 1, NULL, global_work_size, NULL, 1, &producer_sync_event, &consumer_sync_event );
172 test_error_ret(err, " clEnqueueNDRangeKernel failed", -1);
173
174 err = clEnqueueReadBuffer(queue, buffers[2], true, 0, sizeof(status), &status, 1, &consumer_sync_event, NULL);
175 test_error_ret(err, " clEnqueueReadBuffer failed", -1);
176
177 if (status == 0)
178 {
179 log_error("test_pipe_readwrite_errors failed\n");
180 return -1;
181 }
182 else
183 {
184 status = 0;
185 }
186
187 log_info("test_pipe_readwrite_errors passed\n");
188
189 return 0;
190 }
191