Repository navigation
cuFFTDx 2d FFT explanation of logical remapping for shared memory based global memory to register memory IO #307
Description
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 FFTthreadIdx.y→local_fft_idblockIdx.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:
- 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.ydirectly for global memory indexing as in the non-smem variants? - After the transpose, the smem index is computed as
tidx + tidy * block_dim().x. Then after__syncthreads(), the read phase usesthreadIdx.x + threadIdx.y * block_dim().xwith 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? - Is my understanding of
blockIdx.ycorrect — is it unused / always 0 in this kernel launch configuration, or does it carry meaning?
Thanks in advance for your time!
Activity
drzejan2 commentedon Mar 18, 2026
Hi @Coll1ns-cult
Here are answers to your questions, to best of my knowledge:
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.
Update: the original reply was incorrect. Sorry for that. I will try to get back to your question later with the details.
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.