| // |
| // 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 = checkCorrectnessStep(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. |
| } |
| |
| static 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"); |
| } |
