From 8579e61401d66d9c1e9e2903cf4eaad1c8d5afdc Mon Sep 17 00:00:00 2001 From: James Price Date: Mon, 1 Jun 2020 09:31:50 -0400 Subject: [PATCH] Improve async build callback testing (#797) * Improve async build callback testing Check that the program build status can be queried from inside the program completion callback, and also add a test to ensure that the completion callback is called when the build fails. * Address review comments * Fix formatting --- .../compiler/test_async_build.cpp | 177 ++++++++++++------ 1 file changed, 122 insertions(+), 55 deletions(-) diff --git a/test_conformance/compiler/test_async_build.cpp b/test_conformance/compiler/test_async_build.cpp index 30024228..d1533624 100644 --- a/test_conformance/compiler/test_async_build.cpp +++ b/test_conformance/compiler/test_async_build.cpp @@ -1,6 +1,6 @@ // -// Copyright (c) 2017 The Khronos Group Inc. -// +// Copyright (c) 2017-2020 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 @@ -19,75 +19,142 @@ #include #endif +#include +#include + +namespace { const char *sample_async_kernel[] = { -"__kernel void sample_test(__global float *src, __global int *dst)\n" -"{\n" -" int tid = get_global_id(0);\n" -"\n" -" dst[tid] = (int)src[tid];\n" -"\n" -"}\n" }; + "__kernel void sample_test(__global float *src, __global int *dst)\n" + "{\n" + " size_t tid = get_global_id(0);\n" + "\n" + " dst[tid] = (int)src[tid];\n" + "\n" + "}\n" +}; -volatile int buildNotificationSent; +const char *sample_async_kernel_error[] = { + "__kernel void sample_test(__global float *src, __global int *dst)\n" + "{\n" + " size_t tid = get_global_id(0);\n" + "\n" + " dst[tid] = badcodehere;\n" + "\n" + "}\n" +}; -void CL_CALLBACK test_notify_build_complete( cl_program program, void *userData ) +// Data passed to a program completion callback +struct TestData { - if( userData == NULL || strcmp( (char *)userData, "userData" ) != 0 ) - { - log_error( "ERROR: User data passed in to build notify function was not correct!\n" ); - buildNotificationSent = -1; - } - else - buildNotificationSent = 1; - log_info( "\n <-- program successfully built\n" ); + cl_device_id device; + cl_build_status expectedStatus; +}; + +std::atomic callbackResult; + } -int test_async_build(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) +void CL_CALLBACK test_notify_build_complete(cl_program program, void *userData) { - int error; - cl_program program; + TestData *data = reinterpret_cast(userData); + + // Check user data is valid + if (data == nullptr) + { + log_error("ERROR: User data passed to callback was not valid!\n"); + callbackResult = -1; + return; + } + + // Get program build status cl_build_status status; - - - buildNotificationSent = 0; - - /* First, test by doing the slow method of the individual calls */ - error = create_single_kernel_helper_create_program(context, &program, 1, sample_async_kernel); - test_error(error, "Unable to create program from source"); - - /* Compile the program */ - error = clBuildProgram( program, 1, &deviceID, NULL, test_notify_build_complete, (void *)"userData" ); - test_error( error, "Unable to build program source" ); - - /* Wait for build to complete (just keep polling, since we're just a test */ - if( ( error = clGetProgramBuildInfo( program, deviceID, CL_PROGRAM_BUILD_STATUS, sizeof( status ), &status, NULL ) ) != CL_SUCCESS ) + cl_int err = + clGetProgramBuildInfo(program, data->device, CL_PROGRAM_BUILD_STATUS, + sizeof(cl_build_status), &status, NULL); + if (err != CL_SUCCESS) { - print_error( error, "Unable to get program build status" ); - return -1; - } - while( (int)status == CL_BUILD_IN_PROGRESS ) - { - log_info( "\n -- still waiting for build... (status is %d)", status ); - sleep( 1 ); - error = clGetProgramBuildInfo( program, deviceID, CL_PROGRAM_BUILD_STATUS, sizeof( status ), &status, NULL ); - test_error( error, "Unable to get program build status" ); + log_info("ERROR: failed to get build status from callback\n"); + callbackResult = -1; + return; } - if( status != CL_BUILD_SUCCESS ) - { - log_error( "ERROR: build failed! (status: %d)\n", (int)status ); - return -1; - } + log_info("Program completion callback received build status %d\n", status); - if( buildNotificationSent == 0 ) + // Check program build status matches expectation + if (status != data->expectedStatus) { - log_error( "ERROR: Async build completed, but build notification was not sent!\n" ); - return -1; + log_info("ERROR: build status %d != expected status %d\n", status, + data->expectedStatus); + callbackResult = -1; } + else + { + callbackResult = 1; + } +} - error = clReleaseProgram( program ); - test_error( error, "Unable to release program object" ); +int test_async_build(cl_device_id deviceID, cl_context context, + cl_command_queue queue, int num_elements) +{ + cl_int error; + + struct TestDef + { + const char **source; + cl_build_status expectedStatus; + }; + + TestDef testDefs[] = { { sample_async_kernel, CL_BUILD_SUCCESS }, + { sample_async_kernel_error, CL_BUILD_ERROR } }; + for (TestDef &testDef : testDefs) + { + log_info("\nTesting program that should produce status %d\n", + testDef.expectedStatus); + + // Create the program + clProgramWrapper program; + error = create_single_kernel_helper_create_program(context, &program, 1, + testDef.source); + test_error(error, "Unable to create program from source"); + + // Start an asynchronous build, registering the completion callback + TestData testData = { deviceID, testDef.expectedStatus }; + callbackResult = 0; + error = clBuildProgram(program, 1, &deviceID, NULL, + test_notify_build_complete, (void *)&testData); + // Allow implementations to return synchronous build failures. + // They still need to call the callback. + if (!(error == CL_BUILD_PROGRAM_FAILURE + && testDef.expectedStatus == CL_BUILD_ERROR)) + test_error(error, "Unable to start build"); + + // Wait for callback to fire + int timeout = 20; + while (callbackResult == 0) + { + if (timeout < 0) + { + log_error("Timeout while waiting for callback to fire.\n\n"); + return -1; + } + + log_info(" -- still waiting for callback...\n"); + sleep(1); + timeout--; + } + + // Check the callback result + if (callbackResult == 1) + { + log_error("Test passed.\n\n"); + } + else + { + log_error("Async build callback indicated test failure.\n\n"); + return -1; + } + } return 0; }