| // |
| // 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/os_helpers.h" |
| |
| const char *define_kernel_code[] = { |
| " #define VALUE\n" |
| "__kernel void define_test(__global int *src, __global int *dstA, __global int *dstB)\n" |
| "{\n" |
| " int tid = get_global_id(0);\n" |
| "#ifdef VALUE\n" |
| " dstA[tid] = src[tid] * 2;\n" |
| "#else\n" |
| " dstA[tid] = src[tid] * 4;\n" |
| "#endif\n" |
| "\n" |
| "#undef VALUE\n" |
| "#ifdef VALUE\n" |
| " dstB[tid] = src[tid] * 2;\n" |
| "#else\n" |
| " dstB[tid] = src[tid] * 4;\n" |
| "#endif\n" |
| "\n" |
| "}\n"}; |
| |
| |
| |
| |
| int test_preprocessor_define_udef(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { |
| |
| cl_int error; |
| clKernelWrapper kernel; |
| clProgramWrapper program; |
| clMemWrapper buffer[3]; |
| cl_int *srcData, *resultData; |
| int i; |
| MTdata d; |
| |
| error = create_single_kernel_helper(context, &program, &kernel, 1, define_kernel_code, "define_test"); |
| if (error) |
| return -1; |
| |
| buffer[0] = clCreateBuffer(context, CL_MEM_READ_WRITE, num_elements*sizeof(cl_int), NULL, &error); |
| test_error( error, "clCreateBuffer failed"); |
| buffer[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, num_elements*sizeof(cl_int), NULL, &error); |
| test_error( error, "clCreateBuffer failed"); |
| buffer[2] = clCreateBuffer(context, CL_MEM_READ_WRITE, num_elements*sizeof(cl_int), NULL, &error); |
| test_error( error, "clCreateBuffer failed"); |
| |
| srcData = (cl_int*)malloc(sizeof(cl_int)*num_elements); |
| if (srcData == NULL) { |
| log_error("Failed to allocate storage for source data (%d cl_ints).\n", num_elements); |
| return -1; |
| } |
| |
| d = init_genrand( gRandomSeed ); |
| for (i=0; i<num_elements; i++) |
| srcData[i] = (int)get_random_float(-1024, 1024,d); |
| free_mtdata(d); d = NULL; |
| |
| resultData = (cl_int*)malloc(sizeof(cl_int)*num_elements); |
| if (resultData == NULL) { |
| free(srcData); |
| log_error("Failed to allocate storage for result data (%d cl_ints).\n", num_elements); |
| return -1; |
| } |
| |
| error = clSetKernelArg(kernel, 0, sizeof(buffer[0]), &buffer[0]); |
| test_error(error, "clSetKernelArg failed"); |
| error = clSetKernelArg(kernel, 1, sizeof(buffer[1]), &buffer[1]); |
| test_error(error, "clSetKernelArg failed"); |
| error = clSetKernelArg(kernel, 2, sizeof(buffer[2]), &buffer[2]); |
| test_error(error, "clSetKernelArg failed"); |
| |
| |
| error = clEnqueueWriteBuffer(queue, buffer[0], CL_TRUE, 0, num_elements*sizeof(cl_int), srcData, 0, NULL, NULL); |
| test_error(error, "clEnqueueWriteBuffer failed"); |
| |
| size_t threads[3] = {num_elements, 0, 0}; |
| error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, threads, NULL, 0, NULL, NULL); |
| test_error(error, "clEnqueueNDRangeKernel failed"); |
| |
| error = clEnqueueReadBuffer(queue, buffer[1], CL_TRUE, 0, num_elements*sizeof(cl_int), resultData, 0, NULL, NULL); |
| test_error(error, "clEnqueueReadBuffer failed"); |
| |
| for (i=0; i<num_elements; i++) |
| if (resultData[i] != srcData[i]*2) { |
| free(srcData); |
| free(resultData); |
| return -1; |
| } |
| |
| error = clEnqueueReadBuffer(queue, buffer[2], CL_TRUE, 0, num_elements*sizeof(cl_int), resultData, 0, NULL, NULL); |
| test_error(error, "clEnqueueReadBuffer failed"); |
| |
| for (i=0; i<num_elements; i++) |
| if (resultData[i] != srcData[i]*4) { |
| free(srcData); |
| free(resultData); |
| return -1; |
| } |
| |
| free(srcData); |
| free(resultData); |
| return 0; |
| } |
| |
| |
| const char *include_kernel_code = |
| "#include \"%s\"\n" |
| "__kernel void include_test(__global int *src, __global int *dstA)\n" |
| "{\n" |
| " int tid = get_global_id(0);\n" |
| "#ifdef HEADER_FOUND\n" |
| " dstA[tid] = HEADER_FOUND;\n" |
| "#else\n" |
| " dstA[tid] = 0;\n" |
| "#endif\n" |
| "\n" |
| "}\n"; |
| |
| |
| int test_preprocessor_include(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { |
| |
| cl_int error; |
| clKernelWrapper kernel; |
| clProgramWrapper program; |
| clMemWrapper buffer[2]; |
| cl_int *resultData; |
| int i; |
| |
| char include_dir[4096] = {0}; |
| char include_kernel[4096] = {0}; |
| |
| char const * sep = get_dir_sep(); |
| char const * path = get_exe_dir(); |
| |
| /* Build with the include directory defined */ |
| sprintf(include_dir,"%s%sincludeTestDirectory%stestIncludeFile.h", path, sep, sep); |
| sprintf(include_kernel, include_kernel_code, include_dir); |
| free( (void *) sep ); |
| free( (void *) path ); |
| |
| const char* test_kernel[] = { include_kernel, 0 }; |
| error = create_single_kernel_helper(context, &program, &kernel, 1, test_kernel, "include_test"); |
| if (error) |
| return -1; |
| |
| buffer[0] = clCreateBuffer(context, CL_MEM_READ_WRITE, num_elements*sizeof(cl_int), NULL, &error); |
| test_error( error, "clCreateBuffer failed"); |
| buffer[1] = clCreateBuffer(context, CL_MEM_READ_WRITE, num_elements*sizeof(cl_int), NULL, &error); |
| test_error( error, "clCreateBuffer failed"); |
| |
| resultData = (cl_int*)malloc(sizeof(cl_int)*num_elements); |
| if (resultData == NULL) { |
| log_error("Failed to allocate storage for result data (%d cl_ints).\n", num_elements); |
| return -1; |
| } |
| |
| error = clSetKernelArg(kernel, 0, sizeof(buffer[0]), &buffer[0]); |
| test_error(error, "clSetKernelArg failed"); |
| error = clSetKernelArg(kernel, 1, sizeof(buffer[1]), &buffer[1]); |
| test_error(error, "clSetKernelArg failed"); |
| |
| size_t threads[3] = {num_elements, 0, 0}; |
| error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, threads, NULL, 0, NULL, NULL); |
| test_error(error, "clEnqueueNDRangeKernel failed"); |
| |
| error = clEnqueueReadBuffer(queue, buffer[1], CL_TRUE, 0, num_elements*sizeof(cl_int), resultData, 0, NULL, NULL); |
| test_error(error, "clEnqueueReadBuffer failed"); |
| |
| for (i=0; i<num_elements; i++) |
| if (resultData[i] != 12) { |
| free(resultData); |
| return -1; |
| } |
| |
| free(resultData); |
| return 0; |
| } |
| |
| |
| |
| |
| const char *line_error_kernel_code[] = { |
| "__kernel void line_error_test(__global int *dstA)\n" |
| "{\n" |
| " int tid = get_global_id(0);\n" |
| "#line 124 \"fictitious/file/name.c\" \n" |
| "#error some error\n" |
| " dstA[tid] = tid;\n" |
| "\n" |
| "}\n"}; |
| |
| |
| int test_preprocessor_line_error(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { |
| |
| cl_int error, error2; |
| clKernelWrapper kernel; |
| clProgramWrapper program; |
| clMemWrapper buffer[2]; |
| |
| char buildLog[ 1024 * 128 ]; |
| |
| log_info("test_preprocessor_line_error may report spurious ERRORS in the conformance log.\n"); |
| |
| /* Create the program object from source */ |
| program = clCreateProgramWithSource( context, 1, line_error_kernel_code, NULL, &error ); |
| test_error(error, "clCreateProgramWithSource failed"); |
| |
| /* Compile the program */ |
| error2 = clBuildProgram( program, 0, NULL, NULL, NULL, NULL ); |
| if (error2) { |
| log_info("Build error detected at clBuildProgram."); |
| } else { |
| log_info("Error not reported by clBuildProgram.\n"); |
| } |
| |
| cl_build_status status; |
| error = clGetProgramBuildInfo(program, deviceID, CL_PROGRAM_BUILD_STATUS, sizeof(status), &status, NULL); |
| test_error(error, "clGetProgramBuildInfo failed for CL_PROGRAM_BUILD_STATUS"); |
| if (status != CL_BUILD_ERROR) { |
| log_error("Build status did not return CL_BUILD_ERROR for a program with #error defined.\n"); |
| return -1; |
| } else if (status == CL_BUILD_ERROR || error2) { |
| error2 = clGetProgramBuildInfo( program, deviceID, CL_PROGRAM_BUILD_LOG, sizeof( buildLog ), buildLog, NULL ); |
| test_error( error2, "Unable to get program build log" ); |
| |
| log_info("Build failed as expected with #error in source:\n"); |
| log_info( "Build log is: ------------\n" ); |
| log_info( "%s\n", buildLog ); |
| log_info( "Original source is: ------------\n" ); |
| log_info( "%s", line_error_kernel_code[0] ); |
| log_info( "\n----------\n" ); |
| |
| if (strstr(buildLog, "fictitious/file/name.c")) { |
| log_info("Found file name from #line param in log output.\n"); |
| } else { |
| log_info("WARNING: Did not find file name from #line param in log output.\n"); |
| } |
| |
| if (strstr(buildLog, "124")) { |
| log_info("Found line number from #line param in log output.\n"); |
| } else { |
| log_info("WARNING: Did not find line number from #line param in log output.\n"); |
| } |
| |
| log_info("test_preprocessor_line_error PASSED.\n"); |
| return 0; |
| } |
| |
| /* And create a kernel from it */ |
| kernel = clCreateKernel( program, "line_error_test", &error ); |
| test_error(error, "clCreateKernel failed"); |
| |
| buffer[0] = clCreateBuffer(context, CL_MEM_READ_WRITE, num_elements*sizeof(cl_int), NULL, &error); |
| test_error( error, "clCreateBuffer failed"); |
| |
| error = clSetKernelArg(kernel, 0, sizeof(buffer[0]), &buffer[0]); |
| test_error(error, "clSetKernelArg failed"); |
| |
| size_t threads[3] = {num_elements, 0, 0}; |
| error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, threads, NULL, 0, NULL, NULL); |
| test_error(error, "clEnqueueNDRangeKernel failed"); |
| |
| log_error("Program built and ran with #error defined."); |
| return -1; |
| } |
| |
| |
| |
| const char *pragma_kernel_code[] = { |
| "__kernel void pragma_test(__global int *dstA)\n" |
| "{\n" |
| "#pragma A fool thinks himself to be wise, but a wise man knows himself to be a fool.\n" |
| " int tid = get_global_id(0);\n" |
| "#pragma\n" |
| " dstA[tid] = tid;\n" |
| "#pragma mark Though I am not naturally honest, I am so sometimes by chance.\n" |
| "\n" |
| "}\n"}; |
| |
| |
| int test_preprocessor_pragma(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) { |
| |
| cl_int error; |
| clKernelWrapper kernel; |
| clProgramWrapper program; |
| clMemWrapper buffer[2]; |
| cl_int *resultData; |
| int i; |
| |
| error = create_single_kernel_helper(context, &program, &kernel, 1, pragma_kernel_code, "pragma_test"); |
| if (error) |
| return -1; |
| |
| buffer[0] = clCreateBuffer(context, CL_MEM_READ_WRITE, num_elements*sizeof(cl_int), NULL, &error); |
| test_error( error, "clCreateBuffer failed"); |
| |
| error = clSetKernelArg(kernel, 0, sizeof(buffer[0]), &buffer[0]); |
| test_error(error, "clSetKernelArg failed"); |
| |
| size_t threads[3] = {num_elements, 0, 0}; |
| error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, threads, NULL, 0, NULL, NULL); |
| test_error(error, "clEnqueueNDRangeKernel failed"); |
| |
| resultData = (cl_int*)malloc(sizeof(cl_int)*num_elements); |
| if (resultData == NULL) { |
| log_error("Failed to allocate storage for result data (%d cl_ints).\n", num_elements); |
| return -1; |
| } |
| |
| error = clEnqueueReadBuffer(queue, buffer[0], CL_TRUE, 0, num_elements*sizeof(cl_int), resultData, 0, NULL, NULL); |
| test_error(error, "clEnqueueReadBuffer failed"); |
| |
| for (i=0; i<num_elements; i++) |
| if (resultData[i] != i) { |
| free(resultData); |
| return -1; |
| } |
| |
| free(resultData); |
| return 0; |
| } |
| |
| |
| |