Memory semantic synchronization

The memory semantic synchronization mechanism allows you to implement synchronization based on the general-purpose device memory. Different from the Event/Notify synchronization mechanism, the synchronization mechanism based on memory semantic allows operators to participate in synchronization. That is, an operator can be synchronized with another stream during execution.

The following figure shows the synchronization process between an operator and another stream during operator execution.

The following is the sample code for calling APIs related to memory semantic synchronization, which is for reference only and cannot be directly copied for compilation and running.

Sample code on the device (operator kernel function implementation code)

extern "C" __global__ __aicore__ void myKernel1(GM_ADDR syncMem)
{
    // Operator logic
    ......
    // Write 1 to the memory pointed to by syncMem.
    __gm__  uint64_t* flag = reinterpret_cast<__gm__ uint64_t*>(syncMem);
    *flag = 1;
    dcci(flag, 0, 2);
    // Operator logic
    ......
}
extern "C" __global__ __aicore__ void myKernel2(GM_ADDR syncMem)
{
    // Operator logic
    ...
    // Polling blocking until the value of the memory pointed to by syncMem is 2.
    __gm__ volatile uint64_t* flag = reinterpret_cast<__gm__ uint64_t*>(syncMem);
    dcci(flag, 0, 2);
    while (*flag != 2) {
        dcci(flag, 0, 2);
    }
    // Operator logic
    ......
}

Sample code on the host

// Create a stream.
aclrtStream stream1;
aclrtStream stream2;
aclrtCreateStream(&stream1);
aclrtCreateStream(&stream2);

// Allocate the device memory.
void* syncMem;
aclrtMalloc(&syncMem, sizeof(uint64_t), ACL_MEM_MALLOC_NORMAL_ONLY);
// Deliver a wait task to stream 1. The task is blocked and waits until the value in the memory pointed to by syncMem is 1.
aclrtValueWait(syncMem, 1, ACL_STREAM_WAIT_VALUE_EQ, stream1);
// Deliver myKernel1 to stream 2. The kernel writes 1 to the memory pointed to by syncMem to unblock the wait task on stream 1.
myKernel1<<<numBlocks, nullptr, stream2>>>(syncMem);

// Deliver myKernel2 to stream 1. The kernel internally polls and waits until the value in the memory pointed to by syncMem is 2.
myKernel2<<<numBlocks, nullptr, stream1>>>(syncMem);
// Deliver a write task to stream 2. The task writes 2 to the memory pointed to by syncMem to unblock myKernel2 on stream 1.
aclrtValueWrite(syncMem, 2, 0, stream2);

Note: The memory semantic synchronization mechanism is implemented based on the general-purpose device memory. Therefore, you can use aclrtMemset/aclrtMemsetAsync to initialize and clear the memory used for synchronization.

Parent topic: Synchronization Management