[Book] Programming massively parallel processors(PMPP) - Wen-mei Hwu, reading minute - ch1 - ch6

Reference:

CUDA C++ Programming Guide
NVIDIA Nsight Compute CUDA code optimization

[virtual memory] recap

Intel TBB
Intel TBB Task Scheduler
[oneTBB] accessor note

Error handling recap; TLPI
EINTR and What It Is Good For

kernal namespace recap
https://vsdmars.blogspot.com/2018/12/linuxnamespace-wrap-up.html
https://vsdmars.blogspot.com/2018/12/linuxnamespace-mount-mnt.html
https://vsdmars.blogspot.com/2018/06/linuxkernel-namespace.html

cache
[Pacific++ 2018][re-read] "Designing for Efficient Cache Usage" - Scott McMillan
[Go][design] high through-put low contention in memory cache

Locking
[futex] futex skim through
[Concurrency] [C++][Go] Wrap up - 2018
[C++] memory model in depth; memory spilling.

Algorithm/Implementation
Under my github/leetcode project

Software abstraction (CUDA runtime)

Similar idea e.g.  Golang  ref: [Go][note] Analysis of the Go runtime scheduler paper note, the difference coming from Go/C++ as multi-purpose language running on CPU (sequential), CUDA/OpenCL as C extension running on GPU. While the underneath hardware architecture difference, the abstraction diffs.

SPMD single-program multiple-data 

host CPU based code

kernel GPU Device code / function, more details later. Basically same code / IR run in parallel.

grid threads group

_h host variable in CPU code

_d device variable in CPU code

__host__  callable from host, executed on host, executed by host thread(e.g. linux thread).

__global__ callable from host or device, executed on device, executed by grid of device threads

__device__ callable from device, executed on device, executed by caller device thread.

If function declared with __host__ and __device__ macro, NVCC generates two version, one for host and one for device.

block 32-based size(hardware efficiency reason), all blocks are in same size. Size of how many GPU threads. Threads in a block can execute in any order with respect to each other.

SM Streaming multiprocessor; each SM has several processing units called CUDA cores. It is designed to execute all threads in a warp following the single-instruction multiple-data(SIMD) model.

HBM high-bandwidth memory

Warp a warp groups 32-threads together. Thus a block of threads will be group into warps, which each warp has 32-threads. Scheduling is based on Warp. Also think as single-instruction, multiple-threads.

FLOP floating-point operations

FLOP/B FLOP to byte ratio.

GPU global memory bandwidth: 1555GB/second; 1555 * 0.25(FLOP/B) = 389 GFLOPS

const readonly variables

 blockIdx; area code
 blockDim; row idx
 threadIdx; phoneline
Those three variable gives the kernel realize which data it is running on.

OUR_GLOBAL_FUNC<<<number of block, threads per block>>>(args...);

Thread Scheduling

Block scheduling

When a kernel is called, the CUDA runtime launches a grid of threads execute the kernel code. These threads are assigned to SMs on a block-by-block basis. All threads in a block are simultaneously assigned to the same SM. There are reserved blocks for system to executed, thus a SMs' blocks are not all scheduled to the user kernel.
Multiple blocks are likely to be simultaneously assigned to the same SM. The concept of Warp scheduling is that, those threads inside the same Warp runs the same instruction set(same kernel), thus the fetch of instruction is one time efforts. Also, the data those threads in the same Warp access are linear thus are prefetchable/cache friendly.

Moreover, threads in the same block can interact with each other in ways that threads across different blocks cannot, such as barrier synchronization,

synchronization / transparent scalability

__syncthreads()

  block until every thread in the same block reaches the code location. if a __syncthreads() statement is present, it must be executed by all threads in a block. i.e. 

void incorrect_barries_example(int n) {
	if (threadIdx.x % 2) {
		__syncthreads(); // sync point-1
	} else {
		__syncthreads(); // sync point-2
	}
} 

Wrong due to not all threads runs into the same barrier synchronization points.

Not only do all threads in a block have to be assigned to the same SM, but also they need to to be assigned to that SM simultaneously. i.e. a block can begin execution only when the runtime system has secured all the resources needed by all threads in the block to complete execution.

The ability to execute the same application code on different hardware with different amounts of execution resources is referred to as transparent scalability.

Control Divergence

The execution works well when either all threads in a warp execute the if-path or all execute the else-path. Otherwise, it has to go with the code twice. One run with the core running if path code and the other core with else path is doing noop. (In the same Warp). Another run with the core doing noop on the if path code and the other core wile else path is running. In old architecture, those 2 runs run in sequence. In new architecture, those 2 runs can run in parallel. This is called independent thread scheduling.

Thus, due to this fact, do not use threadIdx for if branching. But use data for divergence control, this also related to data locality in cache. One important fact, the performance impact of control divergence decreases as the size of the vectors being processed increases.

One cannot assume that all threads in a warp have the same execution timing.(even they are running the same fetched instruction). Thus, use __syncwarp() barrier synchronization instead.

Latency tolerance

simple, i.e. CPU, context switch on single code due to limited of resource(registers, cache etc.) and makes sure code runs preemptive-scheduling fashion.

Thus, SM only has enough execution units to execute a subset of all the threads assigned to it at any point in time.

In recent SM, each SM can execute instructions for a small number of warps at any given point in time.

GPU SMs achieves zero-overhead scheduling by holding all the execution states for the assigned warps in the hardware registers so there is no need to save and restore states when switching from one warp to another.

Thus, allows GPU oversubscription of threads to SMs.

Automatic/Local variables declared in the kernel are placed into registers.

Each SM in A100 GPU has 65,536 registers.

65536/2048(threads) = 32 registers per thread/kernel.

In cases, the compiler may perform register spilling to reduce the register requirement per thread and thus elevate the level of occupancy. However, this could increase latency due to need to fetch data from the memory instead directly from the register.

Nsight Compute Occupancy Calculator

cudaDeviceProp struct has bunch of variable represents the hardware SPEC.

e.g.

 multiProcessorCount Number of multiprocessors on device

 clockRate Clock frequency in kilohertz

 regsPerBlock 32-bit registers available per block

 warpSize  Warp size in threads

Variable declaration scope and lifetime

automatic scalar variables    [mem]register    [scope]thread    [lifetime]grid

automatic array variables    [mem]local    [scope]thread    [lifetime]grid

__device__ __shared__    [mem]shared    [scope]block    [lifetime]grid

__device__    [mem]global    [scope]grid    [lifetime]application

__device__ __constant_    [mem]constant    [scope]grid    [lifetime]application

API

cudaMalloc

cudaMemcpy

[Book][Database Internals] reading note part I

Reference:
Transactional storage for geo-replicated systems[walter-sosp11]

Transaction manager

This manager schedules transactions and ensures they cannot leave

the database in a logically inconsistent state.

Lock manager

This manager locks on the database objects for the running

transactions, ensuring that concurrent operations do not violate

physical data integrity.

Access methods (storage structures)

These manage access and organizing data on disk. Access methods

include heap files and storage structures such as B-Trees (“Ubiquitous B-Trees”) 

or LSM Trees.

Buffer manager

This manager caches data pages in memory (see “Buffer

Management”).

Recovery manager

This manager maintains the operation log and restoring the system

state in case of a failure.

checkpointing

Backup <-> Log till backup snapshot.

Storage Structure

Disk-based storage structures often have a form of wide and short trees

, while memory-based implementations can choose from a larger pool of data structures and

perform optimizations that would otherwise be impossible or difficult to

implement on disk [MOLINA92].

Similarly, handling variable-size data on disk requires special attention, while in memory it’s often a matter of referencing the value with a pointer.

On modern CPUs, vectorized instructions can be used to process multiple data points

with a single CPU instruction [DREPPER07].

Wide Column Stores

Column-oriented databases should not be mixed up with wide column

stores, such as BigTable or HBase, where data is represented as a

multidimensional map, columns are grouped into column families (usually

storing data of the same type), and inside each column family, data is

stored row-wise. This layout is best for storing data retrieved by a key or a

sequence of keys.

https://jepsen.io/consistency

Serializability and external consistency (SSSSSS)

Even though there is some overlap in time in which Txn1 and Txn2 are both executing, their commit timestamps c1 and c2 respect a linear transaction order, which means that all effects of the reads and writes of Txn1 appear to have occurred at a single point of time (c1), and all effects of the reads and writes of Txn2 appear to have occurred at a single point of time (c2).

Furthermore, c1 < c2 (which is guaranteed because both Txn1 and Txn2 committed writes; this is true even if the writes happened on different machines), which respects the order of Txn1 happening before Txn2. (However, if Txn2 only did reads in the transaction, then c1 <= c2). 

If Run() or Commit() succeeds

When a call to Run() (if using the Transaction Runner API) or Commit() (if using the Transaction API) returns with an OK status, it means the transaction committed at time T, and:

The writes buffered on that transaction are guaranteed to be applied at T.

All rows read by the transaction are guaranteed to be valid at T.
In other words, all the data you read is consistent because it came from a same single snapshot of the database.

If Run() or Commit() fails

If you're using the Transaction Runner API and a call to Run() fails, the read and write guarantees you have depend on what error the underlying Commit() call failed with:

• Deadline Exceeded (if the client set a deadline on the read options) or Canceled (if the client canceled) - the transaction may or may not have committed (and thus buffered writes may or may not have happened).
• A constraint failure (e.g. RowNotFound, AlreadyExists, BadUsage, etc.) - Writing the buffered mutations encountered some error, e.g. a row that the client is trying to update doesn't exist. In that case, the reads are guaranteed consistent, the writes are guaranteed to not be applied, and the non-existence of the row is guaranteed to be consistent with the reads as well.

Nested Transactions Are Unsafe

Potential dead lock could happen.

Use timeout to controll distributed transaction liveness

Long-running Operations SSSSSS servers automatically time out transactions that have been idle for 10s. The SSSSSS client automatically keeps transactions alive while a transactional read or query is running.

Lock Modes

SSSSSS uses a combination of shared locks and exclusive locks to control access to the data.

By default when you perform a read as part of a transaction, SSSSSS acquires shared read locks, which allows other reads to still access the data until your transaction is ready to commit. When your transaction is committing and writes are being applied, the transaction attempts to upgrade to an exclusive lock for any data you are writing. It blocks new shared read locks on the data, waits for existing shared read locks to clear, then places an exclusive lock for exclusive access to the data.

Locks are taken at the granularity of row-and-column. If transaction T1 has locked column "A" of row "foo", and transaction T2 wants to write column "B" of row "foo" then there is no conflict.

Writes to a data item that don't also read the data being written (aka "blind writes") don't conflict with other blind writers of the same item (the commit timestamp of each write determines the order in which it is applied to the database). A consequence of this is that SSSSSS only needs to upgrade to an exclusive lock if you have read the data you are writing. Otherwise SSSSSS uses a shared lock called a writer shared lock. But at distributed transaction participants even blind writes will take exclusive locks because transactions at different coordinators cannot be guaranteed, in general, to be assigned different commit timestamps.

SSSSSS uses the standard "wound-wait" algorithm to handle deadlock detection: under the hood SSSSSS keeps track of the age of each transaction that requests conflicting locks, and allows older transactions to abort younger transactions (where "older" means the transaction started earlier). By giving priority to older transactions, SSSSSS ensures that eventually every transaction has a chance to acquire locks by virtue of it becoming old enough to have higher priority than other transactions. e.g. a transaction holding a reader shared lock can be aborted by an older transaction wanting a writer shared lock.

[C++20] Nicolai M. Josuttis's C++20 the complete guide reading minute - operator <=>

Comparisons and Operator <=>

Reference:

std::compare_three_way

std::three_way_comparable

std::lexicographical_compare_three_way

C++20 behaviour breaking existing code with equality operator?

C++20 compiler rewrites

1. Operator != with !(a==b)
2. If above doesn't work, change the order of the operands. !(b==a)

And a, b can be different types.

If there's a free-standing operator:

• A free-standing operator!=(TypeA, TypeB)
• A free-standing operator==(TypeA, TypeB)
• A free-standing operator==(TypeB, TypeA)
• A member function TypeA::operator!=(TypeB)
• A member function TypeA::operator==(TypeB)
• A member function TypeB::operator==(TypeA)
  

Compiler can do the rewrite trick.

Thus To compile

x != y

the compiler might now try all of the following:

x.operator!=(y) // calling member operator!= for x
operator!=(x, y) // calling a free-standing operator!= for x and y
!x.operator==(y) // calling member operator== for x
!operator==(x, y) // calling a free-standing operator== for x and y
!x.operator==(y) // calling member operator== generated by operator<=> for x
!y.operator==(x) // calling member operator== generated by operator<=> for y

The last form is tried to support an implicit type conversion for the first operand, which requires that the operand is a parameter.

In general, the compiler tries to call:

• A free-standing operator !=: operator!=(x, y)
• or a member operator !=: x.operator!=(y)

Having both operators != defined is an ambiguity error.

• A free-standing operator ==: !operator==(x, y)
• or a member operator ==: !x.operator==(y)

Note that the member operator == may be generated from a defaulted operator<=> member.

Having both operators == defined is an ambiguity error.
This also applies if the member operator== is generated due to a defaulted operator<=>.

When an implicit type conversion for the first operand v is necessary, 

the compiler also tries to reorder the operands. Consider:

42 != y // 42 implicitly converts to the type of y

In that case, the compiler tries to call in that order:

• A free-standing or member operator !=
• A free-standing or member operator == (note that the member operator == may be generated from a defaulted operator<=> member)

Note that a rewritten expression never tries to call a member operator !=

Thus To compile

x <= y

new operator <=> and compare the result with 0. 

The operator behaves like a three-way comparison function

returning a negative value for less, 0 for equal, and a positive value for greater 

(the returned value is not a numeric value; it is only a value that supports the corresponding comparisons).

The compiler might now try all of the following:

x.operator<=(y) // calling member operator<= for x
operator<=(x, y) // calling a free-standing operator<= for x and y
x.operator<=>(y) <= 0 // calling member operator<=> for x
operator<=>(x, y) <= 0 // calling a free-standing operator<=> for x and y
0 <= y.operator<=>(x) // calling member operator<=> for y

The last form is tried to support an implicit type conversion for the first operand,

for which it has to become a parameter.

Operator <=>

• The return of <=> operator type should be marked as 'auto' and let compiler to deduce the type.
• Operator <=> takes precedence over all other comparison operators; except explicitly user defined.
• Should only call operator <=> directly when implementing operator<=>.
  However, it can be very helpful to know the returned comparison category.

#include <compare>
// order of the members in the class matters.
class Value {
	// defines the ordering and can be used by the relational operators <, <=, >, and >=.
	auto operator<=> (const Value& rhs) const = default;
	// implicitly generated
	// defines equality and can be used by the equality operators == and !=.
	auto operator== (const Value& rhs) const = default; 
};

class Value {
  private:
  	long id;

  public:
	constexpr Value(long i) noexcept
		: id{i} {}

	// for equality operators:
	bool operator== (const Value& rhs) const {
	  return id == rhs.id; // defines equality (== and !=)
	}

	// for relational operators:
	auto operator<=> (const Value& rhs) const {
	  return id <=> rhs.id; // defines ordering (<, <=, >, and >=)
	}
};

Compiler generated operator has following traits

• They are noexcept if comparing the members never throws 
• They are constexpr if comparing the members is possible at compile time 
• Thanks to rewriting, implicit type conversions for the first operand are also supported (This can also be tricky/buggy)

C++20 compiler rewrites

If no operator<=:

x <= y

rewrites with:

  (x <=> y) <= 0;
  // Or:
  0 <= (y <=> x);

• If the value of x<=>y is equal to 0, x and y are equal or equivalent.
• If the value of x<=>y is less than 0, x is less than y.
• If the value of x<=>y is greater than 0, x is greater than y.

However, note that the return type of operator<=> is not an integral value.

Return type is a type that signals the comparison category, which could be 

• strong ordering, 
• weak ordering, 
• or partial ordering.

These types support the comparison with 0 to deal with the result.

Note that operator<=> is for implementing types. Outside the implementation of an operator<=>,

programmers should never invoke <=> directly. Although you can, you should never write 

a<=>b < 0

instead of

a<b

Comparison Category Types

strong ordering (total ordering):

– std::strong_ordering::less

– std::strong_ordering::equal

(also available as std::strong_ordering::equivalent)

– std::strong_ordering::greater

Any value of a given type is less than or equal to or

greater than any other value of this type (including itself).

weak ordering:

– std::weak_ordering::less

– std::weak_ordering::equivalent

– std::weak_ordering::greater

Any value of a given type is less than or equivalent to or greater than any other

value of this type (including itself). However, equivalent values do not have to be equal

 (have the same value).

E.g. "hello" is equivalent to "HELLO"

partial ordering:

– std::partial_ordering::less

– std::partial_ordering::equivalent

– std::partial_ordering::greater

– std::partial_ordering::unordered

Any value of a given type could be less than or equivalent to or greater than any

other value of this type (including itself). 

However, in addition, it may not be possible to specify a specific order between two values at all.

E.g. floating-point types, because they might have the special value

NaN (“not a number”). Any comparison with NaN yields false. Therefore, in this case a comparison

might yield that two values are unordered and the comparison operator might return one of four values.

std::strong_ordering operator<=> (MyType x, MyOtherType y)
{
  if (xIsEqualToY) return std::strong_ordering::equal;
  if (xIsLessThanY) return std::strong_ordering::less;
  return std::strong_ordering::greater;
}

class MyType {
  std::strong_ordering operator<=> (const MyType& rhs) const {
    return value == rhs.value ? std::strong_ordering::equal :
      value < rhs.value ? std::strong_ordering::less :
      std::strong_ordering::greater;
  }
  type value;
};

// often
class MyType {
  auto operator<=> (const MyType& rhs) const {
    return value <=> rhs.value;
  }
  type value;
};

C++20 compiler rewrites

if (!(x < y || y < x)) // might call operator<=> to check for equality
if (x <= y && y <= x) // might call operator<=> to check for equality 

Operator <=> return type mismatch due to multiple data members:

class Person {
std::string name;
double value;

  std::partial_ordering operator<=> (const Person& rhs) const { // OK
    auto cmp1 = name <=> rhs.name;
    if (cmp1 != 0) return cmp1; // strong_ordering converted to return type
    return value <=> rhs.value; // partial_ordering used as the return type
  }
};

// better
class Person {
std::string name;
double value;

  auto operator<=> (const Person& rhs) const 
	-> std::common_comparison_category_t<decltype(name <=> rhs.name),
	decltype(value <=> rhs.value)> {
    auto cmp1 = name <=> rhs.name;
    if (cmp1 != 0) return cmp1; // used as or converted to common comparison type
    return value <=> rhs.value; // used as or converted to common comparison type
  }
};

// convert to same comparison category:
class Person {
std::string name;
double value;

  std::strong_ordering operator<=> (const Person& rhs) const {
    auto cmp1 = name <=> rhs.name;
    if (cmp1 != 0) return cmp1; // return strong_ordering for std::string
    // map floating-point comparison result to strong ordering:
    // https://en.cppreference.com/w/cpp/utility/compare/strong_order
    return std::strong_order(value, rhs.value);
  }
};

std::strong_order() yields a std::strong_ordering value according to the passed arguments as follows:

• Using std::strong_order(val1, val2) for the passed types if defined
• Otherwise, if the passed values are floating-point types, using the value of totalOrder() as specified in ISO/IEC/IEEE 60559 (for which, e.g., -0 is less than +0 and -NaN is less than any non-NAN value and +NaN) 
• Using the new function object std::compare_three_way{}(val1, val2) if defined for the passed types std::compare_three_way use like std::less 

For other types that have a weaker ordering and operators == and < defined, you can use the function

object std::compare_strong_order_fallback(T&& t, U&& u)

```

Performs three-way comparison on subexpressions t and u and produces a result of type std::strong_ordering, even if the operator <=> is unavailable.

```

accordingly:

class Person {
std::string name;
SomeType value;

  std::strong_ordering operator<=> (const Person& rhs) const {
    auto cmp1 = name <=> rhs.name;
    if (cmp1 != 0) return cmp1; // return strong_ordering for std::string
    // map weak/partial comparison result to strong ordering:
    return std::compare_strong_order_fallback(value, rhs.value);
  }
};

Defaulted operator== and operator<=> contract:

Defaulted operator<=> implies Defaulted operator==
Thus the following is enough to support all six comparison operators for objects of the type Coord:

#include <compare>
struct Coord {
  double x{};
  double y{};
  double z{};
  auto operator<=>(const Coord&) const = default;
};

The second parameter as const lvalue reference (const &) in member function. 

Friend functions might alternatively take both parameters by value.

• The defaulted operators require the support of the members and possible base classes
• Defaulted operators == require the support of == in the members and base classes.
• Defaulted operators <=> require the support of == and either an implemented operator < or a defaulted operator <=> in the members and base classes.
• The operator is noexcept if comparing the members guarantees not to throw.
• The operator is constexpr if comparing the members is possible at compile time.

For empty classes, the defaulted operators compare all objects as equal: 

• operators ==, <=, and >= yield true, 
• operators !=, <, and > yield false, 
• and <=> yields std::strong_ordering::equal.

template<typename T>
class Type {
  public:
	[[nodiscard]] virtual std::strong_ordering
		operator<=>(const Type&) const requires(!std::same_as<T,bool>) = default;
};

// compiler generates equivalent to
template<typename T>
class Type {
  public:
	[[nodiscard]] virtual std::strong_ordering
		operator<=> (const Type&) const requires(!std::same_as<T,bool>) = default;
    [[nodiscard]] virtual bool
		operator== (const Type&) const requires(!std::same_as<T,bool>) = default;
};

Implementation of the Defaulted operator<=>:

Contract:

If operator<=> is defaulted and you have members or base classes and you call one of the relational

operators, then the following happens:

• If operator<=> is defined for a member or base class, that operator is called.
• Otherwise, operator== and operator< are called to decide whether (from the point of view of the members or base classes)

– The objects are equal/equivalent (operator== yields true)

– The objects are less or greater

– The objects are unordered (only when partial ordering is checked)

In this case, the return type of the defaulted operator<=> calling these operators cannot be auto.

For example, consider the following declarations:

struct B {
	bool operator==(const B&) const;
	bool operator<(const B&) const;
};

struct D : public B {
  // return type can not be auto due to base type has the operator== and operator< defined
  // because it cannot decide which ordering category the base class has. 
  // In that case, you need operator<=> in the base class too.
  std::strong_ordering operator<=> (const D&) const = default;

  // auto generated by compiler even
  // operator<=> is declared as
  // auto operator<=> (const D&) const = default;
  // which then d1 > d2; does't work but d1 != d2; works.
  bool operator== (const D&) const = default;  
};

// Then:
D d1, d2;
d1 > d2; // calls B::operator== and possibly B::operator<

// If operator== yields true, we know that the result of > is false and that is it. 
// Otherwise, operator< is called to find out whether the expression is true or false.

Compare values of a generic type

Defines a total order for raw pointers.

For forward declare operator<=>() result type; use std::compare_three_way_result_t

template<typename T>
struct Value {
  T val{};
...
  auto operator<=> (const Value& v) const noexcept(noexcept(val<=>val)) {
     return std::compare_three_way{}(val<=>v.val);
  }
};

template<typename T>
struct Value {
  T val{};
...
  std::compare_three_way_result_t<T,T>
    operator<=> (const Value& v) const noexcept(noexcept(val<=>val));
};

Appendix

std::common_comparison_category impl:

namespace detail
{
    template <unsigned int>
    struct common_cmpcat_base      { using type = void; };
    template <>
    struct common_cmpcat_base <0u> { using type = std::strong_ordering; };
    template <>
    struct common_cmpcat_base <2u> { using type = std::partial_ordering; };
    template <>
    struct common_cmpcat_base <4u> { using type = std::weak_ordering; };
    template <>
    struct common_cmpcat_base <6u> { using type = std::partial_ordering; };
} // namespace detail
 
template <class...Ts>
struct common_comparison_category :
    detail::common_cmpcat_base <(0u | ... |
        (std::is_same_v <Ts, std::strong_ordering>  ? 0u :
         std::is_same_v <Ts, std::weak_ordering>    ? 4u :
         std::is_same_v <Ts, std::partial_ordering> ? 2u : 1u)
    )> {};