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/testHarness.h" 17 18 #include <string> 19 20 class CTest { 21 public: 22 virtual int Execute(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) = 0; 23 }; 24 25 #define NL "\n" 26 27 namespace common { 28 static const std::string CONFORMANCE_VERIFY_FENCE = 29 NL 30 NL "// current spec says get_fence can return any valid fence" 31 NL "bool isFenceValid(cl_mem_fence_flags fence) {" 32 NL " if ((fence == 0) || (fence == CLK_GLOBAL_MEM_FENCE) || (fence == CLK_LOCAL_MEM_FENCE) " 33 NL " || (fence == (CLK_GLOBAL_MEM_FENCE | CLK_LOCAL_MEM_FENCE)))" 34 NL " return true;" 35 NL " else" 36 NL " return false;" 37 NL "}" 38 NL; 39 40 static std::string GLOBAL_KERNEL_FUNCTION = CONFORMANCE_VERIFY_FENCE + 41 NL 42 NL "bool helperFunction(uint *ptr, uint tid) {" 43 NL " if (!isFenceValid(get_fence(ptr)))" 44 NL " return false;" 45 NL 46 NL " if (*ptr != tid)" 47 NL " return false;" 48 NL 49 NL " return true;" 50 NL "}" 51 NL 52 NL "__kernel void testKernel(__global uint *results, __global uint *buf) {" 53 NL " uint tid = get_global_id(0);" 54 NL 55 NL " results[tid] = helperFunction(&buf[tid], tid);" 56 NL "}" 57 NL; 58 59 static std::string LOCAL_KERNEL_FUNCTION = CONFORMANCE_VERIFY_FENCE + 60 NL 61 NL "bool helperFunction(uint *ptr, uint tid) {" 62 NL " if (!isFenceValid(get_fence(ptr)))" 63 NL " return false;" 64 NL 65 NL " if (*ptr != tid)" 66 NL " return false;" 67 NL 68 NL " return true;" 69 NL "}" 70 NL 71 NL "__kernel void testKernel(__global uint *results, __local uint *buf) {" 72 NL " uint tid = get_global_id(0);" 73 NL " if (get_local_id(0) == 0) {" 74 NL " for (uint i = 0; i < get_local_size(0); ++i) {" 75 NL " uint idx = get_local_size(0) * get_group_id(0) + i;" 76 NL " buf[idx] = idx;" 77 NL " }" 78 NL " }" 79 NL 80 NL " work_group_barrier(CLK_LOCAL_MEM_FENCE);" 81 NL " results[tid] = helperFunction(&buf[tid], tid);" 82 NL "}" 83 NL; 84 } 85