From b5a6f3fda8bcd7701c6dd75cc151ad56324d933e Mon Sep 17 00:00:00 2001 From: Ross Brunton Date: Mon, 18 Mar 2024 18:02:35 +0000 Subject: [PATCH] [UR] Improve handling of error cases in urProgramLink Note that this change includes a specification change: urProgramLink now requires the output parameter to contain either nullptr or some unspecified binary on failure. As well as this change, a number of bugs have been fixed: * The Level Zero adapter now correctly returns `UR_RESULT_ERROR_PROGRAM_LINK_FAILURE` when linking fails, rather than `UR_RESULT_ERROR_UNKNOWN`. * A workaround has been added for some OpenCL devices that return `CL_INVALID_BINARY` rather than `CL_LINK_PROGRAM_FAILURE` on linker failure. * The `phProgram` handle is wrapped in a loader handle by the loader even if an error would be returned. This is required by Level Zero, which outputs a "dummy" program to store the linker log. Conformance tests have also been added. --- include/ur_api.h | 10 +++ scripts/core/exp-multi-device-compile.yml | 1 + scripts/core/program.yml | 1 + scripts/templates/helper.py | 24 +++++++ scripts/templates/ldrddi.cpp.mako | 6 +- scripts/templates/libapi.cpp.mako | 2 +- scripts/templates/nullddi.cpp.mako | 1 + scripts/templates/trcddi.cpp.mako | 2 +- scripts/templates/valddi.cpp.mako | 2 +- source/adapters/cuda/program.cpp | 9 ++- source/adapters/hip/program.cpp | 15 +++-- source/adapters/level_zero/common.cpp | 2 + source/adapters/level_zero/program.cpp | 9 ++- source/adapters/native_cpu/program.cpp | 9 ++- source/adapters/null/ur_nullddi.cpp | 6 ++ source/adapters/opencl/program.cpp | 12 +++- source/loader/layers/tracing/ur_trcddi.cpp | 6 ++ source/loader/layers/validation/ur_valddi.cpp | 6 ++ source/loader/ur_ldrddi.cpp | 26 ++++---- source/loader/ur_libapi.cpp | 16 +++++ source/ur_api.cpp | 10 +++ test/conformance/device_code/CMakeLists.txt | 11 +++- test/conformance/device_code/linker_error.cpp | 30 +++++++++ .../program/program_adapter_native_cpu.match | 3 + test/conformance/program/urProgramLink.cpp | 66 +++++++++++++++++++ 25 files changed, 252 insertions(+), 33 deletions(-) create mode 100644 test/conformance/device_code/linker_error.cpp diff --git a/include/ur_api.h b/include/ur_api.h index 80df5a6fc0..37a1cb6eb9 100644 --- a/include/ur_api.h +++ b/include/ur_api.h @@ -4244,6 +4244,11 @@ urProgramCompile( /// in `phProgram` will contain a binary of the /// ::UR_PROGRAM_BINARY_TYPE_EXECUTABLE type for each device in /// `hContext`. +/// - If a non-success code is returned and `phProgram` is not `nullptr`, it +/// will contain an unspecified program or `nullptr`. Implementations may +/// use the build log of this program (accessible via +/// ::urProgramGetBuildInfo) to provide an error log for the linking +/// failure. /// /// @remarks /// _Analogues_ @@ -9193,6 +9198,11 @@ urProgramCompileExp( /// in `phProgram` will contain a binary of the /// ::UR_PROGRAM_BINARY_TYPE_EXECUTABLE type for each device in /// `phDevices`. +/// - If a non-success code is returned and `phProgram` is not `nullptr`, it +/// will contain an unspecified program or `nullptr`. Implementations may +/// use the build log of this program (accessible via +/// ::urProgramGetBuildInfo) to provide an error log for the linking +/// failure. /// /// @remarks /// _Analogues_ diff --git a/scripts/core/exp-multi-device-compile.yml b/scripts/core/exp-multi-device-compile.yml index b51f938f7e..8ccba8c623 100644 --- a/scripts/core/exp-multi-device-compile.yml +++ b/scripts/core/exp-multi-device-compile.yml @@ -94,6 +94,7 @@ analogue: details: - "The application may call this function from simultaneous threads." - "Following a successful call to this entry point the program returned in `phProgram` will contain a binary of the $X_PROGRAM_BINARY_TYPE_EXECUTABLE type for each device in `phDevices`." + - "If a non-success code is returned and `phProgram` is not `nullptr`, it will contain an unspecified program or `nullptr`. Implementations may use the build log of this program (accessible via $xProgramGetBuildInfo) to provide an error log for the linking failure." params: - type: $x_context_handle_t name: hContext diff --git a/scripts/core/program.yml b/scripts/core/program.yml index 45f7710d68..7bc32faa9d 100644 --- a/scripts/core/program.yml +++ b/scripts/core/program.yml @@ -223,6 +223,7 @@ analogue: details: - "The application may call this function from simultaneous threads." - "Following a successful call to this entry point the program returned in `phProgram` will contain a binary of the $X_PROGRAM_BINARY_TYPE_EXECUTABLE type for each device in `hContext`." + - "If a non-success code is returned and `phProgram` is not `nullptr`, it will contain an unspecified program or `nullptr`. Implementations may use the build log of this program (accessible via $xProgramGetBuildInfo) to provide an error log for the linking failure." params: - type: $x_context_handle_t name: hContext diff --git a/scripts/templates/helper.py b/scripts/templates/helper.py index 1d539d70fe..fce8eea513 100644 --- a/scripts/templates/helper.py +++ b/scripts/templates/helper.py @@ -1212,6 +1212,30 @@ def get_pfntables(specs, meta, namespace, tags): return tables +""" +Public: + returns an expression setting required output parameters to null on entry +""" +def get_initial_null_set(obj): + cname = obj_traits.class_name(obj) + lvalue = { + ('$xProgram', 'Link'): 'phProgram', + ('$xProgram', 'LinkExp'): 'phProgram', + }.get((cname, obj['name'])) + if lvalue is not None: + return 'if (nullptr != {0}) {{*{0} = nullptr;}}'.format(lvalue) + return "" + +""" +Public: + returns true if the function always wraps output pointers in loader handles +""" +def always_wrap_outputs(obj): + cname = obj_traits.class_name(obj) + return (cname, obj['name']) in [ + ('$xProgram', 'Link'), + ('$xProgram', 'LinkExp'), + ] """ Private: diff --git a/scripts/templates/ldrddi.cpp.mako b/scripts/templates/ldrddi.cpp.mako index bbc7c7c7d0..cd8de18df6 100644 --- a/scripts/templates/ldrddi.cpp.mako +++ b/scripts/templates/ldrddi.cpp.mako @@ -49,7 +49,7 @@ namespace ur_loader { ${x}_result_t result = ${X}_RESULT_SUCCESS;<% add_local = False - %> + %>${th.get_initial_null_set(obj)} %if re.match(r"\w+AdapterGet$", th.make_func_name(n, tags, obj)): @@ -271,7 +271,7 @@ namespace ur_loader del add_local %> %for i, item in enumerate(epilogue): - %if 0 == i: + %if 0 == i and not th.always_wrap_outputs(obj): if( ${X}_RESULT_SUCCESS != result ) return result; @@ -307,7 +307,7 @@ namespace ur_loader ${item['factory']}.getInstance( ${item['name']}[ i ], dditable ) ); %else: // convert platform handle to loader handle - %if item['optional']: + %if item['optional'] or th.always_wrap_outputs(obj): if( nullptr != ${item['name']} ) *${item['name']} = reinterpret_cast<${item['type']}>( ${item['factory']}.getInstance( *${item['name']}, dditable ) ); diff --git a/scripts/templates/libapi.cpp.mako b/scripts/templates/libapi.cpp.mako index 6fe1f3992b..881459891a 100644 --- a/scripts/templates/libapi.cpp.mako +++ b/scripts/templates/libapi.cpp.mako @@ -77,7 +77,7 @@ try { %elif th.obj_traits.is_loader_only(obj): return ur_lib::${th.make_func_name(n, tags, obj)}(${", ".join(th.make_param_lines(n, tags, obj, format=["name"]))} ); %else: - auto ${th.make_pfn_name(n, tags, obj)} = ${x}_lib::context->${n}DdiTable.${th.get_table_name(n, tags, obj)}.${th.make_pfn_name(n, tags, obj)}; + ${th.get_initial_null_set(obj)}auto ${th.make_pfn_name(n, tags, obj)} = ${x}_lib::context->${n}DdiTable.${th.get_table_name(n, tags, obj)}.${th.make_pfn_name(n, tags, obj)}; if( nullptr == ${th.make_pfn_name(n, tags, obj)} ) return ${X}_RESULT_ERROR_UNINITIALIZED; diff --git a/scripts/templates/nullddi.cpp.mako b/scripts/templates/nullddi.cpp.mako index f503d4073c..2adb62e691 100644 --- a/scripts/templates/nullddi.cpp.mako +++ b/scripts/templates/nullddi.cpp.mako @@ -38,6 +38,7 @@ namespace driver ) try { ${x}_result_t result = ${X}_RESULT_SUCCESS; + ${th.get_initial_null_set(obj)} // if the driver has created a custom function, then call it instead of using the generic path auto ${th.make_pfn_name(n, tags, obj)} = d_context.${n}DdiTable.${th.get_table_name(n, tags, obj)}.${th.make_pfn_name(n, tags, obj)}; diff --git a/scripts/templates/trcddi.cpp.mako b/scripts/templates/trcddi.cpp.mako index 6f6579d5ac..a9cacd1993 100644 --- a/scripts/templates/trcddi.cpp.mako +++ b/scripts/templates/trcddi.cpp.mako @@ -36,7 +36,7 @@ namespace ur_tracing_layer ${line} %endfor ) - { + {${th.get_initial_null_set(obj)} auto ${th.make_pfn_name(n, tags, obj)} = context.${n}DdiTable.${th.get_table_name(n, tags, obj)}.${th.make_pfn_name(n, tags, obj)}; if( nullptr == ${th.make_pfn_name(n, tags, obj)} ) diff --git a/scripts/templates/valddi.cpp.mako b/scripts/templates/valddi.cpp.mako index c8905a7e8b..e88c391fe3 100644 --- a/scripts/templates/valddi.cpp.mako +++ b/scripts/templates/valddi.cpp.mako @@ -46,7 +46,7 @@ namespace ur_validation_layer ${line} %endfor ) - { + {${th.get_initial_null_set(obj)} auto ${th.make_pfn_name(n, tags, obj)} = context.${n}DdiTable.${th.get_table_name(n, tags, obj)}.${th.make_pfn_name(n, tags, obj)}; if( nullptr == ${th.make_pfn_name(n, tags, obj)} ) { diff --git a/source/adapters/cuda/program.cpp b/source/adapters/cuda/program.cpp index c5cb763ada..784a8e04ae 100644 --- a/source/adapters/cuda/program.cpp +++ b/source/adapters/cuda/program.cpp @@ -265,7 +265,10 @@ UR_APIEXPORT ur_result_t UR_APICALL urProgramBuild(ur_context_handle_t hContext, UR_APIEXPORT ur_result_t UR_APICALL urProgramLinkExp( ur_context_handle_t, uint32_t, ur_device_handle_t *, uint32_t, - const ur_program_handle_t *, const char *, ur_program_handle_t *) { + const ur_program_handle_t *, const char *, ur_program_handle_t *phProgram) { + if (nullptr != phProgram) { + *phProgram = nullptr; + } return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } @@ -277,6 +280,10 @@ urProgramLink(ur_context_handle_t hContext, uint32_t count, const ur_program_handle_t *phPrograms, const char *pOptions, ur_program_handle_t *phProgram) { ur_result_t Result = UR_RESULT_SUCCESS; + if (nullptr != phProgram) { + *phProgram = nullptr; + } + // All programs must be associated with the same device for (auto i = 1u; i < count; ++i) UR_ASSERT(phPrograms[i]->getDevice() == phPrograms[0]->getDevice(), diff --git a/source/adapters/hip/program.cpp b/source/adapters/hip/program.cpp index 8e3653ee02..681a2a0ec2 100644 --- a/source/adapters/hip/program.cpp +++ b/source/adapters/hip/program.cpp @@ -321,14 +321,19 @@ UR_APIEXPORT ur_result_t UR_APICALL urProgramBuild(ur_context_handle_t, UR_APIEXPORT ur_result_t UR_APICALL urProgramLinkExp( ur_context_handle_t, uint32_t, ur_device_handle_t *, uint32_t, - const ur_program_handle_t *, const char *, ur_program_handle_t *) { + const ur_program_handle_t *, const char *, ur_program_handle_t *phProgram) { + if (nullptr != phProgram) { + *phProgram = nullptr; + } return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } -UR_APIEXPORT ur_result_t UR_APICALL urProgramLink(ur_context_handle_t, uint32_t, - const ur_program_handle_t *, - const char *, - ur_program_handle_t *) { +UR_APIEXPORT ur_result_t UR_APICALL +urProgramLink(ur_context_handle_t, uint32_t, const ur_program_handle_t *, + const char *, ur_program_handle_t *phProgram) { + if (nullptr != phProgram) { + *phProgram = nullptr; + } return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } diff --git a/source/adapters/level_zero/common.cpp b/source/adapters/level_zero/common.cpp index 926d5f4ba4..6b2b6ec82f 100644 --- a/source/adapters/level_zero/common.cpp +++ b/source/adapters/level_zero/common.cpp @@ -61,6 +61,8 @@ ur_result_t ze2urResult(ze_result_t ZeResult) { return UR_RESULT_ERROR_OUT_OF_HOST_MEMORY; case ZE_RESULT_ERROR_UNSUPPORTED_FEATURE: return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; + case ZE_RESULT_ERROR_MODULE_LINK_FAILURE: + return UR_RESULT_ERROR_PROGRAM_LINK_FAILURE; default: return UR_RESULT_ERROR_UNKNOWN; } diff --git a/source/adapters/level_zero/program.cpp b/source/adapters/level_zero/program.cpp index 447721f004..aa8f7d4c4e 100644 --- a/source/adapters/level_zero/program.cpp +++ b/source/adapters/level_zero/program.cpp @@ -300,6 +300,9 @@ UR_APIEXPORT ur_result_t UR_APICALL urProgramLinkExp( ur_program_handle_t *phProgram ///< [out] pointer to handle of program object created. ) { + if (nullptr != phProgram) { + *phProgram = nullptr; + } for (uint32_t i = 0; i < numDevices; i++) { UR_ASSERT(hContext->isValidDevice(phDevices[i]), UR_RESULT_ERROR_INVALID_DEVICE); @@ -445,11 +448,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urProgramLinkExp( // because the ZeBuildLog tells which symbols are unresolved. if (ZeResult == ZE_RESULT_SUCCESS) { ZeResult = checkUnresolvedSymbols(ZeModule, &ZeBuildLog); - if (ZeResult == ZE_RESULT_ERROR_MODULE_LINK_FAILURE) { - UrResult = - UR_RESULT_ERROR_UNKNOWN; // TODO: - // UR_RESULT_ERROR_PROGRAM_LINK_FAILURE; - } else if (ZeResult != ZE_RESULT_SUCCESS) { + if (ZeResult != ZE_RESULT_SUCCESS) { return ze2urResult(ZeResult); } } diff --git a/source/adapters/native_cpu/program.cpp b/source/adapters/native_cpu/program.cpp index ee21a707b7..77edd83bce 100644 --- a/source/adapters/native_cpu/program.cpp +++ b/source/adapters/native_cpu/program.cpp @@ -119,11 +119,13 @@ UR_APIEXPORT ur_result_t UR_APICALL urProgramLink(ur_context_handle_t hContext, uint32_t count, const ur_program_handle_t *phPrograms, const char *pOptions, ur_program_handle_t *phProgram) { + if (nullptr != phProgram) { + *phProgram = nullptr; + } std::ignore = hContext; std::ignore = count; std::ignore = phPrograms; std::ignore = pOptions; - std::ignore = phProgram; DIE_NO_IMPLEMENTATION } @@ -144,7 +146,10 @@ UR_APIEXPORT ur_result_t UR_APICALL urProgramBuildExp(ur_program_handle_t, UR_APIEXPORT ur_result_t UR_APICALL urProgramLinkExp( ur_context_handle_t, uint32_t, ur_device_handle_t *, uint32_t, - const ur_program_handle_t *, const char *, ur_program_handle_t *) { + const ur_program_handle_t *, const char *, ur_program_handle_t *phProgram) { + if (nullptr != phProgram) { + *phProgram = nullptr; + } return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } diff --git a/source/adapters/null/ur_nullddi.cpp b/source/adapters/null/ur_nullddi.cpp index a713a385a7..0401385205 100644 --- a/source/adapters/null/ur_nullddi.cpp +++ b/source/adapters/null/ur_nullddi.cpp @@ -1914,6 +1914,9 @@ __urdlllocal ur_result_t UR_APICALL urProgramLink( *phProgram ///< [out] pointer to handle of program object created. ) try { ur_result_t result = UR_RESULT_SUCCESS; + if (nullptr != phProgram) { + *phProgram = nullptr; + } // if the driver has created a custom function, then call it instead of using the generic path auto pfnLink = d_context.urDdiTable.Program.pfnLink; @@ -5700,6 +5703,9 @@ __urdlllocal ur_result_t UR_APICALL urProgramLinkExp( *phProgram ///< [out] pointer to handle of program object created. ) try { ur_result_t result = UR_RESULT_SUCCESS; + if (nullptr != phProgram) { + *phProgram = nullptr; + } // if the driver has created a custom function, then call it instead of using the generic path auto pfnLinkExp = d_context.urDdiTable.ProgramExp.pfnLinkExp; diff --git a/source/adapters/opencl/program.cpp b/source/adapters/opencl/program.cpp index 14f4da22ec..8b59f059d3 100644 --- a/source/adapters/opencl/program.cpp +++ b/source/adapters/opencl/program.cpp @@ -215,6 +215,13 @@ urProgramLink(ur_context_handle_t hContext, uint32_t count, pOptions, cl_adapter::cast(count), cl_adapter::cast(phPrograms), nullptr, nullptr, &CLResult)); + + if (CL_INVALID_BINARY == CLResult) { + // Some OpenCL drivers incorrectly return CL_INVALID_BINARY here, convert it + // to CL_LINK_PROGRAM_FAILURE + CLResult = CL_LINK_PROGRAM_FAILURE; + } + CL_RETURN_ON_FAILURE(CLResult); return UR_RESULT_SUCCESS; @@ -236,7 +243,10 @@ UR_APIEXPORT ur_result_t UR_APICALL urProgramBuildExp(ur_program_handle_t, UR_APIEXPORT ur_result_t UR_APICALL urProgramLinkExp( ur_context_handle_t, uint32_t, ur_device_handle_t *, uint32_t, - const ur_program_handle_t *, const char *, ur_program_handle_t *) { + const ur_program_handle_t *, const char *, ur_program_handle_t *phProgram) { + if (nullptr != phProgram) { + *phProgram = nullptr; + } return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } diff --git a/source/loader/layers/tracing/ur_trcddi.cpp b/source/loader/layers/tracing/ur_trcddi.cpp index da61c34992..78df89c233 100644 --- a/source/loader/layers/tracing/ur_trcddi.cpp +++ b/source/loader/layers/tracing/ur_trcddi.cpp @@ -2480,6 +2480,9 @@ __urdlllocal ur_result_t UR_APICALL urProgramLink( ur_program_handle_t *phProgram ///< [out] pointer to handle of program object created. ) { + if (nullptr != phProgram) { + *phProgram = nullptr; + } auto pfnLink = context.urDdiTable.Program.pfnLink; if (nullptr == pfnLink) { @@ -7607,6 +7610,9 @@ __urdlllocal ur_result_t UR_APICALL urProgramLinkExp( ur_program_handle_t *phProgram ///< [out] pointer to handle of program object created. ) { + if (nullptr != phProgram) { + *phProgram = nullptr; + } auto pfnLinkExp = context.urDdiTable.ProgramExp.pfnLinkExp; if (nullptr == pfnLinkExp) { diff --git a/source/loader/layers/validation/ur_valddi.cpp b/source/loader/layers/validation/ur_valddi.cpp index 6435cc24e1..d785947e83 100644 --- a/source/loader/layers/validation/ur_valddi.cpp +++ b/source/loader/layers/validation/ur_valddi.cpp @@ -2822,6 +2822,9 @@ __urdlllocal ur_result_t UR_APICALL urProgramLink( ur_program_handle_t *phProgram ///< [out] pointer to handle of program object created. ) { + if (nullptr != phProgram) { + *phProgram = nullptr; + } auto pfnLink = context.urDdiTable.Program.pfnLink; if (nullptr == pfnLink) { @@ -9228,6 +9231,9 @@ __urdlllocal ur_result_t UR_APICALL urProgramLinkExp( ur_program_handle_t *phProgram ///< [out] pointer to handle of program object created. ) { + if (nullptr != phProgram) { + *phProgram = nullptr; + } auto pfnLinkExp = context.urDdiTable.ProgramExp.pfnLinkExp; if (nullptr == pfnLinkExp) { diff --git a/source/loader/ur_ldrddi.cpp b/source/loader/ur_ldrddi.cpp index fb392dd607..f7a8d39bf4 100644 --- a/source/loader/ur_ldrddi.cpp +++ b/source/loader/ur_ldrddi.cpp @@ -2551,6 +2551,9 @@ __urdlllocal ur_result_t UR_APICALL urProgramLink( *phProgram ///< [out] pointer to handle of program object created. ) { ur_result_t result = UR_RESULT_SUCCESS; + if (nullptr != phProgram) { + *phProgram = nullptr; + } // extract platform's function pointer table auto dditable = reinterpret_cast(hContext)->dditable; @@ -2573,14 +2576,12 @@ __urdlllocal ur_result_t UR_APICALL urProgramLink( result = pfnLink(hContext, count, phProgramsLocal.data(), pOptions, phProgram); - if (UR_RESULT_SUCCESS != result) { - return result; - } - try { // convert platform handle to loader handle - *phProgram = reinterpret_cast( - ur_program_factory.getInstance(*phProgram, dditable)); + if (nullptr != phProgram) { + *phProgram = reinterpret_cast( + ur_program_factory.getInstance(*phProgram, dditable)); + } } catch (std::bad_alloc &) { result = UR_RESULT_ERROR_OUT_OF_HOST_MEMORY; } @@ -7884,6 +7885,9 @@ __urdlllocal ur_result_t UR_APICALL urProgramLinkExp( *phProgram ///< [out] pointer to handle of program object created. ) { ur_result_t result = UR_RESULT_SUCCESS; + if (nullptr != phProgram) { + *phProgram = nullptr; + } // extract platform's function pointer table auto dditable = reinterpret_cast(hContext)->dditable; @@ -7913,14 +7917,12 @@ __urdlllocal ur_result_t UR_APICALL urProgramLinkExp( result = pfnLinkExp(hContext, numDevices, phDevicesLocal.data(), count, phProgramsLocal.data(), pOptions, phProgram); - if (UR_RESULT_SUCCESS != result) { - return result; - } - try { // convert platform handle to loader handle - *phProgram = reinterpret_cast( - ur_program_factory.getInstance(*phProgram, dditable)); + if (nullptr != phProgram) { + *phProgram = reinterpret_cast( + ur_program_factory.getInstance(*phProgram, dditable)); + } } catch (std::bad_alloc &) { result = UR_RESULT_ERROR_OUT_OF_HOST_MEMORY; } diff --git a/source/loader/ur_libapi.cpp b/source/loader/ur_libapi.cpp index 2fa318e71c..8d31a7a8ee 100644 --- a/source/loader/ur_libapi.cpp +++ b/source/loader/ur_libapi.cpp @@ -3083,6 +3083,11 @@ ur_result_t UR_APICALL urProgramCompile( /// in `phProgram` will contain a binary of the /// ::UR_PROGRAM_BINARY_TYPE_EXECUTABLE type for each device in /// `hContext`. +/// - If a non-success code is returned and `phProgram` is not `nullptr`, it +/// will contain an unspecified program or `nullptr`. Implementations may +/// use the build log of this program (accessible via +/// ::urProgramGetBuildInfo) to provide an error log for the linking +/// failure. /// /// @remarks /// _Analogues_ @@ -3114,6 +3119,9 @@ ur_result_t UR_APICALL urProgramLink( ur_program_handle_t *phProgram ///< [out] pointer to handle of program object created. ) try { + if (nullptr != phProgram) { + *phProgram = nullptr; + } auto pfnLink = ur_lib::context->urDdiTable.Program.pfnLink; if (nullptr == pfnLink) { return UR_RESULT_ERROR_UNINITIALIZED; @@ -8600,6 +8608,11 @@ ur_result_t UR_APICALL urProgramCompileExp( /// in `phProgram` will contain a binary of the /// ::UR_PROGRAM_BINARY_TYPE_EXECUTABLE type for each device in /// `phDevices`. +/// - If a non-success code is returned and `phProgram` is not `nullptr`, it +/// will contain an unspecified program or `nullptr`. Implementations may +/// use the build log of this program (accessible via +/// ::urProgramGetBuildInfo) to provide an error log for the linking +/// failure. /// /// @remarks /// _Analogues_ @@ -8635,6 +8648,9 @@ ur_result_t UR_APICALL urProgramLinkExp( ur_program_handle_t *phProgram ///< [out] pointer to handle of program object created. ) try { + if (nullptr != phProgram) { + *phProgram = nullptr; + } auto pfnLinkExp = ur_lib::context->urDdiTable.ProgramExp.pfnLinkExp; if (nullptr == pfnLinkExp) { return UR_RESULT_ERROR_UNINITIALIZED; diff --git a/source/ur_api.cpp b/source/ur_api.cpp index b8496a83c7..ad89b03b0b 100644 --- a/source/ur_api.cpp +++ b/source/ur_api.cpp @@ -2626,6 +2626,11 @@ ur_result_t UR_APICALL urProgramCompile( /// in `phProgram` will contain a binary of the /// ::UR_PROGRAM_BINARY_TYPE_EXECUTABLE type for each device in /// `hContext`. +/// - If a non-success code is returned and `phProgram` is not `nullptr`, it +/// will contain an unspecified program or `nullptr`. Implementations may +/// use the build log of this program (accessible via +/// ::urProgramGetBuildInfo) to provide an error log for the linking +/// failure. /// /// @remarks /// _Analogues_ @@ -7280,6 +7285,11 @@ ur_result_t UR_APICALL urProgramCompileExp( /// in `phProgram` will contain a binary of the /// ::UR_PROGRAM_BINARY_TYPE_EXECUTABLE type for each device in /// `phDevices`. +/// - If a non-success code is returned and `phProgram` is not `nullptr`, it +/// will contain an unspecified program or `nullptr`. Implementations may +/// use the build log of this program (accessible via +/// ::urProgramGetBuildInfo) to provide an error log for the linking +/// failure. /// /// @remarks /// _Analogues_ diff --git a/test/conformance/device_code/CMakeLists.txt b/test/conformance/device_code/CMakeLists.txt index 1419604b9d..dba61c0a06 100644 --- a/test/conformance/device_code/CMakeLists.txt +++ b/test/conformance/device_code/CMakeLists.txt @@ -100,7 +100,15 @@ macro(add_device_binary SOURCE_FILE) continue() endif() - add_custom_command(OUTPUT "${BIN_PATH}" + # cuda and hip seem to do linking at compile time (rather than runtime) + if(${TRIPLE} MATCHES "nvptx" AND ${KERNEL_NAME} MATCHES "linker_error") + continue() + endif() + if(${TRIPLE} MATCHES "amd" AND ${KERNEL_NAME} MATCHES "linker_error") + continue() + endif() + + add_custom_command(OUTPUT ${BIN_PATH} COMMAND ${UR_DPCXX} -fsycl -fsycl-targets=${TRIPLE} -fsycl-device-code-split=off ${AMD_TARGET_BACKEND} ${AMD_OFFLOAD_ARCH} ${AMD_NOGPULIB} ${DPCXX_BUILD_FLAGS_LIST} ${SOURCE_FILE} -o ${EXE_PATH} @@ -149,6 +157,7 @@ add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/build_failure.cpp) add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/fixed_wg_size.cpp) add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/standard_types.cpp) add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/subgroup.cpp) +add_device_binary(${CMAKE_CURRENT_SOURCE_DIR}/linker_error.cpp) set(KERNEL_HEADER ${UR_CONFORMANCE_DEVICE_BINARIES_DIR}/kernel_entry_points.h) add_custom_command(OUTPUT ${KERNEL_HEADER} diff --git a/test/conformance/device_code/linker_error.cpp b/test/conformance/device_code/linker_error.cpp new file mode 100644 index 0000000000..5fc7eebf6f --- /dev/null +++ b/test/conformance/device_code/linker_error.cpp @@ -0,0 +1,30 @@ +// Copyright (C) 2024 Intel Corporation +// Part of the Unified-Runtime Project, under the Apache License v2.0 with LLVM Exceptions. +// See LICENSE.TXT +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception + +#include + +SYCL_EXTERNAL void this_function_does_not_exist(); + +int main() { + cl::sycl::queue deviceQueue; + cl::sycl::range<1> numOfItems{1}; + + try { + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto kern = [=](cl::sycl::id<1>) { +#ifdef __SYCL_DEVICE_ONLY__ + this_function_does_not_exist(); +#endif + }; + cgh.parallel_for(numOfItems, kern); + }); + std::cout << "Expected an error compiling the program." << std::endl; + abort(); + } catch (sycl::exception &e) { + // OK + } + + return 0; +} diff --git a/test/conformance/program/program_adapter_native_cpu.match b/test/conformance/program/program_adapter_native_cpu.match index 9a5b0a9830..1f3071a745 100644 --- a/test/conformance/program/program_adapter_native_cpu.match +++ b/test/conformance/program/program_adapter_native_cpu.match @@ -127,6 +127,9 @@ {{OPT}}urProgramLinkTest.InvalidNullPointerProgram/SYCL_NATIVE_CPU___SYCL_Native_CPU_ {{OPT}}urProgramLinkTest.InvalidNullPointerInputPrograms/SYCL_NATIVE_CPU___SYCL_Native_CPU_ {{OPT}}urProgramLinkTest.InvalidSizeCount/SYCL_NATIVE_CPU___SYCL_Native_CPU_ +{{OPT}}urProgramLinkTest.SetOutputOnZeroCount/SYCL_NATIVE_CPU___SYCL_Native_CPU_ +{{OPT}}urProgramLinkErrorTest.LinkFailure/SYCL_NATIVE_CPU___SYCL_Native_CPU_ +{{OPT}}urProgramLinkErrorTest.SetOutputOnLinkError/SYCL_NATIVE_CPU___SYCL_Native_CPU_ {{OPT}}urProgramReleaseTest.Success/SYCL_NATIVE_CPU___SYCL_Native_CPU_ {{OPT}}urProgramReleaseTest.InvalidNullHandleProgram/SYCL_NATIVE_CPU___SYCL_Native_CPU_ {{OPT}}urProgramRetainTest.Success/SYCL_NATIVE_CPU___SYCL_Native_CPU_ diff --git a/test/conformance/program/urProgramLink.cpp b/test/conformance/program/urProgramLink.cpp index ec2fb36073..e14c38d883 100644 --- a/test/conformance/program/urProgramLink.cpp +++ b/test/conformance/program/urProgramLink.cpp @@ -31,6 +31,46 @@ struct urProgramLinkTest : uur::urProgramTest { }; UUR_INSTANTIATE_KERNEL_TEST_SUITE_P(urProgramLinkTest); +struct urProgramLinkErrorTest : uur::urQueueTest { + const std::string linker_error_program_name = "linker_error"; + + void SetUp() override { + UUR_RETURN_ON_FATAL_FAILURE(urQueueTest::SetUp()); + // TODO: This should use a query for urProgramCreateWithIL support or + // rely on UR_RESULT_ERROR_UNSUPPORTED_FEATURE being returned. + ur_platform_backend_t backend; + ASSERT_SUCCESS(urPlatformGetInfo(platform, UR_PLATFORM_INFO_BACKEND, + sizeof(ur_platform_backend_t), + &backend, nullptr)); + if (backend == UR_PLATFORM_BACKEND_HIP) { + GTEST_SKIP(); + } + // Don't know how to produce alinker error on CUDA + if (backend == UR_PLATFORM_BACKEND_CUDA) { + GTEST_SKIP(); + } + + std::shared_ptr> il_binary{}; + UUR_RETURN_ON_FATAL_FAILURE( + uur::KernelsEnvironment::instance->LoadSource( + linker_error_program_name, il_binary)); + ASSERT_SUCCESS(uur::KernelsEnvironment::instance->CreateProgram( + platform, context, device, *il_binary, nullptr, &program)); + ASSERT_SUCCESS(urProgramCompile(context, program, nullptr)); + } + + void TearDown() override { + if (linked_program) { + EXPECT_SUCCESS(urProgramRelease(linked_program)); + } + UUR_RETURN_ON_FATAL_FAILURE(urQueueTest::TearDown()); + } + + ur_program_handle_t program = nullptr; + ur_program_handle_t linked_program = nullptr; +}; +UUR_INSTANTIATE_KERNEL_TEST_SUITE_P(urProgramLinkErrorTest); + TEST_P(urProgramLinkTest, Success) { ASSERT_SUCCESS( urProgramLink(context, 1, &program, nullptr, &linked_program)); @@ -63,3 +103,29 @@ TEST_P(urProgramLinkTest, InvalidSizeCount) { UR_RESULT_ERROR_INVALID_SIZE, urProgramLink(context, 0, &program, nullptr, &linked_program)); } + +TEST_P(urProgramLinkErrorTest, LinkFailure) { + ASSERT_EQ_RESULT( + UR_RESULT_ERROR_PROGRAM_LINK_FAILURE, + urProgramLink(context, 1, &program, nullptr, &linked_program)); +} + +TEST_P(urProgramLinkTest, SetOutputOnZeroCount) { + uintptr_t invalid_pointer; + linked_program = reinterpret_cast(&invalid_pointer); + ASSERT_EQ_RESULT( + UR_RESULT_ERROR_INVALID_SIZE, + urProgramLink(context, 0, &program, nullptr, &linked_program)); + ASSERT_NE(linked_program, + reinterpret_cast(&invalid_pointer)); +} + +TEST_P(urProgramLinkErrorTest, SetOutputOnLinkError) { + uintptr_t invalid_pointer; + linked_program = reinterpret_cast(&invalid_pointer); + ASSERT_EQ_RESULT( + UR_RESULT_ERROR_PROGRAM_LINK_FAILURE, + urProgramLink(context, 1, &program, nullptr, &linked_program)); + ASSERT_NE(linked_program, + reinterpret_cast(&invalid_pointer)); +}