// // 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. // #ifndef _WIN32 #include #endif #include "procs.h" #include "harness/conversions.h" #include "harness/typeWrappers.h" #include "harness/errorHelpers.h" // For global, local, and constant const char *parameter_kernel_long = "%s\n" // optional pragma "kernel void test(global ulong *results, %s %s *mem0, %s %s2 *mem2, %s %s3 *mem3, %s %s4 *mem4, %s %s8 *mem8, %s %s16 *mem16)\n" "{\n" " results[0] = (ulong)&mem0[0];\n" " results[1] = (ulong)&mem2[0];\n" " results[2] = (ulong)&mem3[0];\n" " results[3] = (ulong)&mem4[0];\n" " results[4] = (ulong)&mem8[0];\n" " results[5] = (ulong)&mem16[0];\n" "}\n"; // For private and local const char *local_kernel_long = "%s\n" // optional pragma "kernel void test(global ulong *results)\n" "{\n" " %s %s mem0[3];\n" " %s %s2 mem2[3];\n" " %s %s3 mem3[3];\n" " %s %s4 mem4[3];\n" " %s %s8 mem8[3];\n" " %s %s16 mem16[3];\n" " results[0] = (ulong)&mem0[0];\n" " results[1] = (ulong)&mem2[0];\n" " results[2] = (ulong)&mem3[0];\n" " results[3] = (ulong)&mem4[0];\n" " results[4] = (ulong)&mem8[0];\n" " results[5] = (ulong)&mem16[0];\n" "}\n"; // For constant const char *constant_kernel_long = "%s\n" // optional pragma " constant %s mem0[3] = {0};\n" " constant %s2 mem2[3] = {(%s2)(0)};\n" " constant %s3 mem3[3] = {(%s3)(0)};\n" " constant %s4 mem4[3] = {(%s4)(0)};\n" " constant %s8 mem8[3] = {(%s8)(0)};\n" " constant %s16 mem16[3] = {(%s16)(0)};\n" "\n" "kernel void test(global ulong *results)\n" "{\n" " results[0] = (ulong)&mem0;\n" " results[1] = (ulong)&mem2;\n" " results[2] = (ulong)&mem3;\n" " results[3] = (ulong)&mem4;\n" " results[4] = (ulong)&mem8;\n" " results[5] = (ulong)&mem16;\n" "}\n"; // For global, local, and constant const char *parameter_kernel_no_long = "%s\n" // optional pragma "kernel void test(global uint *results, %s %s *mem0, %s %s2 *mem2, %s %s3 *mem3, %s %s4 *mem4, %s %s8 *mem8, %s %s16 *mem16)\n" "{\n" " results[0] = (uint)&mem0[0];\n" " results[1] = (uint)&mem2[0];\n" " results[2] = (uint)&mem3[0];\n" " results[3] = (uint)&mem4[0];\n" " results[4] = (uint)&mem8[0];\n" " results[5] = (uint)&mem16[0];\n" "}\n"; // For private and local const char *local_kernel_no_long = "%s\n" // optional pragma "kernel void test(global uint *results)\n" "{\n" " %s %s mem0[3];\n" " %s %s2 mem2[3];\n" " %s %s3 mem3[3];\n" " %s %s4 mem4[3];\n" " %s %s8 mem8[3];\n" " %s %s16 mem16[3];\n" " results[0] = (uint)&mem0[0];\n" " results[1] = (uint)&mem2[0];\n" " results[2] = (uint)&mem3[0];\n" " results[3] = (uint)&mem4[0];\n" " results[4] = (uint)&mem8[0];\n" " results[5] = (uint)&mem16[0];\n" "}\n"; // For constant const char *constant_kernel_no_long = "%s\n" // optional pragma " constant %s mem0[3] = {0};\n" " constant %s2 mem2[3] = {(%s2)(0)};\n" " constant %s3 mem3[3] = {(%s3)(0)};\n" " constant %s4 mem4[3] = {(%s4)(0)};\n" " constant %s8 mem8[3] = {(%s8)(0)};\n" " constant %s16 mem16[3] = {(%s16)(0)};\n" "\n" "kernel void test(global uint *results)\n" "{\n" " results[0] = (uint)&mem0;\n" " results[1] = (uint)&mem2;\n" " results[2] = (uint)&mem3;\n" " results[3] = (uint)&mem4;\n" " results[4] = (uint)&mem8;\n" " results[5] = (uint)&mem16;\n" "}\n"; enum AddressSpaces { kGlobal = 0, kLocal, kConstant, kPrivate }; typedef enum AddressSpaces AddressSpaces; #define DEBUG 0 const char * get_explicit_address_name( AddressSpaces address ) { /* Quick method to avoid branching: make sure the following array matches the Enum order */ static const char *sExplicitAddressNames[] = { "global", "local", "constant", "private"}; return sExplicitAddressNames[ address ]; } int test_kernel_memory_alignment(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems, AddressSpaces address ) { const char *constant_kernel; const char *parameter_kernel; const char *local_kernel; if ( gHasLong ) { constant_kernel = constant_kernel_long; parameter_kernel = parameter_kernel_long; local_kernel = local_kernel_long; } else { constant_kernel = constant_kernel_no_long; parameter_kernel = parameter_kernel_no_long; local_kernel = local_kernel_no_long; } ExplicitType vecType[] = { kChar, kUChar, kShort, kUShort, kInt, kUInt, kLong, kULong, kFloat, kDouble }; char *kernel_code = (char*)malloc(4096); cl_kernel kernel; cl_program program; int error; int total_errors = 0; cl_mem results; cl_ulong *results_data; cl_mem mem0, mem2, mem3, mem4, mem8, mem16; results_data = (cl_ulong*)malloc(sizeof(cl_ulong)*6); results = clCreateBuffer(context, 0, sizeof(cl_ulong)*6, NULL, &error); test_error(error, "clCreateBuffer failed"); mem0 = clCreateBuffer(context, 0, sizeof(cl_long), NULL, &error); test_error(error, "clCreateBuffer failed"); mem2 = clCreateBuffer(context, 0, sizeof(cl_long)*2, NULL, &error); test_error(error, "clCreateBuffer failed"); mem3 = clCreateBuffer(context, 0, sizeof(cl_long)*4, NULL, &error); test_error(error, "clCreateBuffer failed"); mem4 = clCreateBuffer(context, 0, sizeof(cl_long)*4, NULL, &error); test_error(error, "clCreateBuffer failed"); mem8 = clCreateBuffer(context, 0, sizeof(cl_long)*8, NULL, &error); test_error(error, "clCreateBuffer failed"); mem16 = clCreateBuffer(context, 0, sizeof(cl_long)*16, NULL, &error); test_error(error, "clCreateBuffer failed"); // For each type // Calculate alignment mask for each size // For global, local, constant, private // If global, local or constant -- do parameter_kernel // If private or local -- do local_kernel // If constant -- do constant kernel int numConstantArgs; clGetDeviceInfo(device, CL_DEVICE_MAX_CONSTANT_ARGS, sizeof(numConstantArgs), &numConstantArgs, NULL); int typeIndex; for (typeIndex = 0; typeIndex < 10; typeIndex++) { // Skip double tests if we don't support doubles if (vecType[typeIndex] == kDouble && !is_extension_available(device, "cl_khr_fp64")) { log_info("Extension cl_khr_fp64 not supported; skipping double tests.\n"); continue; } if (( vecType[ typeIndex ] == kLong || vecType[ typeIndex ] == kULong ) && !gHasLong ) continue; log_info("Testing %s...\n", get_explicit_type_name(vecType[typeIndex])); // Determine the expected alignment masks. // E.g., if it is supposed to be 4 byte aligned, we should get 4-1=3 = ... 000011 // We can then and the returned address with that and we should have 0. cl_ulong alignments[6]; alignments[0] = get_explicit_type_size(vecType[typeIndex])-1; alignments[1] = (get_explicit_type_size(vecType[typeIndex])<<1)-1; alignments[2] = (get_explicit_type_size(vecType[typeIndex])<<2)-1; alignments[3] = (get_explicit_type_size(vecType[typeIndex])<<2)-1; alignments[4] = (get_explicit_type_size(vecType[typeIndex])<<3)-1; alignments[5] = (get_explicit_type_size(vecType[typeIndex])<<4)-1; // Parameter kernel if (address == kGlobal || address == kLocal || address == kConstant) { log_info("\tTesting parameter kernel...\n"); if ( (gIsEmbedded) && (address == kConstant) && (numConstantArgs < 6)) { sprintf(kernel_code, parameter_kernel, vecType[typeIndex] == kDouble ? "#pragma OPENCL EXTENSION cl_khr_fp64 : enable" : "", get_explicit_address_name(address), get_explicit_type_name(vecType[typeIndex]), get_explicit_address_name(address), get_explicit_type_name(vecType[typeIndex]), get_explicit_address_name(address), get_explicit_type_name(vecType[typeIndex]), get_explicit_address_name(address), get_explicit_type_name(vecType[typeIndex]), get_explicit_address_name(address), get_explicit_type_name(vecType[typeIndex]), get_explicit_address_name(address), get_explicit_type_name(vecType[typeIndex]) ); } else { sprintf(kernel_code, parameter_kernel, vecType[typeIndex] == kDouble ? "#pragma OPENCL EXTENSION cl_khr_fp64 : enable" : "", get_explicit_address_name(address), get_explicit_type_name(vecType[typeIndex]), get_explicit_address_name(address), get_explicit_type_name(vecType[typeIndex]), get_explicit_address_name(address), get_explicit_type_name(vecType[typeIndex]), get_explicit_address_name(address), get_explicit_type_name(vecType[typeIndex]), get_explicit_address_name(address), get_explicit_type_name(vecType[typeIndex]), get_explicit_address_name(address), get_explicit_type_name(vecType[typeIndex]) ); } //printf("Kernel is: \n%s\n", kernel_code); // Create the kernel error = create_single_kernel_helper(context, &program, &kernel, 1, (const char **)&kernel_code, "test"); test_error(error, "create_single_kernel_helper failed"); // Initialize the results memset(results_data, 0, sizeof(cl_long)*5); error = clEnqueueWriteBuffer(queue, results, CL_TRUE, 0, sizeof(cl_long)*6, results_data, 0, NULL, NULL); test_error(error, "clEnqueueWriteBuffer failed"); // Set the arguments error = clSetKernelArg(kernel, 0, sizeof(results), &results); test_error(error, "clSetKernelArg failed"); if (address != kLocal) { error = clSetKernelArg(kernel, 1, sizeof(mem0), &mem0); test_error(error, "clSetKernelArg failed"); error = clSetKernelArg(kernel, 2, sizeof(mem2), &mem2); test_error(error, "clSetKernelArg failed"); error = clSetKernelArg(kernel, 3, sizeof(mem3), &mem3); test_error(error, "clSetKernelArg failed"); error = clSetKernelArg(kernel, 4, sizeof(mem4), &mem4); test_error(error, "clSetKernelArg failed"); error = clSetKernelArg(kernel, 5, sizeof(mem8), &mem8); test_error(error, "clSetKernelArg failed"); error = clSetKernelArg(kernel, 6, sizeof(mem16), &mem16); test_error(error, "clSetKernelArg failed"); } else { error = clSetKernelArg(kernel, 1, get_explicit_type_size(vecType[typeIndex]), NULL); test_error(error, "clSetKernelArg failed"); error = clSetKernelArg(kernel, 2, get_explicit_type_size(vecType[typeIndex])*2, NULL); test_error(error, "clSetKernelArg failed"); error = clSetKernelArg(kernel, 3, get_explicit_type_size(vecType[typeIndex])*4, NULL); test_error(error, "clSetKernelArg failed"); error = clSetKernelArg(kernel, 4, get_explicit_type_size(vecType[typeIndex])*4, NULL); test_error(error, "clSetKernelArg failed"); error = clSetKernelArg(kernel, 5, get_explicit_type_size(vecType[typeIndex])*8, NULL); test_error(error, "clSetKernelArg failed"); error = clSetKernelArg(kernel, 6, get_explicit_type_size(vecType[typeIndex])*16, NULL); test_error(error, "clSetKernelArg failed"); } // Enqueue the kernel size_t global_size = 1; error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global_size, NULL, 0, NULL, NULL); test_error(error, "clEnqueueNDRangeKernel failed"); // Read back the results error = clEnqueueReadBuffer(queue, results, CL_TRUE, 0, sizeof(cl_ulong)*6, results_data, 0, NULL, NULL); test_error(error, "clEnqueueReadBuffer failed"); // Verify the results if (gHasLong) { for (int i = 0; i < 6; i++) { if ((results_data[i] & alignments[i]) != 0) { total_errors++; log_error("\tVector size %d failed: 0x%llx is not properly aligned.\n", 1 << i, results_data[i]); } else { if (DEBUG) log_info("\tVector size %d passed: 0x%llx is properly aligned.\n", 1 << i, results_data[i]); } } } // Verify the results on devices that do not support longs else { cl_uint *results_data_no_long = (cl_uint *)results_data; for (int i = 0; i < 6; i++) { if ((results_data_no_long[i] & alignments[i]) != 0) { total_errors++; log_error("\tVector size %d failed: 0x%llx is not properly aligned.\n", 1 << i, results_data_no_long[i]); } else { if (DEBUG) log_info("\tVector size %d passed: 0x%llx is properly aligned.\n", 1 << i, results_data_no_long[i]); } } } clReleaseKernel(kernel); clReleaseProgram(program); } // Local kernel if (address == kLocal || address == kPrivate) { log_info("\tTesting local kernel...\n"); sprintf(kernel_code, local_kernel, vecType[typeIndex] == kDouble ? "#pragma OPENCL EXTENSION cl_khr_fp64 : enable" : "", get_explicit_address_name(address), get_explicit_type_name(vecType[typeIndex]), get_explicit_address_name(address), get_explicit_type_name(vecType[typeIndex]), get_explicit_address_name(address), get_explicit_type_name(vecType[typeIndex]), get_explicit_address_name(address), get_explicit_type_name(vecType[typeIndex]), get_explicit_address_name(address), get_explicit_type_name(vecType[typeIndex]), get_explicit_address_name(address), get_explicit_type_name(vecType[typeIndex]) ); //printf("Kernel is: \n%s\n", kernel_code); // Create the kernel error = create_single_kernel_helper(context, &program, &kernel, 1, (const char **)&kernel_code, "test"); test_error(error, "create_single_kernel_helper failed"); // Initialize the results memset(results_data, 0, sizeof(cl_long)*5); error = clEnqueueWriteBuffer(queue, results, CL_TRUE, 0, sizeof(cl_long)*5, results_data, 0, NULL, NULL); test_error(error, "clEnqueueWriteBuffer failed"); // Set the arguments error = clSetKernelArg(kernel, 0, sizeof(results), &results); test_error(error, "clSetKernelArg failed"); // Enqueue the kernel size_t global_size = 1; error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global_size, NULL, 0, NULL, NULL); test_error(error, "clEnqueueNDRangeKernel failed"); // Read back the results error = clEnqueueReadBuffer(queue, results, CL_TRUE, 0, sizeof(cl_ulong)*5, results_data, 0, NULL, NULL); test_error(error, "clEnqueueReadBuffer failed"); // Verify the results if (gHasLong) { for (int i = 0; i < 5; i++) { if ((results_data[i] & alignments[i]) != 0) { total_errors++; log_error("\tVector size %d failed: 0x%llx is not properly aligned.\n", 1 << i, results_data[i]); } else { if (DEBUG) log_info("\tVector size %d passed: 0x%llx is properly aligned.\n", 1 << i, results_data[i]); } } } // Verify the results on devices that do not support longs else { cl_uint *results_data_no_long = (cl_uint *)results_data; for (int i = 0; i < 5; i++) { if ((results_data_no_long[i] & alignments[i]) != 0) { total_errors++; log_error("\tVector size %d failed: 0x%llx is not properly aligned.\n", 1 << i, results_data_no_long[i]); } else { if (DEBUG) log_info("\tVector size %d passed: 0x%llx is properly aligned.\n", 1 << i, results_data_no_long[i]); } } } clReleaseKernel(kernel); clReleaseProgram(program); } // Constant kernel if (address == kConstant) { log_info("\tTesting constant kernel...\n"); sprintf(kernel_code, constant_kernel, vecType[typeIndex] == kDouble ? "#pragma OPENCL EXTENSION cl_khr_fp64 : enable" : "", get_explicit_type_name(vecType[typeIndex]), get_explicit_type_name(vecType[typeIndex]), get_explicit_type_name(vecType[typeIndex]), get_explicit_type_name(vecType[typeIndex]), get_explicit_type_name(vecType[typeIndex]), get_explicit_type_name(vecType[typeIndex]), get_explicit_type_name(vecType[typeIndex]), get_explicit_type_name(vecType[typeIndex]), get_explicit_type_name(vecType[typeIndex]), get_explicit_type_name(vecType[typeIndex]), get_explicit_type_name(vecType[typeIndex]), get_explicit_type_name(vecType[typeIndex]) ); //printf("Kernel is: \n%s\n", kernel_code); // Create the kernel error = create_single_kernel_helper(context, &program, &kernel, 1, (const char **)&kernel_code, "test"); test_error(error, "create_single_kernel_helper failed"); // Initialize the results memset(results_data, 0, sizeof(cl_long)*5); error = clEnqueueWriteBuffer(queue, results, CL_TRUE, 0, sizeof(cl_long)*5, results_data, 0, NULL, NULL); test_error(error, "clEnqueueWriteBuffer failed"); // Set the arguments error = clSetKernelArg(kernel, 0, sizeof(results), &results); test_error(error, "clSetKernelArg failed"); // Enqueue the kernel size_t global_size = 1; error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global_size, NULL, 0, NULL, NULL); test_error(error, "clEnqueueNDRangeKernel failed"); // Read back the results error = clEnqueueReadBuffer(queue, results, CL_TRUE, 0, sizeof(cl_ulong)*5, results_data, 0, NULL, NULL); test_error(error, "clEnqueueReadBuffer failed"); // Verify the results if (gHasLong) { for (int i = 0; i < 5; i++) { if ((results_data[i] & alignments[i]) != 0) { total_errors++; log_error("\tVector size %d failed: 0x%llx is not properly aligned.\n", 1 << i, results_data[i]); } else { if (DEBUG) log_info("\tVector size %d passed: 0x%llx is properly aligned.\n", 1 << i, results_data[i]); } } } // Verify the results on devices that do not support longs else { cl_uint *results_data_no_long = (cl_uint *)results_data; for (int i = 0; i < 5; i++) { if ((results_data_no_long[i] & alignments[i]) != 0) { total_errors++; log_error("\tVector size %d failed: 0x%llx is not properly aligned.\n", 1 << i, results_data_no_long[i]); } else { if (DEBUG) log_info("\tVector size %d passed: 0x%llx is properly aligned.\n", 1 << i, results_data_no_long[i]); } } } clReleaseKernel(kernel); clReleaseProgram(program); } } clReleaseMemObject(results); clReleaseMemObject(mem0); clReleaseMemObject(mem2); clReleaseMemObject(mem3); clReleaseMemObject(mem4); clReleaseMemObject(mem8); clReleaseMemObject(mem16); free( kernel_code ); free( results_data ); if (total_errors != 0) return -1; return 0; } int test_kernel_memory_alignment_local(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems ) { return test_kernel_memory_alignment( device, context, queue, n_elems, kLocal ); } int test_kernel_memory_alignment_global(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems ) { return test_kernel_memory_alignment( device, context, queue, n_elems, kGlobal ); } int test_kernel_memory_alignment_constant(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems ) { // There is a class of approved OpenCL 1.0 conformant devices out there that in some circumstances // are unable to meaningfully take (or more precisely use) the address of constant data by virtue // of limitations in their ISA design. This feature was not tested in 1.0, so they were declared // conformant by Khronos. The failure is however caught here. // // Unfortunately, determining whether or not these devices are 1.0 conformant is not the jurisdiction // of the 1.1 tests -- We can't fail them from 1.1 conformance here because they are not 1.1 // devices. They are merely 1.0 conformant devices that interop with 1.1 devices in a 1.1 platform. // To add new binding tests now to conformant 1.0 devices would violate the workingroup requirement // of no new tests for 1.0 devices. So certain allowances have to be made in intractable cases // such as this one. // // There is some precedent. Similar allowances are made for other 1.0 hardware features such as // local memory size. The minimum required local memory size grew from 16 kB to 32 kB in OpenCL 1.1. // Detect 1.0 devices // Get CL_DEVICE_VERSION size size_t string_size = 0; int err; if( (err = clGetDeviceInfo( device, CL_DEVICE_VERSION, 0, NULL, &string_size ) ) ) { log_error( "FAILURE: Unable to get size of CL_DEVICE_VERSION string!" ); return -1; } //Allocate storage to hold the version string char *version_string = (char*) malloc(string_size); if( NULL == version_string ) { log_error( "FAILURE: Unable to allocate memory to hold CL_DEVICE_VERSION string!" ); return -1; } // Get CL_DEVICE_VERSION string if( (err = clGetDeviceInfo( device, CL_DEVICE_VERSION, string_size, version_string, NULL ) ) ) { log_error( "FAILURE: Unable to read CL_DEVICE_VERSION string!" ); return -1; } // easy out for 1.0 devices const char *string_1_0 = "OpenCL 1.0 "; if( 0 == strncmp( version_string, string_1_0, strlen(string_1_0)) ) { log_info( "WARNING: Allowing device to escape testing of difficult constant memory alignment case.\n\tDevice is not a OpenCL 1.1 device. CL_DEVICE_VERSION: \"%s\"\n", version_string ); free(version_string); return 0; } log_info( "Device version string: \"%s\"\n", version_string ); free(version_string); // Everyone else is to be ground mercilessly under the wheels of progress return test_kernel_memory_alignment( device, context, queue, n_elems, kConstant ); } int test_kernel_memory_alignment_private(cl_device_id device, cl_context context, cl_command_queue queue, int n_elems ) { return test_kernel_memory_alignment( device, context, queue, n_elems, kPrivate ); }