Istruzioni HLO asincroni

1. Aggiungere operazioni asincrone a HLO è complicata (ad es. all-reduce-start e all-reduce-done).
2. La suddivisione iniziale e completata potrebbe essere inadeguata per parte dell'uso asincrono d'uso diversi.

Per colmare il primo problema, proponiamo di introdurre un'ultima serie di nuovi codici operativi asincroni: kAsyncStart, kAsyncUpdate e kAsyncDone. L'idea consiste nel creare un opcode generico asincrono in grado di aggregare qualsiasi istruzione HLO. L'operazione effettiva che verrà eseguita in modo asincrono sarà codificata utilizzando chiamato calcolo che ha solo l'istruzione come radice e qualsiasi per gli input. Gestione e aliasing del buffer di input/output in corso possono essere condivisi per qualsiasi operazione asincrona. Il comando async-start la forma di output sarà quindi una tupla di operandi di input, valori di output ed eventuali stato intermedio necessario per async-update o async-done istruzioni.

%async_op {
  %param0 = f32[64] parameter(0)
  ROOT %op = f32[32] op(f32[64] %param0), op_specific_attr=”foo”
}

%async-start = (f32[64], f32[32], s32[]) async-start(f32[64] %operand),
                                         calls=%async_op
%async-done = f32[32] async-done((f32[64], f32[32], s32[]) %async-start)

Nella rappresentazione precedente, solo async-start ha un calcolo chiamato dal è banale scoprire che cosa fa async-done seguendo il suo operando trova il valore async-start corrispondente per trovare il calcolo chiamato.

Nota anche che il primo elemento nella tupla di output di async-start alias con operando, in modo che il buffer rimanga attivo fino almeno fino all'istruzione asincrona. Allo stesso modo, gli alias del secondo elemento con l'output async-done e Il terzo elemento è lo stato del contesto utilizzato per tenere traccia un'operazione asincrona. Questa rappresentazione supporta anche più tensori l'input e/o l'output dell'operazione asincrona e l'aliasing funziona allo stesso modo modo:

%async_op {
  %param0 = f32[64] parameter(0)
  %param1 = f32[64] parameter(1)
  ROOT %op = (f32[32], f32[32]) op(f32[64] %param0, f32[64] %param1),
                                op_specific_attr=”foo”
}

%async-start = ((f32[64], f32[64]), (f32[32], f32[32]), s32[])
               async-start(f32[64] %operand0, f32[64] %operand1),
               calls=%async_op
%async-done = (f32[32], f32[32]) async-done(%async-start)

Inoltre, l'operazione può essere ulteriormente scomposta in zero o più async-update che eseguono calcoli intermedi. L'aliasing di input/output funziona allo stesso modo con l'istruzione async-update e ogni async-start e Le istruzioni di async-update devono avere un utente che sia un altro async-update o async-done:

%async_op {
  %param0 = f32[64] parameter(0)
  ROOT %op = f32[32] op(f32[64] %param0), op_specific_attr=”foo”
}

%async-start = (f32[64], f32[32], s32[]) async-start(f32[64] %operand),
                                         calls=%async_op
%async-update0 = (f32[64], f32[32], s32[]) async-update(
                           (f32[64], f32[32], s32[]) %async-start)
%async-update1 = (f32[64], f32[32], s32[]) async-update(
                           (f32[64], f32[32], s32[]) %async-update0)
%async-done = f32[32] async-done((f32[64], f32[32], s32[]) %async-update1)

Sintassi dello zucchero

Dato che disporre di un calcolo separato per definire l'operazione che eseguita in modo asincrono è un po' ingombrante, proponiamo anche una sintassi stampano e analizzano automaticamente le operazioni asincrone come se fossero di prima qualità opcode. L'idea è trattare i suffissi "-start", "-update" e "-done" in particolare creando automaticamente il calcolo e l'istruzione (senza ) durante l'analisi. Ad esempio, lo snippet di codice riportato sopra può essere nei campi seguenti e i due possono essere analizzati nella stessa rappresentazione:

%op-start = (f32[64], f32[32], s32[]) op-start(f32[64] %operand),
                                      op_specific_attr=”foo”
%op-update0 = (f32[64], f32[32], s32[]) op-update(
                        (f32[64], f32[32], s32[]) %op-start),
                        op_specific_attr=”foo”
%op-update1 = (f32[64], f32[32], s32[]) op-update(
                        (f32[64], f32[32], s32[]) %op-update0)
%op-done = f32[32] op-done((f32[64], f32[32], s32[]) %op-update1)

Per non creare ambiguità, lo strumento di verifica non consente a un'operazione di essere aggregato con async-start se abbiamo definito esplicitamente un codice operativo con i suffissi "-start" e/o "-done". Anche questa è una modalità di interpretazione letterale si schiudono nel caso in cui abbiamo istruzioni che richiedono un trattamento a livello di HLO che non si adatta al modello descritto sopra (ad es. l'input/output dell'aliasing buffer). All'inizio, quindi, copy-start/copy-done, collective-permute-start/collective-permute-done e così via continueranno a utilizzare i rispettivi opcode di prima classe anziché i nuovi Codici operativi async-start/async-done finché non puliamo il codice per rimuoverli i codici operativi "-start"/"-done".