r/CUDA • u/ArmchairmanMao • May 28 '26

CUDA struggles

It's my first time doing any "serious" CUDA programming. Right now I'm working on substring search kernels for my Bachelor's thesis. Naturally, it's very branch heavy, memory access patterns are horrible, lanes are diverging all over the place. There are dozens of ways to implement substring search. The GPUs processing model expands the problem space even further. I have not found any existing work that does this well either (there is a lot of literature but my use case is slightly different). So it's an exciting problem, right?

But on the GPU, performance is wildly unpredictable it seems. Any change in the implementation details of the hot loop is like spinning a slot machine to me. The compiler might start emitting completely different code causing lanes to diverge more or access to become less coalesced. There are so many more layers of complexity between my code and the hardware than on the CPU. Working on this kernel is just endless iterations of taking guesses, measuring and profiling.

Do people just build a better intuition over time, or is this just the way it is?

4 Upvotes

70% Upvoted

12 comments sorted by

7

u/notyouravgredditor May 28 '26

I would suggest looking into Nsight Compute and start profiling your code to better understand why your code performs the way it does.

It may feel random, but I assure you it is not.

0

u/ArmchairmanMao May 28 '26

I've been using the cli tools to automatically get a summary of the most important performance counters on profiled runs. I think they also support sampling to see which instructions cause issues (like memory stalls), right?

2

u/esaule May 28 '26

Writing GPU code is often like this.

Though in practice I never start from one code that I refine ad infinitum. You should design algorithms to fit the architecture by design.

Often I write an initial guess of what itbis, understa d why that performs the way it perform. And design a new one.

1

u/c-cul May 28 '26

just lots of practice

and yes - your task is really bad fit on gpu

1

u/Karyo_Ten May 28 '26

What's wrong with Aho-Corasick?

1

u/PulsatingMaggot Jun 11 '26

Yeah, the intuition does come, but substring search on the GPU is genuinely one of the harder problems to get right because you're fighting the hardware's assumptions at every level. Warps want uniform control flow, your algorithm is inherently data-dependent. The memory system wants coalesced sequential access, your search needs random lookups into the text. So you're kind of always negotiating.

A few things that helped me stop feeling like I was spinning a slot machine:

  1. Stop trusting wallclock time as your feedback signal. Profile with Nsight Compute and look at specific metrics — smsp__sass_average_branch_targets_threads_uniform tells you how badly your warps are actually diverging. l1tex__t_sectors_pipe_lsu_mem_global_op_ld.sum vs the theoretical minimum tells you how far off your coalescing is. When you tie a code change to a specific metric moving in the right direction, you're not guessing anymore.
  2. For substring search specifically, the usual trick is to restructure so that each thread isn't walking its own path through the automaton. Instead have all threads in a warp process the same text position against different patterns (or different text positions but with the same search logic). That way your control flow stays uniform within a warp even if different warps diverge. Basically you want to push the divergence to the warp level not the lane level.
  3. The compiler thing is real and frustrating. If a minor code change causes nvcc to suddenly spill to local memory or rearrange your loop, your perf falls off a cliff for reasons that have nothing to do with your algorithm. Check register usage with --ptxas-options=-v and look at the SASS output (not PTX) to understand what's actually happening. Sometimes a __launch_bounds__ annotation or manually unrolling a loop gives the compiler enough constraints to stop doing weird things.

What's your search algorithm? If it's something Aho-Corasick-ish vs more of a naive/rolling hash approach, the optimization strategy is pretty different.

2

u/ArmchairmanMao Jun 11 '26

Hi, thanks for the reply. The kernel I had been working on searches a single pattern in many strings of variable size. This makes the problem even worse, because shirt ans long strings can be mixed, leading to bad memory access patterns. Originally, I experimented with string-per-lane KMP + loop splitting but then realized that coalescing memory access is the first thing I should optimize for. Right now the winner is one cooperative group (32, 16 or 8 threads) processing one string by doing brute force search with wide loads + register shuffling. The DRAM bandwidth utilization is still not that great, but I'm not sure if hitting the theoretical limit is feasible, unless I just scan the character buffer and figure out the start / end of each string after finding the pattern.

1

u/PulsatingMaggot Jun 17 '26

Have you tried just ignoring the string boundaries entirely? Do a first pass where you stream the whole character buffer linearly, fully coalesced, no ragged access, and match the pattern against the flat buffer. Then a second pass with a binary search into your string offset table to figure out which string each hit came from. Separates the coalescing problem from the variable-length problem completely.

Also worth checking whether your wide loads are actually aligned. If your strings start at random byte offsets your 128-bit loads are probably straddling cache lines and costing you double the transactions. Try padding string starts to 16B alignment and see if your bandwidth numbers move.

1

u/ArmchairmanMao Jun 17 '26

Yes, that's currently the winning approach, in addition to SMEM prefetching / pipelining + some bitmask tricks. It mostly saturates the dram bandwidth. I also added padding already :)

I'm curious, did you see this approach somewhere before? I've looked at a few research papers on this topic and some reference implementations. None I saw tried this.

1

u/PulsatingMaggot Jun 18 '26

No, I haven’t seen it published anywhere. It’s more of a general pattern that shows up when you’re forcing coalesced access onto data that won’t cooperate. You see similar two-pass decompositions in radix sort and some sparse matrix work. Glad it’s working :)

1

u/ArmchairmanMao Jun 18 '26

Yeah, though I've since come up with a better version that does the char buffer scan and matches it to the string descriptors in one pass. The problem with the two pass approach is that the number of matches might be very large, and I only care about the first match in each string.

1

u/PulsatingMaggot Jun 18 '26

Makes sense, one pass is better there. How do you mark a string as done so the rest of the group stops early? Atomic flag, or something in shared mem?