Replacing _Thrd_sleep with _Thrd_sleep_for causes synchronous exception on SYCL applications with openCL

[iyerrakesh]

With #3914, we now expose a "void __stdcall _Thrd_sleep_for" function instead of "CRTIMP2_PURE void __cdecl _Thrd_sleep". When this is defined as an extern SYCL_EXTERNAL function and called inside SYCL kernel, it gives a synchronous exception when the compiled exe is run.

C:\Temp>type repro.cpp
#include <sycl/sycl.hpp>
#include <iostream>
#include <thread>
#include <xtimec.h>
#include <string>
#include <sstream>

#define CPU 1
#define GPU 2
#define OFFLOAD_TYPE CPU


extern SYCL_EXTERNAL _CRTIMP2_PURE _Thrd_id_t __cdecl _Thrd_id();
extern SYCL_EXTERNAL _CRTIMP2_PURE long long __cdecl _Query_perf_frequency();
extern SYCL_EXTERNAL _CRTIMP2_PURE long long __cdecl _Query_perf_counter();
extern SYCL_EXTERNAL _CRTIMP2_PURE long long __cdecl _Xtime_get_ticks();

//extern SYCL_EXTERNAL _CRTIMP2_PURE _Thrd_sleep(const xtime*); // VS2019
extern SYCL_EXTERNAL void __stdcall _Thrd_sleep_for(unsigned long /*ms*/) noexcept; // VS2022
extern SYCL_EXTERNAL int stallThread(int item);

int main()
{

#if OFFLOAD_TYPE == CPU
    constexpr size_t LENGTH = 3;
    // Length of the work-group.
    constexpr int threadCount = LENGTH;

    // Variables input array.
    int input[LENGTH] = { 12, 111, 3 };

    // Thread ids output array.
    std::thread::id threadIds[LENGTH];

    // Variable values output array.
    int variableValues[LENGTH];

    // Create queue with CPU selector.
    sycl::cpu_selector selector;
#elif OFFLOAD_TYPE == GPU
    constexpr size_t LENGTH = 12;
// Length of the work-group.
    constexpr int threadCount = 2;

    // Variables input array.
    int input[threadCount];

    // Initialize input.
    for (unsigned int i = 0; i < threadCount; i++)
        input[i] = i + 100;

    // Thread ids output array.
    int threadIds[threadCount];
    // Variable values output array.
    int variableValues[threadCount];

    // Create queue with GPU selector.
    sycl::gpu_selector selector;
#endif // OFFLOAD_TYPE == CPU

    try {
        sycl::queue q(selector);
        std::stringstream offloadMessageStream;
		offloadMessageStream << "[SYCL] Using device: ["
			<< q.get_device().get_info<sycl::info::device::name>()
			<< "] with backend: [" 
			<< q.get_backend()
			<< "]"
			<< std::endl;
		auto offloadMessage = offloadMessageStream.str();
		std::cout << offloadMessage << std::endl;
        sycl::range data_range{ threadCount }; // OFFLOAD_MESSAGE_BREAKPOINT
        sycl::buffer buffer_in{ input, data_range };
        sycl::buffer buffer_threadIds{ threadIds, data_range };
        sycl::buffer buffer_variableValues{ variableValues, data_range };

        std::chrono::steady_clock::time_point start = std::chrono::steady_clock::time_point::min();
        q.submit([&](sycl::handler& cgh) {
            // Memory accessors for buffers.
            sycl::accessor in(buffer_in, cgh, sycl::read_only);
            sycl::accessor out_threadIds(buffer_threadIds, cgh, sycl::write_only);
            sycl::accessor out_variableValues(buffer_variableValues, cgh, sycl::write_only);

            // Timer start.
            start = std::chrono::high_resolution_clock::now();

            // NOTE: We are assuming that the iterator starts from 0 and increments by 1
            // kernel start.
            cgh.parallel_for(sycl::range{ LENGTH }, [=](sycl::item <1> it) {
#if OFFLOAD_TYPE == CPU
                std::thread::id this_id = std::this_thread::get_id(); // Get current thread id.
                out_threadIds[it.get_id()] = this_id; // Assign thread id to the out_threadIds acessor.
                int currentWorkItem = in[it.get_id()]; // Get current work item value.
                out_variableValues[it.get_id()] = stallThread(in[it.get_id()]); // BREAKPOINT_1
#elif OFFLOAD_TYPE == GPU
                int threadId = 0;
                int index = it.get_id() / 8; // Devide by 8 to get correct index value.
                int currentWorkItem = in[index]; // Get current work item value.
                out_variableValues[index] = currentWorkItem;
                out_threadIds[index] = threadId; // BREAKPOINT_2
#endif // OFFLOAD_TYPE == CPU
                });
            }); // kernel end

         //Wait for queue end execution.
        q.wait_and_throw();

        // Stop the timer, calculate execution time and print it to the console.
        auto stop = std::chrono::high_resolution_clock::now();
        auto duration = std::chrono::duration_cast<std::chrono::milliseconds>(stop - start);
        std::cout << "\nTime taken by kernel: " << duration.count() << " milliseconds\n" << std::endl;
    }
    catch (sycl::exception const& e) {
        std::cout << "fail; synchronous exception occurred: " << e.what() << "\n";
        return -1;
    }

    // Print the thread ids and corresponding variable values.
    std::cout << "Output:" << std::endl;
    for (int i = 0; i < threadCount; i++) {
        std::cout << "Current thread id: " << threadIds[i] << " Value: " << variableValues[i] << "\n";
    }

    // Verify the variableValues output array.
    for (int i = 0; i < threadCount; i++) {
        if (variableValues[i] != input[i]) {
            std::cout << "\nfail; element " << i << " is " << variableValues[i] << std::endl;
            return -1;
        }
    }
    std::cout << "\nsuccess; result is correct.\n";
    return 0;
 }

// NOTE: Below function is used to stall thread to make sure that uniqe thread is used for each work item.
//make a thread sleep for 5 seconds
 extern SYCL_EXTERNAL int stallThread(int item) {
     _Thrd_sleep_for(5000);
     return item;
 }

#Steps to reproduce

Perquisites:
A OpenCL compatible CPU for running the SYCL application.
Install Intel oneAPI base toolkit if using Intel device and for SYCL libraries.

• Use the github repo - https://github.com/intel-rganesh/DPCPPConsoleApplication.
• Open DPCPPConsoleApplication.sln with VS2022 17.12.4.
• Build the solution and try to run the exe or debug the application.

Expected behavior

We should be able to use _Thrd_sleep_for with VS 17.12.4 to stall an thread for a specified duration when called from inside a SYCL kernel.
Sleeping or pausing a kernel execution is still possible in VS2019 with 16.11 where we use "_Thrd_sleep(const xtime*)".

[sylveon]

Tapping into STL implementation details like this is not supported. Do you have a repro where this occurs with std::thread::sleep_until?

[davidmrdavid]

We merged #3914 to fix a standard conformance issue, and to preserve binary compatibility we had to introduce _Thrd_sleep_for. This is an internal STL implementation detail (as noted by the _Ugly_casing), which we do not support replacing.

Naturally, as an OSS project, you're free to modify the internals beyond our implementation, but Microsoft does not provide official support to that.

[iyerrakesh]

We do not have function declaration of sleep_until in stl/inc/xthreads.h. We need to declare the function as SYCL_EXTERNAL before calling it from sycl kernel.

Since we keep the _Thrd_sleep_for for binary compatibility, it would be nice to leave the declaration in stl/inc/xthreads.h too (although I understand why it was preferred to be removed). The alternative to stall a thread from a sycl kernel is to perform a dummy operation with a long workload, but this is not the best solution.

// Perform a dummy operation to simulate a delay
for (int64_t i = 0; i < 100000000; i++)
{
\\Do some operation
}