| // |
| // 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" |
| |
| |
| size_t get_align(size_t vecSize) |
| { |
| if (vecSize == 3) |
| { |
| return 4; |
| } |
| return vecSize; |
| } |
| |
| /* // Lots of conditionals means this is not gonna be an optimal min on intel. |
| */ |
| /* // That's okay, make sure we only call a few times per test, not for every */ |
| /* // element */ |
| /* size_t min_of_nonzero(size_t a, size_t b) */ |
| /* { */ |
| /* if(a != 0 && (a<=b || b==0)) */ |
| /* { */ |
| /* return a; */ |
| /* } */ |
| /* if(b != 0 && (b<a || a==0)) */ |
| /* { */ |
| /* return b; */ |
| /* } */ |
| /* return 0; */ |
| /* } */ |
| |
| |
| /* size_t get_min_packed_alignment(size_t preSize, size_t typeMultiplePreSize, |
| */ |
| /* size_t postSize, size_t typeMultiplePostSize, */ |
| /* ExplicitType kType, size_t vecSize) */ |
| /* { */ |
| /* size_t pre_min = min_of_nonzero(preSize, */ |
| /* typeMultiplePreSize* */ |
| /* get_explicit_type_size(kType)); */ |
| /* size_t post_min = min_of_nonzero(postSize, */ |
| /* typeMultiplePostSize* */ |
| /* get_explicit_type_size(kType)); */ |
| /* size_t struct_min = min_of_nonzero(pre_min, post_min); */ |
| /* size_t result = min_of_nonzero(struct_min, get_align(vecSize) */ |
| /* *get_explicit_type_size(kType)); */ |
| /* return result; */ |
| |
| /* } */ |
| |
| |
| int test_vec_internal(cl_device_id deviceID, cl_context context, |
| cl_command_queue queue, const char* pattern, |
| const char* testName, size_t bufSize, size_t preSize, |
| size_t typeMultiplePreSize, size_t postSize, |
| size_t typeMultiplePostSize) |
| { |
| int err; |
| int typeIdx, vecSizeIdx; |
| |
| char tmpBuffer[2048]; |
| char srcBuffer[2048]; |
| |
| size_t preSizeBytes, postSizeBytes, typeSize, totSize; |
| |
| clState* pClState = newClState(deviceID, context, queue); |
| bufferStruct* pBuffers = newBufferStruct( |
| bufSize, bufSize * sizeof(cl_uint) / sizeof(cl_char), pClState); |
| |
| if (pBuffers == NULL) |
| { |
| destroyClState(pClState); |
| vlog_error("%s : Could not create buffer\n", testName); |
| return -1; |
| } |
| |
| for (typeIdx = 0; types[typeIdx] != kNumExplicitTypes; ++typeIdx) |
| { |
| |
| // Skip doubles if it is not supported otherwise enable pragma |
| if (types[typeIdx] == kDouble) |
| { |
| if (!is_extension_available(deviceID, "cl_khr_fp64")) |
| { |
| continue; |
| } |
| else |
| { |
| doReplace(tmpBuffer, 2048, pattern, ".PRAGMA.", |
| "#pragma OPENCL EXTENSION cl_khr_fp64: ", ".STATE.", |
| "enable"); |
| } |
| } |
| else |
| { |
| if (types[typeIdx] == kLong || types[typeIdx] == kULong) |
| { |
| if (gIsEmbedded) continue; |
| } |
| |
| doReplace(tmpBuffer, 2048, pattern, ".PRAGMA.", " ", ".STATE.", |
| " "); |
| } |
| |
| typeSize = get_explicit_type_size(types[typeIdx]); |
| preSizeBytes = preSize + typeSize * typeMultiplePreSize; |
| postSizeBytes = postSize + typeSize * typeMultiplePostSize; |
| |
| |
| for (vecSizeIdx = 1; vecSizeIdx < NUM_VECTOR_SIZES; ++vecSizeIdx) |
| { |
| |
| totSize = preSizeBytes + postSizeBytes |
| + typeSize * get_align(g_arrVecSizes[vecSizeIdx]); |
| |
| doReplace(srcBuffer, 2048, tmpBuffer, ".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; |
| } |
| |
| // log_info("Buffer is \"\n%s\n\"\n", srcBuffer); |
| // fflush(stdout); |
| |
| 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; |
| } |
| |
| // log_info("About to Run kernel\n"); fflush(stdout); |
| // now we run the kernel |
| err = runKernel( |
| pClState, |
| bufSize |
| / (g_arrVecSizes[vecSizeIdx] * g_arrTypeSizes[typeIdx])); |
| 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; |
| } |
| |
| // log_info("About to retrieve results\n"); fflush(stdout); |
| 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; |
| } |
| |
| |
| if (preSizeBytes + postSizeBytes == 0) |
| { |
| // log_info("About to Check Correctness\n"); fflush(stdout); |
| err = checkCorrectnessAlign(pBuffers, pClState, |
| get_align(g_arrVecSizes[vecSizeIdx]) |
| * typeSize); |
| } |
| else |
| { |
| // we're checking for an aligned struct |
| err = checkPackedCorrectness(pBuffers, pClState, totSize, |
| preSizeBytes); |
| } |
| |
| 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; |
| } |
| |
| clStateDestroyProgramAndKernel(pClState); |
| } |
| } |
| |
| 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[] = { |
| ".PRAGMA..STATE.\n" |
| "__kernel void test_vec_align_array(.SRC_SCOPE. .TYPE..NUM. *source, " |
| ".DST_SCOPE. uint *dest)\n" |
| "{\n" |
| " int tid = get_global_id(0);\n" |
| " dest[tid] = (uint)((.SRC_SCOPE. uchar *)(source+tid));\n" |
| "}\n", |
| ".PRAGMA..STATE.\n" |
| "typedef struct myUnpackedStruct { \n" |
| ".PRE." |
| " .TYPE..NUM. vec;\n" |
| ".POST." |
| "} testStruct;\n" |
| "__kernel void test_vec_align_struct(__constant .TYPE..NUM. *source, " |
| ".DST_SCOPE. uint *dest)\n" |
| "{\n" |
| " .SRC_SCOPE. testStruct test;\n" |
| " int tid = get_global_id(0);\n" |
| " dest[tid] = (uint)((.SRC_SCOPE. uchar *)&(test.vec));\n" |
| "}\n", |
| ".PRAGMA..STATE.\n" |
| "typedef struct __attribute__ ((packed)) myPackedStruct { \n" |
| ".PRE." |
| " .TYPE..NUM. vec;\n" |
| ".POST." |
| "} testStruct;\n" |
| "__kernel void test_vec_align_packed_struct(__constant .TYPE..NUM. " |
| "*source, .DST_SCOPE. uint *dest)\n" |
| "{\n" |
| " .SRC_SCOPE. testStruct test;\n" |
| " int tid = get_global_id(0);\n" |
| " dest[tid] = (uint)((.SRC_SCOPE. uchar *)&(test.vec) - (.SRC_SCOPE. " |
| "uchar *)&test);\n" |
| "}\n", |
| ".PRAGMA..STATE.\n" |
| "typedef struct myStruct { \n" |
| ".PRE." |
| " .TYPE..NUM. vec;\n" |
| ".POST." |
| "} testStruct;\n" |
| "__kernel void test_vec_align_struct_arr(.SRC_SCOPE. testStruct *source, " |
| ".DST_SCOPE. uint *dest)\n" |
| "{\n" |
| " int tid = get_global_id(0);\n" |
| " dest[tid] = (uint)((.SRC_SCOPE. uchar *)&(source[tid].vec));\n" |
| "}\n", |
| ".PRAGMA..STATE.\n" |
| "typedef struct __attribute__ ((packed)) myPackedStruct { \n" |
| ".PRE." |
| " .TYPE..NUM. vec;\n" |
| ".POST." |
| "} testStruct;\n" |
| "__kernel void test_vec_align_packed_struct_arr(.SRC_SCOPE. testStruct " |
| "*source, .DST_SCOPE. uint *dest)\n" |
| "{\n" |
| " int tid = get_global_id(0);\n" |
| " dest[tid] = (uint)((.SRC_SCOPE. uchar *)&(source[tid].vec) - " |
| "(.SRC_SCOPE. uchar *)&(source[0]));\n" |
| "}\n", |
| // __attribute__ ((packed)) |
| }; |
| |
| |
| const char* pre_substitution_arr[] = { "", |
| "char c;\n", |
| "short3 s;", |
| ".TYPE.3 tPre;\n", |
| ".TYPE. arrPre[5];\n", |
| ".TYPE. arrPre[12];\n", |
| NULL }; |
| |
| |
| // alignments of everything in pre_substitution_arr as raw alignments |
| // 0 if such a thing is meaningless |
| size_t pre_align_arr[] = { 0, |
| sizeof(cl_char), |
| 4 * sizeof(cl_short), |
| 0, // taken care of in type_multiple_pre_align_arr |
| 0, |
| 0 }; |
| |
| // alignments of everything in pre_substitution_arr as multiples of |
| // sizeof(.TYPE.) |
| // 0 if such a thing is meaningless |
| size_t type_multiple_pre_align_arr[] = { 0, 0, 0, 4, 5, 12 }; |
| |
| const char* post_substitution_arr[] = { "", |
| "char cPost;\n", |
| ".TYPE. arrPost[3];\n", |
| ".TYPE. arrPost[5];\n", |
| ".TYPE.3 arrPost;\n", |
| ".TYPE. arrPost[12];\n", |
| NULL }; |
| |
| |
| // alignments of everything in post_substitution_arr as raw alignments |
| // 0 if such a thing is meaningless |
| size_t post_align_arr[] = { 0, sizeof(cl_char), |
| 0, // taken care of in type_multiple_post_align_arr |
| 0, 0, |
| 0 }; |
| |
| // alignments of everything in post_substitution_arr as multiples of |
| // sizeof(.TYPE.) |
| // 0 if such a thing is meaningless |
| size_t type_multiple_post_align_arr[] = { 0, 0, 3, 5, 4, 12 }; |
| |
| // there hsould be a packed version of this? |
| int test_vec_align_array(cl_device_id deviceID, cl_context context, |
| cl_command_queue queue, int num_elements) |
| { |
| char tmp[2048]; |
| int result; |
| |
| log_info("Testing global\n"); |
| doReplace(tmp, (size_t)2048, patterns[0], ".SRC_SCOPE.", "__global", |
| ".DST_SCOPE.", "__global"); // |
| result = test_vec_internal(deviceID, context, queue, tmp, |
| "test_vec_align_array", BUFFER_SIZE, 0, 0, 0, 0); |
| return result; |
| } |
| |
| |
| int test_vec_align_struct(cl_device_id deviceID, cl_context context, |
| cl_command_queue queue, int num_elements) |
| { |
| char tmp1[2048], tmp2[2048]; |
| int result = 0; |
| int preIdx, postIdx; |
| |
| log_info("testing __private\n"); |
| doReplace(tmp2, (size_t)2048, patterns[1], ".SRC_SCOPE.", "__private", |
| ".DST_SCOPE.", "__global"); // |
| |
| for (preIdx = 0; pre_substitution_arr[preIdx] != NULL; ++preIdx) |
| { |
| for (postIdx = 0; post_substitution_arr[postIdx] != NULL; ++postIdx) |
| { |
| doReplace(tmp1, (size_t)2048, tmp2, ".PRE.", |
| pre_substitution_arr[preIdx], ".POST.", |
| post_substitution_arr[postIdx]); |
| |
| result = |
| test_vec_internal(deviceID, context, queue, tmp1, |
| "test_vec_align_struct", 512, 0, 0, 0, 0); |
| if (result != 0) |
| { |
| return result; |
| } |
| } |
| } |
| |
| log_info("testing __local\n"); |
| doReplace(tmp2, (size_t)2048, patterns[1], ".SRC_SCOPE.", "__local", |
| ".DST_SCOPE.", "__global"); // |
| |
| for (preIdx = 0; pre_substitution_arr[preIdx] != NULL; ++preIdx) |
| { |
| for (postIdx = 0; post_substitution_arr[postIdx] != NULL; ++postIdx) |
| { |
| doReplace(tmp1, (size_t)2048, tmp2, ".PRE.", |
| pre_substitution_arr[preIdx], ".POST.", |
| post_substitution_arr[postIdx]); |
| |
| result = |
| test_vec_internal(deviceID, context, queue, tmp1, |
| "test_vec_align_struct", 512, 0, 0, 0, 0); |
| if (result != 0) |
| { |
| return result; |
| } |
| } |
| } |
| return 0; |
| } |
| |
| int test_vec_align_packed_struct(cl_device_id deviceID, cl_context context, |
| cl_command_queue queue, int num_elements) |
| { |
| char tmp1[2048], tmp2[2048]; |
| int result = 0; |
| int preIdx, postIdx; |
| |
| |
| log_info("Testing __private\n"); |
| doReplace(tmp2, (size_t)2048, patterns[2], ".SRC_SCOPE.", "__private", |
| ".DST_SCOPE.", "__global"); // |
| |
| for (preIdx = 0; pre_substitution_arr[preIdx] != NULL; ++preIdx) |
| { |
| for (postIdx = 0; post_substitution_arr[postIdx] != NULL; ++postIdx) |
| { |
| doReplace(tmp1, (size_t)2048, tmp2, ".PRE.", |
| pre_substitution_arr[preIdx], ".POST.", |
| post_substitution_arr[postIdx]); |
| |
| result = test_vec_internal( |
| deviceID, context, queue, tmp1, "test_vec_align_packed_struct", |
| 512, pre_align_arr[preIdx], type_multiple_pre_align_arr[preIdx], |
| post_align_arr[postIdx], type_multiple_post_align_arr[postIdx]); |
| if (result != 0) |
| { |
| return result; |
| } |
| } |
| } |
| |
| log_info("testing __local\n"); |
| doReplace(tmp2, (size_t)2048, patterns[2], ".SRC_SCOPE.", "__local", |
| ".DST_SCOPE.", "__global"); // |
| |
| for (preIdx = 0; pre_substitution_arr[preIdx] != NULL; ++preIdx) |
| { |
| for (postIdx = 0; post_substitution_arr[postIdx] != NULL; ++postIdx) |
| { |
| doReplace(tmp1, (size_t)2048, tmp2, ".PRE.", |
| pre_substitution_arr[preIdx], ".POST.", |
| post_substitution_arr[postIdx]); |
| |
| result = test_vec_internal( |
| deviceID, context, queue, tmp1, "test_vec_align_packed_struct", |
| 512, pre_align_arr[preIdx], type_multiple_pre_align_arr[preIdx], |
| post_align_arr[postIdx], type_multiple_post_align_arr[postIdx]); |
| if (result != 0) |
| { |
| return result; |
| } |
| } |
| } |
| return 0; |
| } |
| |
| int test_vec_align_struct_arr(cl_device_id deviceID, cl_context context, |
| cl_command_queue queue, int num_elements) |
| { |
| char tmp1[2048], tmp2[2048]; |
| int result = 0; |
| int preIdx, postIdx; |
| |
| |
| log_info("testing __global\n"); |
| doReplace(tmp2, (size_t)2048, patterns[3], ".SRC_SCOPE.", "__global", |
| ".DST_SCOPE.", "__global"); // |
| |
| for (preIdx = 0; pre_substitution_arr[preIdx] != NULL; ++preIdx) |
| { |
| for (postIdx = 0; post_substitution_arr[postIdx] != NULL; ++postIdx) |
| { |
| doReplace(tmp1, (size_t)2048, tmp2, ".PRE.", |
| pre_substitution_arr[preIdx], ".POST.", |
| post_substitution_arr[postIdx]); |
| |
| result = test_vec_internal(deviceID, context, queue, tmp1, |
| "test_vec_align_struct_arr", BUFFER_SIZE, |
| 0, 0, 0, 0); |
| if (result != 0) |
| { |
| return result; |
| } |
| } |
| } |
| return 0; |
| } |
| |
| int test_vec_align_packed_struct_arr(cl_device_id deviceID, cl_context context, |
| cl_command_queue queue, int num_elements) |
| { |
| char tmp1[2048], tmp2[2048]; |
| int result = 0; |
| int preIdx, postIdx; |
| |
| |
| log_info("Testing __global\n"); |
| doReplace(tmp2, (size_t)2048, patterns[4], ".SRC_SCOPE.", "__global", |
| ".DST_SCOPE.", "__global"); // |
| |
| for (preIdx = 0; pre_substitution_arr[preIdx] != NULL; ++preIdx) |
| { |
| for (postIdx = 0; post_substitution_arr[postIdx] != NULL; ++postIdx) |
| { |
| doReplace(tmp1, (size_t)2048, tmp2, ".PRE.", |
| pre_substitution_arr[preIdx], ".POST.", |
| post_substitution_arr[postIdx]); |
| |
| result = test_vec_internal( |
| deviceID, context, queue, tmp1, |
| "test_vec_align_packed_struct_arr", BUFFER_SIZE, |
| pre_align_arr[preIdx], type_multiple_pre_align_arr[preIdx], |
| post_align_arr[postIdx], type_multiple_post_align_arr[postIdx]); |
| if (result != 0) return result; |
| } |
| } |
| return 0; |
| } |
