cuFFTDx 2d FFT explanation of logical remapping for shared memory based global memory to register memory IO

[Coll1ns-cult]

Hi,

I am trying to understand the thread remapping logic in the shared memory variants of load_strided / store_strided in block_io_generic_strided.hpp of mathdx\25.12\example\cufftdx\05_fft_Xd folder.

My current understanding of the natural thread assignment (used in the non-shared-memory variants and during FFT computation for 1D FFT in introduction_example.cu) is:

• threadIdx.x → thread's position within its FFT
• threadIdx.y → local_fft_id
• blockIdx.x → block's batch offset

However, in the shared memory variants, the code performs the following remapping before doing the global memory I/O:

const unsigned int tid          = threadIdx.x + FFT::working_group::block_dim().x * threadIdx.y;
const unsigned int tidx         = tid / FFT::working_group::block_dim().y;
const unsigned int tidy         = tid % FFT::working_group::block_dim().y;
unsigned int       smem_index   = tidx + tidy * FFT::working_group::block_dim().x;

I tried to understand it through with AI, however, I yet to succeed with a clear explanation of this logical mapping.
What I still don't fully understand is:

1. What is the concrete motivation for performing this logical mapping at all? In other words, what problem does it solve compared to just using threadIdx.x / threadIdx.y directly for global memory indexing as in the non-smem variants?
2. After the transpose, the smem index is computed as tidx + tidy * block_dim().x. Then after __syncthreads(), the read phase uses threadIdx.x + threadIdx.y * block_dim().x with the same formula. Could you walk through concretely what data each thread writes vs. reads, and why those two different identities (tidx/tidy vs threadIdx.x/threadIdx.y) indexing the same smem layout performs the intended transpose?
3. Is my understanding of blockIdx.y correct — is it unused / always 0 in this kernel launch configuration, or does it carry meaning?

Thanks in advance for your time!

[drzejan2]

Hi @Coll1ns-cult

Here are answers to your questions, to best of my knowledge:

1. What is the concrete motivation for performing this logical mapping at all? In other words, what problem does it solve compared to just using threadIdx.x / threadIdx.y directly for global memory indexing as in the non-smem variants?

The motivation for applying this mapping for shared memory is improving performance of data loading by enforcing coalesced memory access to global data. By using shared memory buffer, as an intermediate step, we enforce a performant data loading pattern from global memory.

In the next step, as we have the data in shared memory buffer, we pick pieces of data that are required by each of threads to process with FFT signals with different data access pattern. As all the data is in shared memory buffer the cost of accessing it, possibly converting to complex type, and then storing in buffer local to the thread should be lower when compared to the approach where shared memory is not used. This becomes more important when we deal with FFT that has more than 1 dimension and pieces of data needed by single thread are scattered across large buffer.

2. After the transpose, the smem index is computed as tidx + tidy * block_dim().x. Then after __syncthreads(), the read phase uses threadIdx.x + threadIdx.y * block_dim().x with the same formula. Could you walk through concretely what data each thread writes vs. reads, and why those two different identities (tidx/tidy vs threadIdx.x/threadIdx.y) indexing the same smem layout performs the intended transpose?

Update: the original reply was incorrect. Sorry for that. I will try to get back to your question later with the details.

3. Is my understanding of blockIdx.y correct — is it unused / always 0 in this kernel launch configuration, or does it carry meaning?

Update: In the example example/cufftdx/05_fft_Xd/fft_3d.cu, the grid defined in a way that y/z dimensions are set to 1. It is not used in the helpers.

Thank you for reaching out for details when something parts of prepared examples were not clear.