You can not select more than 25 topics
Topics must start with a letter or number, can include dashes ('-') and can be up to 35 characters long.
102 lines
3.3 KiB
102 lines
3.3 KiB
//
|
|
// Copyright (c) 2017 The Khronos Group Inc.
|
|
//
|
|
// Licensed under the Apache License, Version 2.0 (the "License");
|
|
// you may not use this file except in compliance with the License.
|
|
// You may obtain a copy of the License at
|
|
//
|
|
// http://www.apache.org/licenses/LICENSE-2.0
|
|
//
|
|
// Unless required by applicable law or agreed to in writing, software
|
|
// distributed under the License is distributed on an "AS IS" BASIS,
|
|
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
|
// See the License for the specific language governing permissions and
|
|
// limitations under the License.
|
|
//
|
|
#include "harness/compat.h"
|
|
|
|
#include <stdio.h>
|
|
#include <stdlib.h>
|
|
#include <string.h>
|
|
#include <sys/types.h>
|
|
#include <sys/stat.h>
|
|
|
|
|
|
#include "procs.h"
|
|
|
|
const char *constant_source_kernel_code[] = {
|
|
"__constant int outVal = 42;\n"
|
|
"__constant int outIndex = 7;\n"
|
|
"__constant int outValues[ 16 ] = { 17, 01, 11, 12, 1955, 11, 5, 1985, 113, 1, 24, 1984, 7, 23, 1979, 97 };\n"
|
|
"\n"
|
|
"__kernel void constant_kernel( __global int *out )\n"
|
|
"{\n"
|
|
" int tid = get_global_id(0);\n"
|
|
"\n"
|
|
" if( tid == 0 )\n"
|
|
" {\n"
|
|
" out[ 0 ] = outVal;\n"
|
|
" out[ 1 ] = outValues[ outIndex ];\n"
|
|
" }\n"
|
|
" else\n"
|
|
" {\n"
|
|
" out[ tid + 1 ] = outValues[ tid ];\n"
|
|
" }\n"
|
|
"}\n" };
|
|
|
|
int test_constant_source(cl_device_id device, cl_context context, cl_command_queue queue, int num_elements)
|
|
{
|
|
clProgramWrapper program;
|
|
clKernelWrapper kernel;
|
|
|
|
clMemWrapper outStream;
|
|
cl_int outValues[ 17 ];
|
|
cl_int expectedValues[ 17 ] = { 42, 1985, 01, 11, 12, 1955, 11, 5, 1985, 113, 1, 24, 1984, 7, 23, 1979, 97 };
|
|
|
|
cl_int error;
|
|
|
|
|
|
// Create a kernel to test with
|
|
error = create_single_kernel_helper( context, &program, &kernel, 1, constant_source_kernel_code, "constant_kernel" );
|
|
test_error( error, "Unable to create testing kernel" );
|
|
|
|
// Create our output buffer
|
|
outStream = clCreateBuffer( context, CL_MEM_WRITE_ONLY, sizeof( outValues ), NULL, &error );
|
|
test_error( error, "Unable to create output buffer" );
|
|
|
|
// Set the argument
|
|
error = clSetKernelArg( kernel, 0, sizeof( outStream ), &outStream );
|
|
test_error( error, "Unable to set kernel argument" );
|
|
|
|
// Run test kernel
|
|
size_t threads[ 1 ] = { 16 };
|
|
error = clEnqueueNDRangeKernel( queue, kernel, 1, NULL, threads, NULL, 0, NULL, NULL );
|
|
test_error( error, "Unable to enqueue kernel" );
|
|
|
|
// Read results
|
|
error = clEnqueueReadBuffer( queue, outStream, CL_TRUE, 0, sizeof( outValues ), outValues, 0, NULL, NULL );
|
|
test_error( error, "Unable to read results" );
|
|
|
|
// Verify results
|
|
for( int i = 0; i < 17; i++ )
|
|
{
|
|
if( expectedValues[ i ] != outValues[ i ] )
|
|
{
|
|
if( i == 0 )
|
|
log_error( "ERROR: Output value %d from constant source global did not validate! (Expected %d, got %d)\n", i, expectedValues[ i ], outValues[ i ] );
|
|
else if( i == 1 )
|
|
log_error( "ERROR: Output value %d from constant-indexed constant array did not validate! (Expected %d, got %d)\n", i, expectedValues[ i ], outValues[ i ] );
|
|
else
|
|
log_error( "ERROR: Output value %d from variable-indexed constant array did not validate! (Expected %d, got %d)\n", i, expectedValues[ i ], outValues[ i ] );
|
|
return -1;
|
|
}
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
|
|
|
|
|
|
|