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 "testBase.h"
17 #include "harness/conversions.h"
18
19 #define TEST_SIZE 512
20
21 enum OpKonstants
22 {
23 kIncrement = 0,
24 kDecrement,
25 kBoth
26 };
27
28 const char *testKernel =
29 "__kernel void test( __global %s *inOut, __global char * control )\n"
30 "{\n"
31 " size_t tid = get_global_id(0);\n"
32 "\n"
33 " %s%s inOutVal = %s;\n"
34 "\n"
35 " if( control[tid] == 0 )\n"
36 " inOutVal++;\n"
37 " else if( control[tid] == 1 )\n"
38 " ++inOutVal;\n"
39 " else if( control[tid] == 2 )\n"
40 " inOutVal--;\n"
41 " else // if( control[tid] == 3 )\n"
42 " --inOutVal;\n"
43 "\n"
44 " %s;\n"
45 "}\n";
46
47 typedef int (*OpVerifyFn)( void * actualPtr, void * inputPtr, size_t vecSize, size_t numVecs, cl_char * controls );
48
test_unary_op(cl_command_queue queue,cl_context context,OpKonstants whichOp,ExplicitType vecType,size_t vecSize,MTdata d,OpVerifyFn verifyFn)49 int test_unary_op( cl_command_queue queue, cl_context context, OpKonstants whichOp,
50 ExplicitType vecType, size_t vecSize,
51 MTdata d, OpVerifyFn verifyFn )
52 {
53 clProgramWrapper program;
54 clKernelWrapper kernel;
55 clMemWrapper streams[2];
56 cl_long inData[TEST_SIZE * 16], outData[TEST_SIZE * 16];
57 cl_char controlData[TEST_SIZE];
58 int error;
59 size_t i;
60 size_t threads[1], localThreads[1];
61 char kernelSource[10240];
62 char *programPtr;
63
64
65 // Create the source
66 char loadLine[ 1024 ], storeLine[ 1024 ];
67 if( vecSize == 1 )
68 {
69 sprintf( loadLine, "inOut[tid]" );
70 sprintf( storeLine, "inOut[tid] = inOutVal" );
71 }
72 else
73 {
74 sprintf( loadLine, "vload%ld( tid, inOut )", vecSize );
75 sprintf( storeLine, "vstore%ld( inOutVal, tid, inOut )", vecSize );
76 }
77
78 char sizeNames[][4] = { "", "", "2", "3", "4", "", "", "", "8", "", "", "", "", "", "", "", "16" };
79 sprintf( kernelSource, testKernel, get_explicit_type_name( vecType ), /*sizeNames[ vecSize ],*/
80 get_explicit_type_name( vecType ), sizeNames[ vecSize ],
81 loadLine, storeLine );
82
83 // Create the kernel
84 programPtr = kernelSource;
85 if( create_single_kernel_helper( context, &program, &kernel, 1, (const char **)&programPtr, "test" ) )
86 {
87 log_error( "ERROR: Unable to create test program!\n" );
88 return -1;
89 }
90
91 // Generate two streams. The first is our random data to test against, the second is our control stream
92 generate_random_data( vecType, vecSize * TEST_SIZE, d, inData );
93 streams[0] = clCreateBuffer(
94 context, CL_MEM_COPY_HOST_PTR,
95 get_explicit_type_size(vecType) * vecSize * TEST_SIZE, inData, &error);
96 test_error( error, "Creating input data array failed" );
97
98 cl_uint bits;
99 for( i = 0; i < TEST_SIZE; i++ )
100 {
101 size_t which = i & 7;
102 if( which == 0 )
103 bits = genrand_int32(d);
104
105 controlData[ i ] = ( bits >> ( which << 1 ) ) & 0x03;
106 if( whichOp == kDecrement )
107 // For sub ops, the min control value is 2. Otherwise, it's 0
108 controlData[ i ] |= 0x02;
109 else if( whichOp == kIncrement )
110 // For addition ops, the max control value is 1. Otherwise, it's 3
111 controlData[ i ] &= ~0x02;
112 }
113 streams[1] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
114 sizeof(controlData), controlData, &error);
115 test_error( error, "Unable to create control stream" );
116
117 // Assign streams and execute
118 error = clSetKernelArg( kernel, 0, sizeof( streams[0] ), &streams[0] );
119 test_error( error, "Unable to set indexed kernel arguments" );
120 error = clSetKernelArg( kernel, 1, sizeof( streams[1] ), &streams[1] );
121 test_error( error, "Unable to set indexed kernel arguments" );
122
123
124 // Run the kernel
125 threads[0] = TEST_SIZE;
126
127 error = get_max_common_work_group_size( context, kernel, threads[0], &localThreads[0] );
128 test_error( error, "Unable to get work group size to use" );
129
130 error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, localThreads, 0, NULL, NULL );
131 test_error( error, "Unable to execute test kernel" );
132
133
134 // Read the results
135 error = clEnqueueReadBuffer( queue, streams[0], CL_TRUE, 0,
136 get_explicit_type_size( vecType ) * TEST_SIZE * vecSize,
137 outData, 0, NULL, NULL );
138 test_error( error, "Unable to read output array!" );
139
140 // Now verify the results
141 return verifyFn( outData, inData, vecSize, TEST_SIZE, controlData );
142 }
143
VerifyFn(void * actualPtr,void * inputPtr,size_t vecSize,size_t numVecs,cl_char * controls)144 template<typename T> int VerifyFn( void * actualPtr, void * inputPtr, size_t vecSize, size_t numVecs, cl_char * controls )
145 {
146 T * actualData = (T *)actualPtr;
147 T * inputData = (T *)inputPtr;
148
149 size_t index = 0;
150 for( size_t i = 0; i < numVecs; i++ )
151 {
152 for( size_t j = 0; j < vecSize; j++, index++ )
153 {
154 T nextVal = inputData[ index ];
155 if( controls[ i ] & 0x02 )
156 nextVal--;
157 else
158 nextVal++;
159
160 if( actualData[ index ] != nextVal )
161 {
162 log_error( "ERROR: Validation failed on vector %ld:%ld (expected %lld, got %lld)", i, j,
163 (cl_long)nextVal, (cl_long)actualData[ index ] );
164 return -1;
165 }
166 }
167 }
168 return 0;
169 }
170
test_unary_op_set(cl_command_queue queue,cl_context context,OpKonstants whichOp)171 int test_unary_op_set( cl_command_queue queue, cl_context context, OpKonstants whichOp )
172 {
173 ExplicitType types[] = { kChar, kUChar, kShort, kUShort, kInt, kUInt, kLong, kULong, kNumExplicitTypes };
174 OpVerifyFn verifys[] = { VerifyFn<cl_char>, VerifyFn<cl_uchar>, VerifyFn<cl_short>, VerifyFn<cl_ushort>, VerifyFn<cl_int>, VerifyFn<cl_uint>, VerifyFn<cl_long>, VerifyFn<cl_ulong>, NULL };
175 unsigned int vecSizes[] = { 1, 2, 3, 4, 8, 16, 0 };
176 unsigned int index, typeIndex;
177 int retVal = 0;
178 RandomSeed seed(gRandomSeed );
179
180 for( typeIndex = 0; types[ typeIndex ] != kNumExplicitTypes; typeIndex++ )
181 {
182 if ((types[ typeIndex ] == kLong || types[ typeIndex ] == kULong) && !gHasLong)
183 continue;
184
185 for( index = 0; vecSizes[ index ] != 0; index++ )
186 {
187 if( test_unary_op( queue, context, whichOp, types[ typeIndex ], vecSizes[ index ], seed, verifys[ typeIndex ] ) != 0 )
188 {
189 log_error( " Vector %s%d FAILED\n", get_explicit_type_name( types[ typeIndex ] ), vecSizes[ index ] );
190 retVal = -1;
191 }
192 }
193 }
194
195 return retVal;
196 }
197
test_unary_ops_full(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)198 int test_unary_ops_full(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
199 {
200 return test_unary_op_set( queue, context, kBoth );
201 }
202
test_unary_ops_increment(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)203 int test_unary_ops_increment(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
204 {
205 return test_unary_op_set( queue, context, kIncrement );
206 }
207
test_unary_ops_decrement(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)208 int test_unary_ops_decrement(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
209 {
210 return test_unary_op_set( queue, context, kDecrement );
211 }
212