// // 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 "testBase.h" #include "harness/conversions.h" #include "harness/typeWrappers.h" #include "harness/testHarness.h" #include "structs.h" #include "defines.h" #include "type_replacer.h" /* test_step_type, test_step_var, test_step_typedef_type, test_step_typedef_var, */ int test_step_internal(cl_device_id deviceID, cl_context context, cl_command_queue queue, const char * pattern, const char * testName) { int err; int typeIdx, vecSizeIdx; char tempBuffer[2048]; clState * pClState = newClState(deviceID, context, queue); bufferStruct * pBuffers = newBufferStruct(BUFFER_SIZE, BUFFER_SIZE, pClState); if(pBuffers == NULL) { destroyClState(pClState); vlog_error("%s : Could not create buffer\n", testName); return -1; } //detect whether profile of the device is embedded char profile[1024] = ""; err = clGetDeviceInfo(deviceID, CL_DEVICE_PROFILE, sizeof(profile), profile, NULL); if (err) { print_error(err, "clGetDeviceInfo for CL_DEVICE_PROFILE failed\n" ); return -1; } gIsEmbedded = NULL != strstr(profile, "EMBEDDED_PROFILE"); for(typeIdx = 0; types[typeIdx] != kNumExplicitTypes; ++typeIdx) { if( types[ typeIdx ] == kDouble ) { // If we're testing doubles, we need to check for support first if( !is_extension_available( deviceID, "cl_khr_fp64" ) ) { log_info( "Not testing doubles (unsupported on this device)\n" ); continue; } } if( types[ typeIdx ] == kLong || types[ typeIdx ] == kULong ) { // If we're testing long/ulong, we need to check for embedded support if( gIsEmbedded && !is_extension_available( deviceID, "cles_khr_int64") ) { log_info( "Not testing longs (unsupported on this embedded device)\n" ); continue; } } char srcBuffer[2048]; doSingleReplace(tempBuffer, 2048, pattern, ".EXTENSIONS.", types[typeIdx] == kDouble ? "#pragma OPENCL EXTENSION cl_khr_fp64 : enable" : ""); for(vecSizeIdx = 0; vecSizeIdx < NUM_VECTOR_SIZES; ++vecSizeIdx) { doReplace(srcBuffer, 2048, tempBuffer, ".TYPE.", g_arrTypeNames[typeIdx], ".NUM.", g_arrVecSizeNames[vecSizeIdx]); if(srcBuffer[0] == '\0') { vlog_error("%s: failed to fill source buf for type %s%s\n", testName, g_arrTypeNames[typeIdx], g_arrVecSizeNames[vecSizeIdx]); destroyBufferStruct(pBuffers, pClState); destroyClState(pClState); return -1; } err = clStateMakeProgram(pClState, srcBuffer, testName ); if (err) { vlog_error("%s: Error compiling \"\n%s\n\"", testName, srcBuffer); destroyBufferStruct(pBuffers, pClState); destroyClState(pClState); return -1; } err = pushArgs(pBuffers, pClState); if(err != 0) { vlog_error("%s: failed to push args %s%s\n", testName, g_arrTypeNames[typeIdx], g_arrVecSizeNames[vecSizeIdx]); destroyBufferStruct(pBuffers, pClState); destroyClState(pClState); return -1; } // now we run the kernel err = runKernel(pClState, 1024); if(err != 0) { vlog_error("%s: runKernel fail (%ld threads) %s%s\n", testName, pClState->m_numThreads, g_arrTypeNames[typeIdx], g_arrVecSizeNames[vecSizeIdx]); destroyBufferStruct(pBuffers, pClState); destroyClState(pClState); return -1; } err = retrieveResults(pBuffers, pClState); if(err != 0) { vlog_error("%s: failed to retrieve results %s%s\n", testName, g_arrTypeNames[typeIdx], g_arrVecSizeNames[vecSizeIdx]); destroyBufferStruct(pBuffers, pClState); destroyClState(pClState); return -1; } err = checkCorrectness(pBuffers, pClState, g_arrTypeSizes[typeIdx], g_arrVecSizes[vecSizeIdx]); if(err != 0) { vlog_error("%s: incorrect results %s%s\n", testName, g_arrTypeNames[typeIdx], g_arrVecSizeNames[vecSizeIdx]); vlog_error("%s: Source was \"\n%s\n\"", testName, srcBuffer); destroyBufferStruct(pBuffers, pClState); destroyClState(pClState); return -1; } } } destroyBufferStruct(pBuffers, pClState); destroyClState(pClState); // vlog_error("%s : implementation incomplete : FAIL\n", testName); return 0; // -1; // fails on account of not being written. } const char * patterns[] = { ".EXTENSIONS.\n" "__kernel void test_step_type(__global .TYPE..NUM. *source, __global int *dest)\n" "{\n" " int tid = get_global_id(0);\n" " dest[tid] = vec_step(.TYPE..NUM.);\n" "\n" "}\n", ".EXTENSIONS.\n" "__kernel void test_step_var(__global .TYPE..NUM. *source, __global int *dest)\n" "{\n" " int tid = get_global_id(0);\n" " dest[tid] = vec_step(source[tid]);\n" "\n" "}\n", ".EXTENSIONS.\n" " typedef .TYPE..NUM. TypeToTest;\n" "__kernel void test_step_typedef_type(__global TypeToTest *source, __global int *dest)\n" "{\n" " int tid = get_global_id(0);\n" " dest[tid] = vec_step(TypeToTest);\n" "\n" "}\n", ".EXTENSIONS.\n" " typedef .TYPE..NUM. TypeToTest;\n" "__kernel void test_step_typedef_var(__global TypeToTest *source, __global int *dest)\n" "{\n" " int tid = get_global_id(0);\n" " dest[tid] = vec_step(source[tid]);\n" "\n" "}\n", }; /* test_step_type, test_step_var, test_step_typedef_type, test_step_typedef_var, */ int test_step_type(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { return test_step_internal(deviceID, context, queue, patterns[0], "test_step_type"); } int test_step_var(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { return test_step_internal(deviceID, context, queue, patterns[1], "test_step_var"); } int test_step_typedef_type(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { return test_step_internal(deviceID, context, queue, patterns[2], "test_step_typedef_type"); } int test_step_typedef_var(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { return test_step_internal(deviceID, context, queue, patterns[3], "test_step_typedef_var"); }