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
27 const char *constant_source_kernel_code[] = {
28 "__constant int outVal = 42;\n"
29 "__constant int outIndex = 7;\n"
30 "__constant int outValues[ 16 ] = { 17, 01, 11, 12, 1955, 11, 5, 1985, 113, 1, 24, 1984, 7, 23, 1979, 97 };\n"
31 "\n"
32 "__kernel void constant_kernel( __global int *out )\n"
33 "{\n"
34 " int tid = get_global_id(0);\n"
35 "\n"
36 " if( tid == 0 )\n"
37 " {\n"
38 " out[ 0 ] = outVal;\n"
39 " out[ 1 ] = outValues[ outIndex ];\n"
40 " }\n"
41 " else\n"
42 " {\n"
43 " out[ tid + 1 ] = outValues[ tid ];\n"
44 " }\n"
45 "}\n" };
46
test_constant_source(cl_device_id device,cl_context context,cl_command_queue queue,int num_elements)47 int test_constant_source(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
48 {
49 clProgramWrapper program;
50 clKernelWrapper kernel;
51
52 clMemWrapper outStream;
53 cl_int outValues[ 17 ];
54 cl_int expectedValues[ 17 ] = { 42, 1985, 01, 11, 12, 1955, 11, 5, 1985, 113, 1, 24, 1984, 7, 23, 1979, 97 };
55
56 cl_int error;
57
58
59 // Create a kernel to test with
60 error = create_single_kernel_helper( context, &program, &kernel, 1, constant_source_kernel_code, "constant_kernel" );
61 test_error( error, "Unable to create testing kernel" );
62
63 // Create our output buffer
64 outStream = clCreateBuffer( context, CL_MEM_WRITE_ONLY, sizeof( outValues ), NULL, &error );
65 test_error( error, "Unable to create output buffer" );
66
67 // Set the argument
68 error = clSetKernelArg( kernel, 0, sizeof( outStream ), &outStream );
69 test_error( error, "Unable to set kernel argument" );
70
71 // Run test kernel
72 size_t threads[ 1 ] = { 16 };
73 error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, NULL, 0, NULL, NULL );
74 test_error( error, "Unable to enqueue kernel" );
75
76 // Read results
77 error = clEnqueueReadBuffer( queue, outStream, CL_TRUE, 0, sizeof( outValues ), outValues, 0, NULL, NULL );
78 test_error( error, "Unable to read results" );
79
80 // Verify results
81 for( int i = 0; i < 17; i++ )
82 {
83 if( expectedValues[ i ] != outValues[ i ] )
84 {
85 if( i == 0 )
86 log_error( "ERROR: Output value %d from constant source global did not validate! (Expected %d, got %d)\n", i, expectedValues[ i ], outValues[ i ] );
87 else if( i == 1 )
88 log_error( "ERROR: Output value %d from constant-indexed constant array did not validate! (Expected %d, got %d)\n", i, expectedValues[ i ], outValues[ i ] );
89 else
90 log_error( "ERROR: Output value %d from variable-indexed constant array did not validate! (Expected %d, got %d)\n", i, expectedValues[ i ], outValues[ i ] );
91 return -1;
92 }
93 }
94
95 return 0;
96 }
97
98
99
100
101
102