Kernel Generation

[ScriptLineStudios]

The end goal here is a vector of Kernel objects which will be derived from the tree. Each kernel will derive from a base BruteForceKernel and store all needed data for cuda emission, critically, this emission must be setup in such a way such that multiple kernels can have their logic stitched together, the emission made by the BruteForceKernel will look something along the lines of:

bool bruteforce_pool_0(uint64_t lootseed) {
    // in reality this will probably be a lootseed in some other state, but lets just say for simplicity's sake that it is the loot seed :)
    uint64_t rng = lootseed;
    set_seed(&rng, rng);
    int rolls = next_int_bounded(&rng, 4, 8);
    for (int r = 0; r < rolls; r++) {
        ...
     }
}

The emission made by the SingleItemReversalKernel must will look something along the lines of:

bool filter_forward(uint64_t state) {
    // blah blah blah
}

__global__ void kernel(uint64_t offset) {
    uint64_t tid = (uint64_t)blockDim.x * blockIdx.x + (uint64_t)threadIdx.x + offset;
    uint32_t lower17 = tid & 0x1ffff;
    for (int r = MIN_REMAINDER; r <= MAX_REMAINDER; r++) {
      uint64_t upper31 = (tid >> 17) * NEXT_INT_BOUND + r;
      uint64_t state = (upper31 << 17) | lower17;
  
      // The code below illustrates the general process we're going for:
      // - filter the pool we started in forward 
      // - go back to the very beginning of the loot generation process and check the full loot table.

      if (!filter_forward(state)) {
          return;
      }

      for (int i = MIN_GO_BACK; i < MAX_GO_BACK; i++) {
           go_back(&state);
           uint64_t loot_seed = state ^ 0x5DEECE66D;
           if (filter_forward_pools_0_2(loot_seed)) {
               add_result(loot_seed);
           }
      }
    }

This kernel has 2 responsibilities: first of all it must fill out the above template, it also needs to be self-aware of the functions that have been made accessible by the parent Kernel and fill in these functions calls in the appropriate places. As discussed yesterday, all kernels will follow this rather predictable pattern:

Meaning that this can probably be achieved without much abstraction. Things we still need to discuss include: How do we want to store kernel templates? Perhaps we introduce some sort of extension like .cu_temp and come up with a bunch of special keywords which the template generator will be responsible for replacing (this is probably the closest we're gonna get to an lsm revival)