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             size_t numItems = 1;
95             for( size_t j = 0; j < dim; j++ )
96             {
97                 // All of our thread sizes should be within the max local sizes, since they're all <= 20
98                 threads[ j ] = (size_t)random_in_range( 1, 20, d );
99                 localThreads[ j ] = threads[ j ] / (size_t)random_in_range( 1, (int)threads[ j ], d );
100                 while( localThreads[ j ] > 1 && ( threads[ j ] % localThreads[ j ] != 0 ) )
101                     localThreads[ j ]--;
102 
103                 numItems *= threads[ j ];
104 
105                 // Hack for now: localThreads > 1 are iffy
106                 localThreads[ j ] = 1;
107             }
108             error = clEnqueueNDRangeKernel( queue, kernel, (cl_uint)dim, NULL, threads, localThreads, 0, NULL, NULL );
109             test_error( error, "Unable to run kernel" );
110 
111             error = clEnqueueReadBuffer( queue, outData, CL_TRUE, 0, sizeof( testData ), testData, 0, NULL, NULL );
112             test_error( error, "Unable to read results" );
113 
114             // Validate
115             for( size_t q = 0; q < threads[0]; q++ )
116             {
117                 // We can't really validate the actual value of each one, but we can validate that they're within a sane range
118                 if( testData[ q ].workDim != (cl_uint)dim )
119                 {
120                     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 );
121                     free_mtdata(d);
122                     return -1;
123                 }
124                 for( size_t j = 0; j < dim; j++ )
125                 {
126                     if( testData[ q ].globalSize[ j ] != (cl_uint)threads[ j ] )
127                     {
128                         log_error( "ERROR: get_global_size(%d) did not return proper value for %d dimensions (expected %d, got %d)\n",
129                                     (int)j, (int)dim, (int)threads[ j ], (int)testData[ q ].globalSize[ j ] );
130                         free_mtdata(d);
131                         return -1;
132                     }
133                     if (testData[q].globalID[j] >= (cl_uint)threads[j])
134                     {
135                         log_error( "ERROR: get_global_id(%d) did not return proper value for %d dimensions (max %d, got %d)\n",
136                                   (int)j, (int)dim, (int)threads[ j ], (int)testData[ q ].globalID[ j ] );
137                         free_mtdata(d);
138                         return -1;
139                     }
140                     if( testData[ q ].localSize[ j ] != (cl_uint)localThreads[ j ] )
141                     {
142                         log_error( "ERROR: get_local_size(%d) did not return proper value for %d dimensions (expected %d, got %d)\n",
143                                   (int)j, (int)dim, (int)localThreads[ j ], (int)testData[ q ].localSize[ j ] );
144                         free_mtdata(d);
145                         return -1;
146                     }
147                     if (testData[q].localID[j] >= (cl_uint)localThreads[j])
148                     {
149                         log_error( "ERROR: get_local_id(%d) did not return proper value for %d dimensions (max %d, got %d)\n",
150                                   (int)j, (int)dim, (int)localThreads[ j ], (int)testData[ q ].localID[ j ] );
151                         free_mtdata(d);
152                         return -1;
153                     }
154                     size_t groupCount = ( threads[ j ] + localThreads[ j ] - 1 ) / localThreads[ j ];
155                     if( testData[ q ].numGroups[ j ] != (cl_uint)groupCount )
156                     {
157                         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",
158                                   (int)j, (int)dim, (int)groupCount, (int)threads[ j ], (int)localThreads[ j ], (int)testData[ q ].numGroups[ j ] );
159                         free_mtdata(d);
160                         return -1;
161                     }
162                     if (testData[q].groupID[j] >= (cl_uint)groupCount)
163                     {
164                         log_error( "ERROR: get_group_id(%d) did not return proper value for %d dimensions (max %d, got %d)\n",
165                                   (int)j, (int)dim, (int)groupCount, (int)testData[ q ].groupID[ j ] );
166                         free_mtdata(d);
167                         return -1;
168                     }
169                 }
170             }
171         }
172     }
173 
174     free_mtdata(d);
175     return 0;
176 }
177 
178 
179