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 <stdlib.h>
20 #include <string.h>
21 #include <sys/types.h>
22 #include <sys/stat.h>
23
24
25 #include "procs.h"
26 #include "harness/conversions.h"
27 #include "harness/typeWrappers.h"
28
29 struct work_item_data
30 {
31 cl_uint workDim;
32 cl_uint globalSize[ 3 ];
33 cl_uint globalID[ 3 ];
34 cl_uint localSize[ 3 ];
35 cl_uint localID[ 3 ];
36 cl_uint numGroups[ 3 ];
37 cl_uint groupID[ 3 ];
38 };
39
40 static const char *workItemKernelCode =
41 "typedef struct {\n"
42 " uint workDim;\n"
43 " uint globalSize[ 3 ];\n"
44 " uint globalID[ 3 ];\n"
45 " uint localSize[ 3 ];\n"
46 " uint localID[ 3 ];\n"
47 " uint numGroups[ 3 ];\n"
48 " uint groupID[ 3 ];\n"
49 " } work_item_data;\n"
50 "\n"
51 "__kernel void sample_kernel( __global work_item_data *outData )\n"
52 "{\n"
53 " int id = get_global_id(0);\n"
54 " outData[ id ].workDim = (uint)get_work_dim();\n"
55 " for( uint i = 0; i < get_work_dim(); i++ )\n"
56 " {\n"
57 " outData[ id ].globalSize[ i ] = (uint)get_global_size( i );\n"
58 " outData[ id ].globalID[ i ] = (uint)get_global_id( i );\n"
59 " outData[ id ].localSize[ i ] = (uint)get_local_size( i );\n"
60 " outData[ id ].localID[ i ] = (uint)get_local_id( i );\n"
61 " outData[ id ].numGroups[ i ] = (uint)get_num_groups( i );\n"
62 " outData[ id ].groupID[ i ] = (uint)get_group_id( i );\n"
63 " }\n"
64 "}";
65
66 #define NUM_TESTS 1
67
test_work_item_functions(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)68 int test_work_item_functions(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
69 {
70 int error;
71
72 clProgramWrapper program;
73 clKernelWrapper kernel;
74 clMemWrapper outData;
75 work_item_data testData[ 10240 ];
76 size_t threads[3], localThreads[3];
77 MTdata d;
78
79
80 error = create_single_kernel_helper( context, &program, &kernel, 1, &workItemKernelCode, "sample_kernel" );
81 test_error( error, "Unable to create testing kernel" );
82
83 outData = clCreateBuffer( context, CL_MEM_READ_WRITE, sizeof( testData ), NULL, &error );
84 test_error( error, "Unable to create output buffer" );
85
86 error = clSetKernelArg( kernel, 0, sizeof( outData ), &outData );
87 test_error( error, "Unable to set kernel arg" );
88
89 d = init_genrand( gRandomSeed );
90 for( size_t dim = 1; dim <= 3; dim++ )
91 {
92 for( int i = 0; i < NUM_TESTS; i++ )
93 {
94 for( size_t j = 0; j < dim; j++ )
95 {
96 // All of our thread sizes should be within the max local sizes, since they're all <= 20
97 threads[ j ] = (size_t)random_in_range( 1, 20, d );
98 localThreads[ j ] = threads[ j ] / (size_t)random_in_range( 1, (int)threads[ j ], d );
99 while( localThreads[ j ] > 1 && ( threads[ j ] % localThreads[ j ] != 0 ) )
100 localThreads[ j ]--;
101
102 // Hack for now: localThreads > 1 are iffy
103 localThreads[ j ] = 1;
104 }
105 error = clEnqueueNDRangeKernel( queue, kernel, (cl_uint)dim, NULL, threads, localThreads, 0, NULL, NULL );
106 test_error( error, "Unable to run kernel" );
107
108 error = clEnqueueReadBuffer( queue, outData, CL_TRUE, 0, sizeof( testData ), testData, 0, NULL, NULL );
109 test_error( error, "Unable to read results" );
110
111 // Validate
112 for( size_t q = 0; q < threads[0]; q++ )
113 {
114 // We can't really validate the actual value of each one, but we can validate that they're within a sane range
115 if( testData[ q ].workDim != (cl_uint)dim )
116 {
117 log_error( "ERROR: get_work_dim() did not return proper value for %d dimensions (expected %d, got %d)\n", (int)dim, (int)dim, (int)testData[ q ].workDim );
118 free_mtdata(d);
119 return -1;
120 }
121 for( size_t j = 0; j < dim; j++ )
122 {
123 if( testData[ q ].globalSize[ j ] != (cl_uint)threads[ j ] )
124 {
125 log_error( "ERROR: get_global_size(%d) did not return proper value for %d dimensions (expected %d, got %d)\n",
126 (int)j, (int)dim, (int)threads[ j ], (int)testData[ q ].globalSize[ j ] );
127 free_mtdata(d);
128 return -1;
129 }
130 if (testData[q].globalID[j] >= (cl_uint)threads[j])
131 {
132 log_error( "ERROR: get_global_id(%d) did not return proper value for %d dimensions (max %d, got %d)\n",
133 (int)j, (int)dim, (int)threads[ j ], (int)testData[ q ].globalID[ j ] );
134 free_mtdata(d);
135 return -1;
136 }
137 if( testData[ q ].localSize[ j ] != (cl_uint)localThreads[ j ] )
138 {
139 log_error( "ERROR: get_local_size(%d) did not return proper value for %d dimensions (expected %d, got %d)\n",
140 (int)j, (int)dim, (int)localThreads[ j ], (int)testData[ q ].localSize[ j ] );
141 free_mtdata(d);
142 return -1;
143 }
144 if (testData[q].localID[j] >= (cl_uint)localThreads[j])
145 {
146 log_error( "ERROR: get_local_id(%d) did not return proper value for %d dimensions (max %d, got %d)\n",
147 (int)j, (int)dim, (int)localThreads[ j ], (int)testData[ q ].localID[ j ] );
148 free_mtdata(d);
149 return -1;
150 }
151 size_t groupCount = ( threads[ j ] + localThreads[ j ] - 1 ) / localThreads[ j ];
152 if( testData[ q ].numGroups[ j ] != (cl_uint)groupCount )
153 {
154 log_error( "ERROR: get_num_groups(%d) did not return proper value for %d dimensions (expected %d with global dim %d and local dim %d, got %d)\n",
155 (int)j, (int)dim, (int)groupCount, (int)threads[ j ], (int)localThreads[ j ], (int)testData[ q ].numGroups[ j ] );
156 free_mtdata(d);
157 return -1;
158 }
159 if (testData[q].groupID[j] >= (cl_uint)groupCount)
160 {
161 log_error( "ERROR: get_group_id(%d) did not return proper value for %d dimensions (max %d, got %d)\n",
162 (int)j, (int)dim, (int)groupCount, (int)testData[ q ].groupID[ j ] );
163 free_mtdata(d);
164 return -1;
165 }
166 }
167 }
168 }
169 }
170
171 free_mtdata(d);
172 return 0;
173 }
174
175
176