8. Memory and Atomics#

8.1. `fatom`#

Perform a floating-point atomic operation.

Syntax:

fatom.addrspace<.memorder><.memscope><.cc>.op.type dst, [addr], src0 <, src1>

  .addrspace = { .global, .generic, .shared }
  .memorder = { .relaxed, .release, .acquire, .acq_rel }
  .memscope = { .workgroup, .gpu, .system }
  .cc = { .uc, .L2wb, .L3wb }
  .op = { .add, .sub, .min, .max, .cas }
  .type = { .bf, .hf, .f, .df }


  - addr is an address operand, in the form MemRR or MemRI
  - dst, src0, src1 are registers

Restrictions

bitwidth of dst, src0, and src1 must match the bitwidth of .type
operand src1 is defined only when qualifier .op is .cas
qualifier .cc must not be specified when qualifier .addrspace is .shared
qualifier .memscope can be specified only when qualifier .memorder != .relaxed
qualifier .type cannot be .df unless qualifier .addrspace is .global and qualifier .op is one of [.add, .sub]

PISA Notes

introduced in PISA version 0.1
requires PISA target 100 or compatible

Semantics:

atomic {
  old = *addr
  switch (op) {
    case cas:
      new = (old == src0) ? src1 : old
      break
    case add:
      new = old + src0
      break
    case sub:
      new = old - src0
      break
    case min:
      new = fmin(old, src0)
      break
    case max:
      new = fmax(old, src0)
      break
  }
  *addr = new
  dst = old
}

Notes:

memory access requirements must be observed.
.generic address must point to either .global or .shared address space.
atomic cache controls may be utilized.

**Supported combinations for floating-point atomic operations**#
`.addrspace`	`.op`	`.type`
.global	.min, .max, .cas	.hf, .bf, .f
.global	.add, .sub	.hf, .bf, .f, .df
.shared	.min, .max, .cas, .add, .sub	.hf, .bf, .f
.generic	.min, .max, .cas, .add, .sub	.hf, .bf, .f

Examples:

.reg .32b %dst, %src0, %addr;

// atomic addition of single-precision floating point and memory value
fatom.shared.add.f  %dst, [%addr], %src0;

8.2. `iatom`#

Perform an integer atomic operation.

Syntax:

iatom.addrspace<.memorder><.memscope><.cc>.op.type dst, [addr] <, src0> <, src1>

  .addrspace = { .global, .shared, .generic }
  .memorder = { .relaxed, .release, .acquire, .acq_rel }
  .memscope = { .workgroup, .gpu, .system }
  .cc = { .uc, .L2wb, .L3wb }
  .op = { .add, .sub, .smin, .smax, .umin, .umax, .and, .or, .xor, .xchg, .cas, .inc,
          .dec, .incwrap, .decwrap }
  .type = { .16b, .32b, .64b, .128b }


  - addr is an address operand, in the form MemRR or MemRI
  - dst, src0, src1 are registers

Restrictions

bitwidth of dst, src0, and src1 must match the bitwidth of .type
operand src0 is not defined when qualifier .op is one of [.inc, .dec]
operand src1 is defined only when qualifier .op is .cas
qualifier .type cannot be .16b when qualifier .op is one of [.incwrap, .decwrap]
qualifier .type cannot be .128b unless qualifier .op is one of [.xchg, .cas]
qualifier .type cannot be .64b unless qualifier .op is .cas or qualifier .addrspace is .global
qualifier .cc must not be specified when qualifier .addrspace is .shared
qualifier .memscope can be specified only when qualifier .memorder != .relaxed

PISA Notes

introduced in PISA version 0.1
requires PISA target 100 or compatible

Semantics:

atomic {
  old = *addr
  switch (op) {
    case cas:
      new = (old == src0) ? src1 : old
      break
    case add:
      new = old + src0
      break
    case sub:
      new = old - src0
      break
    case inc:
      new = old + 1
      break
    case dec:
      new = old - 1
      break
    case incwrap:
      new = (old >= src0) ? 0 : (old + 1); // unsigned
      break
    case decwrap:
      new = ((old == 0) || (old > src0)) ? src0 : (old - 1); // unsigned
      break
    case smin:
      new = smin(old, src0)
      break
    case smax:
      new = smax(old, src0)
      break
    case umin:
      new = umin(old, src0)
      break
    case umax:
      new = umax(old, src0)
      break
    case and:
      new = old & src0
      break
    case or:
      new = old | src0
      break
    case xor:
      new = old ^ src0
      break
    case xchg:
      new = src0
      break
  }
  *addr = new
  dst = old
}

Notes:

memory access requirements must be observed.
.generic address must point to either .global or .shared address space.
atomic cache controls may be utilized.

**Supported combinations for integer atomic operations**#
`.addrspace`	`.op`	`.type`
.global	.add, .sub, .smin, .smax, .umin, .umax, .and, .or, .xor, .xchg, .cas, .inc, .dec	.16b, .32b, .64b
.global	.incwrap, .decwrap	.32b, .64b
.shared, .generic	.add, .sub, .smin, .smax, .umin, .umax, .and, .or, .xor, .xchg, .inc, .dec	.16b, .32b
.shared, .generic	.cas	.16b, .32b, .64b
.shared, .generic	.incwrap, .decwrap	.32b
.global, .shared, .generic	.cas, .xchg	.128b

Examples:

.reg .32b %dst, %src0, %addr;

// atomic addition of 32-bit integer and memory value
iatom.shared.add.32b  %dst, [%addr], %src0;

8.3. `ld`#

Load data from memory.

Load data from memory (non-atomic):

ld<.addrspace><.memorder><.cc><.vec>.type dst, [addr]

  .addrspace = { .global, .generic, .shared, .private, .const }
  .memorder = { .weak }
  .cc = { .L1uc.L2uc.L3uc, .L1uc.L2uc.L3c, .L1uc.L2c.L3uc, .L1uc.L2c.L3c,
          .L1c.L2uc.L3uc, .L1c.L2uc.L3c, .L1c.L2c.L3uc, .L1c.L2c.L3c,
          .L1s.L2uc.L3uc, .L1s.L2uc.L3c, .L1s.L2c.L3uc, .L1s.L2c.L3c,
          .ri }
  .vec = { .v2, .v3, .v4, .v8 }
  .type = { .8b, .16b, .32b, .64b }


  - addr is an address operand, in the form MemRR or MemRI
  - dst is a register

Restrictions

bitwidth of dst must match the bitwidth of .type
number of elements of dst must match value of .vec
qualifier .vec cannot be .v3 unless qualifier .type is one of [.32b, .64b]
qualifier .vec cannot be .v8 unless qualifier .type is .32b
qualifier .cc must not be specified when qualifier .addrspace is .shared
qualifier .memorder must not be specified when qualifier .addrspace is one of [.private, .const]

PISA Notes

introduced in PISA version 0.1
requires PISA target 100 or compatible

Load data from memory (atomic):

ld<.addrspace>.memorder<.memscope><.cc>.type dst, [addr]

  .addrspace = { .global, .generic, .shared }
  .memorder = { .relaxed, .acquire, .seq_cst }
  .memscope = { .workgroup, .gpu, .system }
  .cc = { .uc, .L2wb, .L3wb }
  .type = { .16b, .32b, .64b, .128b }


  - addr is an address operand, in the form MemRR or MemRI
  - dst is a register

Restrictions

bitwidth of dst must match the bitwidth of .type
qualifier .cc must not be specified when qualifier .addrspace is .shared
qualifier .memscope can be specified only when qualifier .memorder is one of [.acquire, .seq_cst]
qualifier .type cannot be .64b unless qualifier .addrspace is .global

PISA Notes

introduced in PISA version 0.1
requires PISA target 100 or compatible

Semantics:

scalar load#

dst = *addr

vector load#

for (i = 0; i < element_count(.vec); i++)
   dst[i] = *(addr + i * sizeof(.type))

Notes:

memory access requirements must be observed.
memory data size restrictions must be observed.
.generic address in atomic load must point to either .global or .shared address space.
atomic load may utilize atomic cache controls.

**Cache control definitions for non-atomic load instruction**#
Cache control	Description
.L1c	Cache the data in L1.
.L1uc	Do not cache the data in L1.
.L1s	Cache the data in L1, but make it more likely to be invalidated later.
.L2c	Cache the data in L2.
.L2uc	Do not cache the data in L2.
.L3c	Cache the data in L3.
.L3uc	Do not cache the data in L3.
.ri	Read-invalidate (e.g. last-use) on all levels of cache.

Examples:

.reg .16b %r, %b;
.shared .16b @S[100];
.reg .64b %a, %g;
.reg .v4.32b %v;

// load a 16-bit element into register %r from shared variable @S
// and cache it in all levels
ld.shared.L1c.L2c.L3c.16b %r, [@S + %b];

// load a 32-bit vector of 4 elements into vector %v from global memory
ld.global.v4.32b %v, [%g + 16];

// load a 64-bit element into register %a from generic address
// considering acquire memory ordering
ld.generic.acquire.64b %a, [%g];

8.4. `ld.param`#

Load data from kernel parameter.

Syntax:

ld.param<.vec>.type dst, [vaddr]

  .vec = { .v2, .v3, .v4, .v8 }
  .type = { .8b, .16b, .32b, .64b }


  - dst is a register

Restrictions

bitwidth of dst must match the bitwidth of .type
number of elements of dst must match value of .vec
qualifier .vec cannot be .v3 unless qualifier .type is one of [.32b, .64b]
qualifier .vec cannot be .v8 unless qualifier .type is .32b

PISA Notes

introduced in PISA version 0.1
requires PISA target 100 or compatible

Notes:

ld.param supports the following address syntax:

[@var] - references a kernel argument
[@var + %reg] - references a kernel argument with a signed 32-bit offset specified by a 32-bit register
[@var + imm] - references a kernel argument with a signed 32-bit offset specified by an immediate

Examples:

.kernel @foo(.param[4] %arg0, .param[8] %arg1) {
  .reg .32b %larg0, %larg1;

  // load contents of %arg0
  ld.param.32b %larg0, [%arg0];
  // load contents of %arg1[4]
  ld.param.32b %larg1, [%arg1 + 4];
}

8.5. `st`#

Store data to memory.

Store data to memory (non-atomic):

st<.addrspace><.memorder><.cc><.vec>.type  [addr], src0

  .addrspace = { .global, .generic, .shared, .private }
  .memorder = { .weak }
  .cc = { .L1uc.L2uc.L3uc, .L1uc.L2uc.L3wb, .L1uc.L2wb.L3uc, .L1uc.L2wb.L3wb,
          .L1wt.L2uc.L3uc, .L1wt.L2uc.L3wb, .L1wt.L2wb.L3uc, .L1wt.L2wb.L3wb,
          .L1s.L2uc.L3uc, .L1s.L2uc.L3wb, .L1s.L2wb.L3uc, .L1wb.L2uc.L3uc,
          .L1wb.L2wb.L3uc, .L1wb.L2uc.L3wb }
  .vec = { .v2, .v3, .v4, .v8 }
  .type = { .8b, .16b, .32b, .64b }


  - addr is an address operand, in the form MemRR or MemRI
  - src0 is a register

Restrictions

bitwidth of src0 must match the bitwidth of .type
number of elements of src0 must match value of .vec
qualifier .vec cannot be .v3 unless qualifier .type is one of [.32b, .64b]
qualifier .vec cannot be .v8 unless qualifier .type is .32b
qualifier .cc must not be specified when qualifier .addrspace is .shared
qualifier .memorder must not be specified when qualifier .addrspace is .private

PISA Notes

introduced in PISA version 0.1
requires PISA target 100 or compatible

Store data to memory (atomic):

st<.addrspace>.memorder<.memscope><.cc>.type  [addr], src0

  .addrspace = { .global, .generic, .shared }
  .memorder = { .relaxed, .release, .seq_cst }
  .memscope = { .workgroup, .gpu, .system }
  .cc = { .uc, .L2wb, .L3wb }
  .type = { .16b, .32b, .64b, .128b }


  - addr is an address operand, in the form MemRR or MemRI
  - src0 is a register

Restrictions

bitwidth of src0 must match the bitwidth of .type
qualifier .cc must not be specified when qualifier .addrspace is .shared
qualifier .memscope can be specified only when qualifier .memorder is one of [.release, .seq_cst]
qualifier .type cannot be .64b unless qualifier .addrspace is .global

PISA Notes

introduced in PISA version 0.1
requires PISA target 100 or compatible

Semantics:

scalar store#

*addr = src0

vector store#

for (i = 0; i < element_count(.vec); i++)
   *(addr + i * sizeof(.type)) = src0[i]

Notes:

memory access requirements must be observed.
memory data size restrictions must be observed.
.generic address in atomic store must point to either .global or .shared address space.
atomic store may utilize atomic cache controls.

**Cache control definitions for non-atomic store instruction**#
Cache control	Description
.L1uc	Bypass L1 and write directly to the next level cache.
.L1wb	Write data to L1 but not the next level cache.
.L1wt	Write data to both L1 and the next level cache.
.L1s	Streaming write to L1.
.L2uc	Bypass L2 and write directly to the next level cache.
.L2wb	Write data to L2 but not the next level cache.
.L3uc	Bypass L3 and write directly to the next level cache.
.L3wb	Write data to L3 but not the next level cache.

Examples:

.reg .16b %r;
.reg .32b %off;
.private .16b @P[64];
.reg .64b %a, %g;
.reg .v2.32b %v;

// store a 16-bit element from register %r
// into private array @P
st.private.16b [@P + %off], %r;

// store 32-bit vector of 2 elements from register %v
// into global address [%g]
st.global.v2.32b [%g], %v;

// store 64-bit element from register %a into generic address [%g]
st.generic.64b [%g], %a;

Memory and Atomics

Contents

8. Memory and Atomics#

8.1. fatom#

8.2. iatom#

8.3. ld#

8.4. ld.param#

8.5. st#

8.1. `fatom`#

8.2. `iatom`#

8.3. `ld`#

8.4. `ld.param`#

8.5. `st`#