Functional/Non-Functional C++ Tile Evaluation

[fbusato]

Functional/Non-Functional C++ Tile Evaluation

Tile-based GPU programming is a promising direction for both providing performance and maintainable code. As a new technology, it is fundamental to explore and evaluate capabilities and limitations in the context of CCCL algorithms.

In particular, this issue lists several functional and non-functional properties that should be covered to get a robust understanding of the new programming approach. This step is critical to understand future directions and development effort.

We will focus on the following routine for the exploration:

• Implement cub::DeviceTransform as tile kernel #8087

Functional

• Correctness

  • What (current) cuTile limitations prevent algorithm implementation?
  • Numerical behavior compared to CUDA C++ implementation.
  • Could GPU arch-dependencies or constrains impact correctness?

• API surface

  • What new APIs are introduced? and they could be coupled together with existing one? (are they stable, experimental, internal?)
  • Which input sizes, ranks, layouts, data types, execution spaces, or iterator/category are supported?

• Interoperability

  • Interaction with existing CCCL components: CUB, Thrust, libcudacxx, CMake targets, tests, benchmarks.

Non-functional properties

Implementation Complexity

• Is Tile documented enough?
• Does the Tile implementation reduce or increase the amount work?
• Are diagnostics good enough to understand the problem?
• How well it support direct or automatic tuning?

Maintainability

• Are implementation details isolated from public APIs?
• What could be enforced in Debug mode?

Performance

• Is Tile mature enough to compete with C++ implementation in terms of perfomance metrics? (do they have the same distribution properties, see NVBench)
• Launch overhead? Impact of tuning?
• Compilation intermediate steps inspection, e.g. IR

Ecosystem

• How current tools Nsight Compute/Systems, nvdisasm, cuobjdump interact with Tile?

Compile time

• How Tile compared to C++ implementation?
• Linking time?

Binary size

• Any changes in final or intermediate artifacts, e.g object files, fatbin

[nanan-nvidia]

A tile kernel was implemented and did some test runs on RTX A6000 ADA.

Of the 22 CUB::DeviceTransform benches ran on RTX A6000 ADA, 20 of them is on the same performance as CUB. The ones that showed sigificant performance degrade are two cases: fibonacci and heavy.

1. fibonacci

Measurements

size      tile GElem/s   cub GElem/s   ratio
2^16        152            144         1.06x
2^20        441            553         0.80x
2^24        505            673         0.75x
2^28        501            684         0.73x

Code on each side

CUB (bench/fib_cub.cu):

[] __device__ (int64_t n) -> int64_t {
    int64_t t1 = 0, t2 = 1;
    if (n < 1) return 0;
    if (n == 1) return 0;
    if (n == 2) return 1;
    for (int64_t k = 3; k <= n; ++k) {
        auto next = t1 + t2;  t1 = t2;  t2 = next;
    }
    return t2;
}

Tile (bench/fib_tile.cu):

struct fib_op {
    template <class N> __tile__ auto operator()(N n) const {
        auto zero = n - n, one = zero + 1;
        auto t1 = zero, t2 = one;
        for (int i = 3; i <= MAX_FIB_N; ++i) {
            auto next = t1 + t2;
            auto cond = n >= i;
            t1 = ct::select(cond, t2, t1);
            t2 = ct::select(cond, next, t2);
        }
        auto out = ct::select(n < 1,  zero, t2);
        out     = ct::select(n == 1, zero, out);
        out     = ct::select(n == 2, one,  out);
        return out;
    }
};

The fundemental reason for performance difference is because how the tile semantics works : we have to operate on the whole tile. Therefore, we have to use fully unroll the kernel with ct::select which does a lot more wasted compute.

2. heavy

The tile kernel fails to launch at any size beyond 6K elements because Device-side malloc failed.

The tile function is as follows:

template <int N>
struct heavy_op {
    template <class T> __tile__ auto operator()(T data) const {
        T reg[N];
        reg[0] = data;
        for (int i = 1; i < N; ++i) reg[i] = reg[i-1] * reg[i-1] + 1u;
        for (int i = 0; i < N; ++i) reg[i] = (reg[i] * reg[i]) % 19u;
        for (int i = 0; i < N; ++i) reg[i] = reg[N-i-1] * reg[i];
        T x = data - data;
        for (int i = 0; i < N; ++i) x = x + reg[i];
        return x;
    }
};

At compile time, lets say the input is int and with N=8 and TileSize=256, T is tile<int, shape<256>>. Here, with normal cuda semantics this should be 8 registers per thread; however, the nvcc frontend lowers this to

%7 = alloca num_elem = 8192, alignment = 4 : tile<ptr<i8>>

For N=8 with TileSize=256, that is 8 x 256 x 4 = 8192 bytes. This gets lowered to a malloc call on GPU:

...                                         ; execute only on thread0
/*0090*/  MOV R0, 0x8
/*00a0*/  MOV R4, 0x2000                    ; malloc size = 8192
/*00b0*/  LDC.64 R2, c[0x4][R0]             ; R2 should now point to __cuda_syscall_aligned_malloc
/*0160*/  CALL.ABS.NOINC R2                 ; CALL
...

Therefore, reg[N] is entirely on global memory and allocated on heap (I did not even know you can call malloc on GPUs, LOL), which causes sigificant performance degrade and failures when there are too many elements to malloc.

Not sure if this is the expected behavior of the front end, since for those cases we'd expect them to be on registers.