Legacy translator: Incorrect stride for map indexing in CUDA kernels

[mattbuergler]

I'm observing a mismatch between host-side and device-side indirect map access in generated CUDA kernels. The root cause is an incorrect stride used in the generated CUDA kernel when indexing the transposed map_data_d array.

Summary

• The function op_decl_map() in op2/src/cuda/op_cuda_decl.cpp transposes host maps to the device layout and each column is padded to a multiple of 32 using round32(set_size).
• However, the generated CUDA kernels continue to index map_data_d using the unpadded stride set_size = set->size + set->exec_size, e.g.:

map1idx = opDat1Map[n + set_size * 0];
map2idx = opDat1Map[n + set_size * 1];
map3idx = opDat1Map[n + set_size * 2];

• This produces misaligned column access and the kernel reads into the padding region (zeros / uninitialized values) instead of the next column.
• By debugging, I confirmed that:
  • Column 0 data is correct (e.g., all map1idx in the above example).
  • At indices [set_size .. round32(set_size)-1] the map contains padding zeros.
  • Column 1 starts at round32(set_size), not at set_size, resulting in wrong map2idx and map3idx .

[mattbuergler]

I proposed solution in #256.

[bozbez]

@mattbuergler Could you give #257 a spin to check it's also working for you? Apologies for this, it's come from a bad merge when we re-added the legacy translator to the new translator branch. There should be a few percent performance improvement from the padded buffers.

[mattbuergler]

@bozbez thanks for preparing this fix. I tested the patch, but ran into the following compilation error:

error: calling a constexpr __host__ function("round32") 
from a __global__ function("op_cuda_kernel_TOPOLOGY_NodeStringHydStates") 
is not allowed. 
The experimental flag '--expt-relaxed-constexpr' can be used to allow this.

map3idx = opDat2Map[n + round32(set_size) * 1];

I am not sure if it makes sense to call round32() on the device, since this also slows down the kernel. I think it would be better to:

• compute the padded stride (round32(set_size)) on the host,
• either pass id as a new arguement to the __global__ CUDA kernel function or store it in a per-kernel __device__ __constant__ symbol (e.g. opDatX_kernelName_stride_OP2CONSTANT),
• and use that constant inside the CUDA kernel instead of calling round32().

I appreciate your effort to get the legacy translator working again. Let me know if you want me to contribute by proposing a draft.

[reguly]

@mattbuergler hopefully fixed now - there was some complexity around which data structures are in a strided format, and which ones are not. Let us know if this works!

[mattbuergler]

@reguly Commit d6c3a2b works for us. Thanks a lot for fixing this!