2020-10-07

However, here is one problem that's causing us to use extra memory allocations: Double buffering. Because double buffering inserts happens at the end of the kernel, we can have cases where something that needs to be written to the double-buffer is kept alive unnecessarily long. It looks like this:

let x_mem = alloc(...)
let x_res = ...
-- something else
let x_res_double_buffer = copy(x_res)

If we could somehow move the copy to x_res_double_buffer up, x_mem would be free to be used for something else.

Okay, let's try to make sense of what's going on here.

__kernel void payoffGPU(
                        __constant LoopROScalars* ro_scal,      // RO SCALARS
                        __constant int*           sobol_v_dir,  // RO SOBOL
                        __constant UCHAR*         fix_index,
                        // RO Brownian Bridge
                        __constant int*           bb_ia,
                        __constant REAL*          bb_coefs,
                        // RO MODELS DATA
                        __constant REAL*          model_coefs,
                        // WO (GLOBAL) VHAT
                        __global   REAL*          model_vhat,
                        // LOCAL ARRAYS
                        __local    REAL*          md_z,
                        //__local    REAL*          inst_trajWF,
                        __local    REAL*          vhat_local
                        ) {

Here's the definition of the main payoff function. Most of the arguments are read-only constants. There is one global array, model_vhat, and two local (shared) arrays, md_z and vhat_local.

By inspecting the kernel, we can see that model_vhat is only ever written to at the very end, so it must contain the result of running the kernel.

It looks like vhat_local is primarily used when running payoffFunction. I reason that it's used as a result buffer. Afterwards, a segmented scan is performed on vhat_local and it is written to model_vhat.

The md_z shared array is used as a result buffer for the call to mlfi_brownianbridge_wiener_pathNoTransGPU, which computes the brownian bridge. The result is then used to compute trajWF.

We conclude that these three allocations have rather limited use, and no re-use takes place at all.

Now, there are further arrays in use in the payoffGPU kernel:

int  sobol_last_num_vec[15];//UINT[ro_scal->sobol_dim];
REAL md_zd              [15]; //[ro_scal->sobol_dim];

Each of these are private arrays to each execution unit, although the compiler will probably expand those automatically, so they correspond to shared memory. Let's see what they are used for.

Fairly quickly, it can be seen that sobol_last_num_vec is only ever used as an intermediate array in the calls to mlfi_genmatrix_uniformGPUind and friends. The results are never used at any later point, and neither is memory.

md_zd, or at least the memory it's allocated in, is used extensively throughout the kernel, both in those calls to mlfi_genmatrix_uniform_GPUind and friends, which explains why we must have both md_zd and sobol_last_num_vec¹. It is first used in mlfi_genmatrix_uniform_GPUind and friends as a result buffer, then in mlfi_ugaussian_Pinv_vector as both an input and result buffer and finally in mlfi_brownianbridge_wiener_pathNoTransGPU as an input buffer. Then something interesting happens:

REAL* trajWF    = md_zd;

The memory md_zd is in, is reused for the trajWF array! trajWF is filled with numbers computed using md_z and the result is used as an input array to payoffFunction.

In conclusion, there are 5 interesting arrays in the payoffGPU kernel from GenericPricingPrivOpt.cl:

• model_vhat is the result buffer for the whole kernel. Resides in global memory.
• vhat_local is the result buffer for the payoffFunction. Resides in shared memory.
• md_z is the result buffer for mlfi_brownianbridge_wiener_pathNoTransGPU. Resides in shared memory.
• sobol_last_num_vec is an intermediate array used inside mlfi_genmatrix_uniformGPUind. It is private to each execution unit, but expanded to reside in shared memory.

• md_zd (and the co-located trajWF) is used for various things throughout the kernel as a result and intermediate buffer.

  Now, the interesting question is: why are there so many buffers used in the Futhark implementation, if this is all we need?

Let's take a look at the code generated from OptionPricing.fut using the command. Note that we're not reusing allocations here, just to keep it simple.

futhark-reuse-allocations dev --kernels -a -e --cse -e --double-buffer -e --cse -e OptionPricing.fut | bat -l fut

To start with, I've removed everything after the creation of sobol_mat:

1: ...
2: let sobol_mat = map_stream (\chunk (ns: [chunk]i32): [chunk][sobvctsz]f32  ->
3:                               sobolChunk dir_vs (#[unsafe] ns[0]) chunk)
4:                            (iota num_mc_it)
5: in sobol_mat |> flatten

Inspecting the resulting code, we see there are a number of allocations inside the kernels:

1: ...
2: let {mem mem_7497} = alloc(size_7496)
3:                           ...
4: let {mem@local mem_7468} = alloc(bytes_7466, @local)
5:                                 ...
6: let {mem@local mem_7472} = alloc(bytes_7466, @local)
7:                                 ...

The two last memory bloks are used as shared memory within a single kernel. Perhaps they correspond to the md_zd and sobol_last_num_vec from the OpenCL kernel?

The Futhark version of the Sobol generators seem much more complicated than the OpenCL version. What's going on there? Hm, maybe it's not so bad…

Okay, what happens if we introduce the gaussian transformation?

1: let sobol_mat = map_stream (\chunk (ns: [chunk]i32): [chunk][sobvctsz]f32  ->
2:                               sobolChunk dir_vs (#[unsafe] ns[0]) chunk)
3:                            (iota num_mc_it)
4: let gauss_mat = map ugaussian sobol_mat
5: in gauss_mat |> flatten

Seems like there are no additional allocations happening. That's good. Let's try brownian. Okay, that seems to cause a new allocation to happen, inside a segmap_thread that handles the brownian bridge. That sounds about right, the mlfi_brownianbridge_wiener_pathNoTransGPU function from the OpenCL implementation also uses an additional allocation: md_z.

Next is payoffs:

 1: let sobol_mat = map_stream (\chunk (ns: [chunk]i32): [chunk][sobvctsz]f32  ->
 2:                               sobolChunk dir_vs (#[unsafe] ns[0]) chunk)
 3:                            (iota num_mc_it)
 4: let gauss_mat = map ugaussian sobol_mat
 5: let bb_mat    = map (brownianBridge num_und bb_inds bb_data) gauss_mat
 6: let payoffs   = #[incremental_flattening(only_intra)]
 7:                 map (\bb_row: [num_models]f32  ->
 8:                        let bd_row = map4 (blackScholes bb_row) md_cs md_vols md_drifts md_sts
 9:                        in map3 (genericPayoff contract_number) md_discts md_detvals bd_row)
10:                 bb_mat
11: in payoffs |> flatten

I added the incremental flattening attribute to reduce the total amount of code that I have to look at. I believe this is the interesting kernel anyway.

A lot happens to the generated code after adding those few lines. I count 6 new allocations inside kernels (most of them @local). Ah, perhaps some of those are caused by incremental flattening. There are multiple different versions.