| // |
| // 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/testHarness.h" |
| #include "harness/typeWrappers.h" |
| #include "base.h" |
| |
| #include <string> |
| #include <vector> |
| #include <algorithm> |
| |
| class CBasicTest : CTest { |
| public: |
| CBasicTest(const std::vector<std::string>& kernel) : CTest(), _kernels(kernel) { |
| |
| } |
| |
| CBasicTest(const std::string& kernel) : CTest(), _kernels(1, kernel) { |
| |
| } |
| |
| int ExecuteSubcase(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements, const std::string& src) { |
| cl_int error; |
| |
| clProgramWrapper program; |
| clKernelWrapper kernel; |
| |
| const char *srcPtr = src.c_str(); |
| |
| if (create_single_kernel_helper_with_build_options(context, &program, &kernel, 1, &srcPtr, "testKernel", "-cl-std=CL2.0")) { |
| log_error("create_single_kernel_helper failed"); |
| return -1; |
| } |
| |
| size_t bufferSize = num_elements * sizeof(cl_uint); |
| clMemWrapper buffer = clCreateBuffer(context, CL_MEM_WRITE_ONLY, bufferSize, NULL, &error); |
| test_error(error, "clCreateBuffer failed"); |
| |
| error = clSetKernelArg(kernel, 0, sizeof(buffer), &buffer); |
| test_error(error, "clSetKernelArg failed"); |
| |
| size_t globalWorkGroupSize = num_elements; |
| size_t localWorkGroupSize = 0; |
| error = get_max_common_work_group_size(context, kernel, globalWorkGroupSize, &localWorkGroupSize); |
| test_error(error, "Unable to get common work group size"); |
| |
| error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &globalWorkGroupSize, &localWorkGroupSize, 0, NULL, NULL); |
| test_error(error, "clEnqueueNDRangeKernel failed"); |
| |
| // verify results |
| std::vector<cl_uint> results(num_elements); |
| |
| error = clEnqueueReadBuffer(queue, buffer, CL_TRUE, 0, bufferSize, &results[0], 0, NULL, NULL); |
| test_error(error, "clEnqueueReadBuffer failed"); |
| |
| size_t passCount = std::count(results.begin(), results.end(), 1); |
| if (passCount != results.size()) { |
| std::vector<cl_uint>::iterator iter = std::find(results.begin(), results.end(), 0); |
| log_error("Verification on device failed at index %ld\n", std::distance(results.begin(), iter)); |
| log_error("%ld out of %ld failed\n", (results.size()-passCount), results.size()); |
| return -1; |
| } |
| |
| return CL_SUCCESS; |
| } |
| |
| int Execute(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { |
| cl_int result = CL_SUCCESS; |
| |
| for (std::vector<std::string>::const_iterator it = _kernels.begin(); it != _kernels.end(); ++it) { |
| log_info("Executing subcase #%ld out of %ld\n", (it - _kernels.begin() + 1), _kernels.size()); |
| |
| result |= ExecuteSubcase(deviceID, context, queue, num_elements, *it); |
| } |
| |
| return result; |
| } |
| |
| private: |
| const std::vector<std::string> _kernels; |
| }; |
| |
| int test_function_get_fence(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { |
| const std::string KERNEL_FUNCTION = common::CONFORMANCE_VERIFY_FENCE + |
| NL |
| NL "__global int gint = 1;" |
| NL "__global uchar guchar = 3;" |
| NL |
| NL "bool helperFunction(int *intp, float *floatp, uchar *ucharp, ushort *ushortp, long *longp) {" |
| NL " if (!isFenceValid(get_fence(intp)))" |
| NL " return false;" |
| NL " if (!isFenceValid(get_fence(floatp)))" |
| NL " return false;" |
| NL " if (!isFenceValid(get_fence(ucharp)))" |
| NL " return false;" |
| NL " if (!isFenceValid(get_fence(ushortp)))" |
| NL " return false;" |
| NL " if (!isFenceValid(get_fence(longp)))" |
| NL " return false;" |
| NL |
| NL " if (*intp != 1 || *floatp != 2.0f || *ucharp != 3 || *ushortp != 4 || *longp != 5)" |
| NL " return false;" |
| NL |
| NL " return true;" |
| NL "}" |
| NL |
| NL "__kernel void testKernel(__global uint *results) {" |
| NL " uint tid = get_global_id(0);" |
| NL |
| NL " __local float lfloat;" |
| NL " lfloat = 2.0f;" |
| NL " __local ushort lushort;" |
| NL " lushort = 4;" |
| NL " long plong = 5;" |
| NL |
| NL " __global int *gintp = &gint;" |
| NL " __local float *lfloatp = &lfloat;" |
| NL " __global uchar *gucharp = &guchar;" |
| NL " __local ushort *lushortp = &lushort;" |
| NL " __private long *plongp = &plong;" |
| NL |
| NL " results[tid] = helperFunction(gintp, lfloatp, gucharp, lushortp, plongp);" |
| NL "}" |
| NL; |
| |
| CBasicTest test(KERNEL_FUNCTION); |
| |
| return test.Execute(deviceID, context, queue, num_elements); |
| } |
| |
| int test_function_to_address_space(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { |
| const std::string KERNEL_FUNCTION = |
| NL |
| NL "__global int gint = 1;" |
| NL "__global uchar guchar = 3;" |
| NL |
| NL "bool helperFunction(int *gintp, float *lfloatp, uchar *gucharp, ushort *lushortp, long *plongp) {" |
| NL " if (to_global(gintp) == NULL)" |
| NL " return false;" |
| NL " if (to_local(lfloatp) == NULL)" |
| NL " return false;" |
| NL " if (to_global(gucharp) == NULL)" |
| NL " return false;" |
| NL " if (to_local(lushortp) == NULL)" |
| NL " return false;" |
| NL " if (to_private(plongp) == NULL)" |
| NL " return false;" |
| NL |
| NL " if (*gintp != 1 || *lfloatp != 2.0f || *gucharp != 3 || *lushortp != 4 || *plongp != 5)" |
| NL " return false;" |
| NL |
| NL " return true;" |
| NL "}" |
| NL |
| NL "__kernel void testKernel(__global uint *results) {" |
| NL " uint tid = get_global_id(0);" |
| NL |
| NL " __local float lfloat;" |
| NL " lfloat = 2.0f;" |
| NL " __local ushort lushort;" |
| NL " lushort = 4;" |
| NL " long plong = 5;" |
| NL |
| NL " __global int *gintp = &gint;" |
| NL " __local float *lfloatp = &lfloat;" |
| NL " __global uchar *gucharp = &guchar;" |
| NL " __local ushort *lushortp = &lushort;" |
| NL " __private long *plongp = &plong;" |
| NL |
| NL " results[tid] = helperFunction(gintp, lfloatp, gucharp, lushortp, plongp);" |
| NL "}" |
| NL; |
| |
| CBasicTest test(KERNEL_FUNCTION); |
| |
| return test.Execute(deviceID, context, queue, num_elements); |
| } |
| |
| int test_variable_get_fence(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { |
| const std::string KERNEL_FUNCTION = common::CONFORMANCE_VERIFY_FENCE + |
| NL |
| NL "__global int gint = 1;" |
| NL |
| NL "__kernel void testKernel(__global uint *results) {" |
| NL " uint tid = get_global_id(0);" |
| NL |
| NL " __local ushort lushort;" |
| NL " lushort = 2;" |
| NL " float pfloat = 3.0f;" |
| NL |
| NL " // tested pointers" |
| NL " __global int *gintp = &gint;" |
| NL " __local ushort *lushortp = &lushort;" |
| NL " __private float *pfloatp = &pfloat;" |
| NL |
| NL " int failures = 0;" |
| NL " if (!isFenceValid(get_fence(gintp)))" |
| NL " failures++;" |
| NL " if (!isFenceValid(get_fence(lushortp)))" |
| NL " failures++;" |
| NL " if (!isFenceValid(get_fence(pfloatp)))" |
| NL " failures++;" |
| NL " results[tid] = (failures == 0);" |
| NL "}" |
| NL; |
| |
| CBasicTest test(KERNEL_FUNCTION); |
| |
| return test.Execute(deviceID, context, queue, num_elements); |
| } |
| |
| int test_variable_to_address_space(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { |
| const std::string KERNEL_FUNCTION = |
| NL |
| NL "__global int gint = 1;" |
| NL |
| NL "__kernel void testKernel(__global uint *results) {" |
| NL " uint tid = get_global_id(0);" |
| NL |
| NL " __local ushort lushort;" |
| NL " lushort = 2;" |
| NL " float pfloat = 3.0f;" |
| NL |
| NL " // tested pointers" |
| NL " __global int * gintp = &gint;" |
| NL " __local ushort *lushortp = &lushort;" |
| NL " __private float *pfloatp = &pfloat;" |
| NL |
| NL " int failures = 0;" |
| NL " if (to_global(gintp) == NULL)" |
| NL " failures++;" |
| NL " if (to_local(lushortp) == NULL)" |
| NL " failures++;" |
| NL " if (to_private(pfloatp) == NULL)" |
| NL " failures++;" |
| NL " results[tid] = (failures == 0);" |
| NL "}" |
| NL; |
| |
| CBasicTest test(KERNEL_FUNCTION); |
| |
| return test.Execute(deviceID, context, queue, num_elements); |
| } |
| |
| int test_casting(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { |
| std::vector<std::string> KERNEL_FUNCTIONS; |
| |
| // pointers to global, local or private are implicitly convertible to generic |
| KERNEL_FUNCTIONS.push_back(common::CONFORMANCE_VERIFY_FENCE + |
| NL |
| NL "__global int gint = 1;" |
| NL |
| NL "__kernel void testKernel(__global uint *results) {" |
| NL " uint tid = get_global_id(0);" |
| NL |
| NL " __local int lint;" |
| NL " lint = 2;" |
| NL " int pint = 3;" |
| NL |
| NL " // count mismatches with expected fence types" |
| NL " int failures = 0;" |
| NL |
| NL " // tested pointer" |
| NL " // generic can be reassigned to different named address spaces" |
| NL " int * intp;" |
| NL |
| NL " intp = &gint;" |
| NL " failures += !(isFenceValid(get_fence(intp)));" |
| NL " failures += !(to_global(intp));" |
| NL " failures += (*intp != 1);" |
| NL |
| NL " intp = &lint;" |
| NL " failures += !(isFenceValid(get_fence(intp)));" |
| NL " failures += !(to_local(intp));" |
| NL " failures += (*intp != 2);" |
| NL |
| NL " intp = &pint;" |
| NL " failures += !(isFenceValid(get_fence(intp)));" |
| NL " failures += !(to_private(intp));" |
| NL " failures += (*intp != 3);" |
| NL |
| NL " results[tid] = (failures == 0);" |
| NL "}" |
| NL |
| ); |
| |
| // converting from a generic pointer to a named address space is legal only with explicit casting |
| KERNEL_FUNCTIONS.push_back(common::CONFORMANCE_VERIFY_FENCE + |
| NL |
| NL "__global int gint = 1;" |
| NL |
| NL "__kernel void testKernel(__global uint *results) {" |
| NL " uint tid = get_global_id(0);" |
| NL |
| NL " __local int lint;" |
| NL " lint = 2;" |
| NL " int pint = 3;" |
| NL |
| NL " // count mismatches with expected fence types" |
| NL " int failures = 0;" |
| NL |
| NL " // tested pointer" |
| NL " // generic can be reassigned to different named address spaces" |
| NL " int * intp;" |
| NL |
| NL " intp = &gint;" |
| NL " global int * gintp = (global int *)intp;" |
| NL " failures += !(isFenceValid(get_fence(gintp)));" |
| NL " failures += !(to_global(gintp));" |
| NL " failures += (*gintp != 1);" |
| NL |
| NL " intp = &lint;" |
| NL " local int * lintp = (local int *)intp;" |
| NL " failures += !(isFenceValid(get_fence(lintp)));" |
| NL " failures += !(to_local(lintp));" |
| NL " failures += (*lintp != 2);" |
| NL |
| NL " intp = &pint;" |
| NL " private int * pintp = (private int *)intp;" |
| NL " failures += !(isFenceValid(get_fence(pintp)));" |
| NL " failures += !(to_private(pintp));" |
| NL " failures += (*pintp != 3);" |
| NL |
| NL " results[tid] = (failures == 0);" |
| NL "}" |
| NL |
| ); |
| |
| CBasicTest test(KERNEL_FUNCTIONS); |
| |
| return test.Execute(deviceID, context, queue, num_elements); |
| } |
| |
| int test_conditional_casting(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { |
| const std::string KERNEL_FUNCTION = common::CONFORMANCE_VERIFY_FENCE + |
| NL |
| NL "__global int gint = 1;" |
| NL |
| NL "__kernel void testKernel(__global uint *results) {" |
| NL " uint tid = get_global_id(0);" |
| NL |
| NL " int *ptr;" |
| NL " __local int lint;" |
| NL " lint = 2;" |
| NL |
| NL " if (tid % 2)" |
| NL " ptr = &gint;" |
| NL " else" |
| NL " ptr = &lint;" |
| NL |
| NL " barrier(CLK_GLOBAL_MEM_FENCE);" |
| NL |
| NL " if (tid % 2)" |
| NL " results[tid] = (isFenceValid(get_fence(ptr)) && to_global(ptr) && *ptr == 1);" |
| NL " else" |
| NL " results[tid] = (isFenceValid(get_fence(ptr)) && to_local(ptr) && *ptr == 2);" |
| NL "}" |
| NL; |
| |
| CBasicTest test(KERNEL_FUNCTION); |
| |
| return test.Execute(deviceID, context, queue, num_elements); |
| } |
| |
| int test_chain_casting(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { |
| const std::string KERNEL_FUNCTION = common::CONFORMANCE_VERIFY_FENCE + |
| NL |
| NL "__global int gint = 1;" |
| NL |
| NL "int f4(int val, int *ptr) { return (isFenceValid(get_fence(ptr)) && val == *ptr) ? 0 : 1; }" |
| NL "int f3(int val, int *ptr) { return f4(val, ptr); }" |
| NL "int f2(int *ptr, int val) { return f3(val, ptr); }" |
| NL "int f1(int *ptr, int val) { return f2(ptr, val); }" |
| NL |
| NL "__kernel void testKernel(__global uint *results) {" |
| NL " uint tid = get_global_id(0);" |
| NL |
| NL " int *ptr;" |
| NL " __local int lint;" |
| NL " lint = 2;" |
| NL " __private int pint = 3;" |
| NL |
| NL " int failures = 0;" |
| NL " failures += f1(&gint, gint);" |
| NL " failures += f1(&lint, lint);" |
| NL " failures += f1(&pint, pint);" |
| NL |
| NL " results[tid] = (failures == 0);" |
| NL "}" |
| NL; |
| CBasicTest test(KERNEL_FUNCTION); |
| |
| return test.Execute(deviceID, context, queue, num_elements); |
| } |
| |
| int test_ternary_operator_casting(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { |
| const std::string KERNEL_FUNCTION = common::CONFORMANCE_VERIFY_FENCE + |
| NL |
| NL "__global int gint = 1;" |
| NL |
| NL "__kernel void testKernel(__global uint *results) {" |
| NL " uint tid = get_global_id(0);" |
| NL |
| NL " int *ptr;" |
| NL " __local int lint;" |
| NL " lint = 2;" |
| NL |
| NL " ptr = (tid % 2) ? &gint : (int *)&lint; // assuming there is an implicit conversion from named address space to generic" |
| NL |
| NL " barrier(CLK_GLOBAL_MEM_FENCE);" |
| NL |
| NL " if (tid % 2)" |
| NL " results[tid] = (isFenceValid(get_fence(ptr)) && to_global(ptr) && *ptr == gint);" |
| NL " else" |
| NL " results[tid] = (isFenceValid(get_fence(ptr)) && to_local(ptr) && *ptr == lint);" |
| NL "}" |
| NL; |
| |
| CBasicTest test(KERNEL_FUNCTION); |
| |
| return test.Execute(deviceID, context, queue, num_elements); |
| } |
| |
| int test_language_struct(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { |
| std::vector<std::string> KERNEL_FUNCTIONS; |
| |
| // implicit private struct |
| KERNEL_FUNCTIONS.push_back(common::CONFORMANCE_VERIFY_FENCE + |
| NL |
| NL "__global int gint = 1;" |
| NL |
| NL "__kernel void testKernel(__global uint *results) {" |
| NL " uint tid = get_global_id(0);" |
| NL " int failures = 0;" |
| NL |
| NL " __local int lint;" |
| NL " lint = 2;" |
| NL " __private int pint = 3;" |
| NL |
| NL " struct {" |
| NL " __global int *gintp;" |
| NL " __local int *lintp;" |
| NL " __private int *pintp;" |
| NL " } structWithPointers;" |
| NL |
| NL " structWithPointers.gintp = &gint;" |
| NL " structWithPointers.lintp = &lint;" |
| NL " structWithPointers.pintp = &pint;" |
| NL |
| NL " failures += !(isFenceValid(get_fence(structWithPointers.gintp)));" |
| NL " failures += !(isFenceValid(get_fence(structWithPointers.lintp)));" |
| NL " failures += !(isFenceValid(get_fence(structWithPointers.pintp)));" |
| NL |
| NL " failures += !(to_global(structWithPointers.gintp));" |
| NL " failures += !(to_local(structWithPointers.lintp));" |
| NL " failures += !(to_private(structWithPointers.pintp));" |
| NL |
| NL " results[tid] = (failures == 0);" |
| NL "}" |
| NL |
| ); |
| |
| // explicit __private struct |
| KERNEL_FUNCTIONS.push_back(common::CONFORMANCE_VERIFY_FENCE + |
| NL |
| NL "__global int gint = 1;" |
| NL |
| NL "__kernel void testKernel(__global uint *results) {" |
| NL " uint tid = get_global_id(0);" |
| NL " int failures = 0;" |
| NL |
| NL " __local int lint;" |
| NL " lint = 2;" |
| NL " __private int pint = 3;" |
| NL |
| NL " typedef struct {" |
| NL " __global int * gintp;" |
| NL " __local int * lintp;" |
| NL " __private int * pintp;" |
| NL " } S;" |
| NL |
| NL " __private S structWithPointers;" |
| NL " structWithPointers.gintp = &gint;" |
| NL " structWithPointers.lintp = &lint;" |
| NL " structWithPointers.pintp = &pint;" |
| NL |
| NL " failures += !(isFenceValid(get_fence(structWithPointers.gintp)));" |
| NL " failures += !(isFenceValid(get_fence(structWithPointers.lintp)));" |
| NL " failures += !(isFenceValid(get_fence(structWithPointers.pintp)));" |
| NL |
| NL " failures += !(to_global(structWithPointers.gintp));" |
| NL " failures += !(to_local(structWithPointers.lintp));" |
| NL " failures += !(to_private(structWithPointers.pintp));" |
| NL |
| NL " results[tid] = (failures == 0);" |
| NL "}" |
| NL |
| ); |
| |
| KERNEL_FUNCTIONS.push_back(common::CONFORMANCE_VERIFY_FENCE + |
| NL |
| NL "__global int gint = 1;" |
| NL |
| NL "__kernel void testKernel(__global uint *results) {" |
| NL " uint tid = get_global_id(0);" |
| NL " int failures = 0;" |
| NL |
| NL " __local int lint;" |
| NL " lint = 2;" |
| NL " __private int pint = 3;" |
| NL |
| NL " typedef struct {" |
| NL " __global int * gintp;" |
| NL " __local int * lintp;" |
| NL " __private int * pintp;" |
| NL " } S;" |
| NL |
| NL " __local S structWithPointers;" |
| NL " structWithPointers.gintp = &gint;" |
| NL " structWithPointers.lintp = &lint;" |
| NL " structWithPointers.pintp = &pint;" |
| NL |
| NL " failures += !(isFenceValid(get_fence(structWithPointers.gintp)));" |
| NL " failures += !(isFenceValid(get_fence(structWithPointers.lintp)));" |
| NL " failures += !(isFenceValid(get_fence(structWithPointers.pintp)));" |
| NL |
| NL " failures += !(to_global(structWithPointers.gintp));" |
| NL " failures += !(to_local(structWithPointers.lintp));" |
| NL " failures += !(to_private(structWithPointers.pintp));" |
| NL |
| NL " results[tid] = (failures == 0);" |
| NL "}" |
| NL |
| ); |
| |
| KERNEL_FUNCTIONS.push_back(common::CONFORMANCE_VERIFY_FENCE + |
| NL |
| NL "typedef struct {" |
| NL " __global int *gintp;" |
| NL " __local int *lintp;" |
| NL " __private int *pintp;" |
| NL "} S;" |
| NL |
| NL "__global S structWithPointers;" |
| NL "__global int gint = 1;" |
| NL |
| NL "__kernel void testKernel(__global uint *results) {" |
| NL " uint tid = get_global_id(0);" |
| NL " int failures = 0;" |
| NL |
| NL " __local int lint;" |
| NL " lint = 2;" |
| NL " __private int pint = 3;" |
| NL |
| NL " structWithPointers.gintp = &gint;" |
| NL " structWithPointers.lintp = &lint;" |
| NL " structWithPointers.pintp = &pint;" |
| NL |
| NL " failures += !(isFenceValid(get_fence(structWithPointers.gintp)));" |
| NL " failures += !(isFenceValid(get_fence(structWithPointers.lintp)));" |
| NL " failures += !(isFenceValid(get_fence(structWithPointers.pintp)));" |
| NL |
| NL " failures += !(to_global(structWithPointers.gintp));" |
| NL " failures += !(to_local(structWithPointers.lintp));" |
| NL " failures += !(to_private(structWithPointers.pintp));" |
| NL |
| NL " results[tid] = (failures == 0);" |
| NL "}" |
| NL |
| ); |
| |
| CBasicTest test(KERNEL_FUNCTIONS); |
| |
| return test.Execute(deviceID, context, queue, num_elements); |
| } |
| |
| int test_language_union(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { |
| std::vector<std::string> KERNEL_FUNCTIONS; |
| |
| KERNEL_FUNCTIONS.push_back(common::CONFORMANCE_VERIFY_FENCE + |
| NL |
| NL "__global int g = 1;" |
| NL |
| NL "__kernel void testKernel(__global uint *results) {" |
| NL " uint tid = get_global_id(0);" |
| NL " int failures = 0;" |
| NL |
| NL " __local int l;" |
| NL " l = 2;" |
| NL " int p = 3;" |
| NL |
| NL " union {" |
| NL " __global int *gintp;" |
| NL " __local int *lintp;" |
| NL " __private int *pintp;" |
| NL " } u;" |
| NL |
| NL " u.gintp = &g;" |
| NL " failures += !(isFenceValid(get_fence(u.gintp)));" |
| NL " failures += !to_global(u.gintp);" |
| NL " failures += (*(u.gintp) != 1);" |
| NL |
| NL " u.lintp = &l;" |
| NL " failures += !(isFenceValid(get_fence(u.lintp)));" |
| NL " failures += !to_local(u.lintp);" |
| NL " failures += (*(u.lintp) != 2);" |
| NL |
| NL " u.pintp = &p;" |
| NL " failures += !(isFenceValid(get_fence(u.pintp)));" |
| NL " failures += !to_private(u.pintp);" |
| NL " failures += (*(u.pintp) != 3);" |
| NL |
| NL " results[tid] = (failures == 0);" |
| NL "}" |
| NL |
| ); |
| |
| KERNEL_FUNCTIONS.push_back(common::CONFORMANCE_VERIFY_FENCE + |
| NL |
| NL "__global int g = 1;" |
| NL |
| NL "__kernel void testKernel(__global uint *results) {" |
| NL " uint tid = get_global_id(0);" |
| NL " int failures = 0;" |
| NL |
| NL " __local int l;" |
| NL " l = 2;" |
| NL " int p = 3;" |
| NL |
| NL " typedef union {" |
| NL " __global int * gintp;" |
| NL " __local int * lintp;" |
| NL " __private int * pintp;" |
| NL " } U;" |
| NL |
| NL " __local U u;" |
| NL |
| NL " u.gintp = &g;" |
| NL " work_group_barrier(CLK_LOCAL_MEM_FENCE);" |
| NL " failures += !(isFenceValid(get_fence(u.gintp)));" |
| NL " failures += !to_global(u.gintp);" |
| NL " failures += (*(u.gintp) != 1);" |
| NL |
| NL " work_group_barrier(CLK_LOCAL_MEM_FENCE);" |
| NL " u.lintp = &l;" |
| NL " work_group_barrier(CLK_LOCAL_MEM_FENCE);" |
| NL " failures += !(isFenceValid(get_fence(u.lintp)));" |
| NL " failures += !to_local(u.lintp);" |
| NL " failures += (*(u.lintp) != 2);" |
| NL |
| NL " work_group_barrier(CLK_LOCAL_MEM_FENCE);" |
| NL " if(get_local_id(0) == 0) {" |
| NL " u.pintp = &p;" |
| NL " failures += !(isFenceValid(get_fence(u.pintp)));" |
| NL " failures += !to_private(u.pintp);" |
| NL " failures += (*(u.pintp) != 3);" |
| NL " }" |
| NL |
| NL " results[tid] = (failures == 0);" |
| NL "}" |
| NL |
| ); |
| |
| KERNEL_FUNCTIONS.push_back(common::CONFORMANCE_VERIFY_FENCE + |
| NL |
| NL "typedef union {" |
| NL " __global int * gintp;" |
| NL " __local int * lintp;" |
| NL " __private int * pintp;" |
| NL "} U;" |
| NL |
| NL "__global U u;" |
| NL "__global int g = 1;" |
| NL |
| NL "__kernel void testKernel(__global uint *results) {" |
| NL " uint tid = get_global_id(0);" |
| NL |
| NL " // for global unions only one thread should modify union's content" |
| NL " if (tid != 0) {" |
| NL " results[tid] = 1;" |
| NL " return;" |
| NL " }" |
| NL |
| NL " int failures = 0;" |
| NL |
| NL " __local int l;" |
| NL " l = 2;" |
| NL " int p = 3;" |
| NL |
| NL " u.gintp = &g;" |
| NL " failures += !(isFenceValid(get_fence(u.gintp)));" |
| NL " failures += !to_global(u.gintp);" |
| NL " failures += (*(u.gintp) != 1);" |
| NL |
| NL " u.lintp = &l;" |
| NL " failures += !(isFenceValid(get_fence(u.lintp)));" |
| NL " failures += !to_local(u.lintp);" |
| NL " failures += (*(u.lintp) != 2);" |
| NL |
| NL " u.pintp = &p;" |
| NL " failures += !(isFenceValid(get_fence(u.pintp)));" |
| NL " failures += !to_private(u.pintp);" |
| NL " failures += (*(u.pintp) != 3);" |
| NL |
| NL " results[tid] = (failures == 0);" |
| NL "}" |
| NL |
| ); |
| |
| CBasicTest test(KERNEL_FUNCTIONS); |
| |
| return test.Execute(deviceID, context, queue, num_elements); |
| } |
| |
| int test_multiple_calls_same_function(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { |
| const std::string KERNEL_FUNCTION = |
| NL |
| NL "int shift2(const int *ptr, int arg) {" |
| NL " return *ptr << arg;" |
| NL "}" |
| NL |
| NL "__kernel void testKernel(__global uint *results) {" |
| NL " uint tid = get_global_id(0);" |
| NL " int failures = 0;" |
| NL |
| NL " __local int val;" |
| NL " val = get_group_id(0);" |
| NL |
| NL " for (int i = 0; i < 5; i++) {" |
| NL " if (shift2(&val, i) != (val << i))" |
| NL " failures++;" |
| NL " }" |
| NL |
| NL " for (int i = 10; i > 5; i--) {" |
| NL " if (shift2(&val, i) != (val << i))" |
| NL " failures++;" |
| NL " }" |
| NL |
| NL " results[tid] = (failures == 0);" |
| NL "}" |
| NL; |
| |
| CBasicTest test(KERNEL_FUNCTION); |
| |
| return test.Execute(deviceID, context, queue, num_elements); |
| } |
| |
| int test_compare_pointers(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { |
| std::vector<std::string> KERNEL_FUNCTIONS; |
| |
| KERNEL_FUNCTIONS.push_back( |
| NL "__kernel void testKernel(__global uint *results) {" |
| NL " uint tid = get_global_id(0);" |
| NL |
| NL " int *ptr = NULL;" |
| NL |
| NL " results[tid] = (ptr == NULL);" |
| NL "}" |
| NL |
| ); |
| |
| KERNEL_FUNCTIONS.push_back( |
| NL "__kernel void testKernel(__global uint *results) {" |
| NL " uint tid = get_global_id(0);" |
| NL |
| NL " int *ptr = NULL;" |
| NL " __global int *gptr = NULL;" |
| NL |
| NL " results[tid] = (ptr == gptr);" |
| NL "}" |
| NL |
| ); |
| |
| KERNEL_FUNCTIONS.push_back( |
| NL "__kernel void testKernel(__global uint *results) {" |
| NL " uint tid = get_global_id(0);" |
| NL |
| NL " int *ptr = NULL;" |
| NL " __local int *lptr = NULL;" |
| NL |
| NL " results[tid] = (ptr == lptr);" |
| NL "}" |
| NL |
| ); |
| |
| KERNEL_FUNCTIONS.push_back( |
| NL "__kernel void testKernel(__global uint *results) {" |
| NL " uint tid = get_global_id(0);" |
| NL |
| NL " int *ptr = NULL;" |
| NL " __private int *pptr = NULL;" |
| NL |
| NL " results[tid] = (ptr == pptr);" |
| NL "}" |
| NL |
| ); |
| |
| KERNEL_FUNCTIONS.push_back( |
| NL "__kernel void testKernel(__global uint *results) {" |
| NL " uint tid = get_global_id(0);" |
| NL |
| NL " int *ptr = NULL;" |
| NL " __local int *lptr = NULL;" |
| NL " __global int *gptr = NULL;" |
| NL |
| NL " ptr = lptr;" |
| NL |
| NL " results[tid] = (gptr == ptr) && (lptr == ptr);" |
| NL "}" |
| NL |
| ); |
| |
| KERNEL_FUNCTIONS.push_back( |
| NL "__kernel void testKernel(__global uint *results) {" |
| NL " uint tid = get_global_id(0);" |
| NL |
| NL " int some_value = 7;" |
| NL " int *ptr = NULL;" |
| NL " __private int *pptr = &some_value;" |
| NL |
| NL " results[tid] = (ptr != pptr);" |
| NL "}" |
| NL |
| ); |
| |
| KERNEL_FUNCTIONS.push_back( |
| NL "__kernel void testKernel(__global uint *results) {" |
| NL " uint tid = get_global_id(0);" |
| NL |
| NL " __local int some_value;" |
| NL " some_value = 7;" |
| NL " int *ptr = NULL;" |
| NL " __local int *lptr = &some_value;" |
| NL |
| NL " results[tid] = (ptr != lptr);" |
| NL "}" |
| NL |
| ); |
| |
| KERNEL_FUNCTIONS.push_back( |
| NL "__global int some_value = 7;" |
| NL |
| NL "__kernel void testKernel(__global uint *results) {" |
| NL " uint tid = get_global_id(0);" |
| NL |
| NL " int *ptr = NULL;" |
| NL " __global int *gptr = &some_value;" |
| NL |
| NL " results[tid] = (ptr != gptr);" |
| NL "}" |
| NL |
| ); |
| |
| KERNEL_FUNCTIONS.push_back( |
| NL "__global int arr[5] = { 0, 1, 2, 3, 4 };" |
| NL |
| NL "__kernel void testKernel(__global uint *results) {" |
| NL " uint tid = get_global_id(0);" |
| NL |
| NL " int *ptr = &arr[1];" |
| NL " __global int *gptr = &arr[3];" |
| NL |
| NL " results[tid] = (gptr >= ptr);" |
| NL "}" |
| NL |
| ); |
| |
| CBasicTest test(KERNEL_FUNCTIONS); |
| |
| return test.Execute(deviceID, context, queue, num_elements); |
| } |
