NVIDIA / NVIDIA/cuCollections

[FEA] Add tests that verify vector loads are generated when expected

Open
#213 0 comments 0 reactions 0 assignees View on GitHub

Nobody has claimed this yet.

P2: Nice to have topic: performance type: feature request
Dominant language
Cuda
Stars
667
Forks
120
Avg merge
7d 5h
Merged PRs (30d)
4

Description

Generating vector load instructions is important for performance.

However, the compiler can be very fussy about actually generating them.

For example, https://godbolt.org/z/sfdaj695P shows prototyping code where all versions of nvcc prior to 11.7 would generate a ld.global.v4.u32 as expected, but in 11.7 it generates 4x ld.global.u32. What's worse is that it checks the condition after each load, so they won't be coalesced.

It would be nice to have a simple test that verifies vector loads (or stores) are generated in code where we expect them to be generated.

We should be sure to verify that vector SASS instructions are generated (and not just PTX) as I've seen cases where PTX will have vector loads but SASS does not.

This could be as simple as using cuobjdump -sass a.out | grep LDG.E.128.SYS, but there's probably more clever ways to do it.

Contributor guide

Open the contributing guide

First steps

  1. Read the whole issue, then the project's contributing guide.
  2. Comment on the issue to say you are picking it up — it saves two people doing the same work.
  3. Fork the repository and make your change on a branch.
  4. Open a pull request that references the issue number.

Research direction

No repository files or existing tests are named; start by locating the CUDA test infrastructure and the build output used with nvcc. Use cuobjdump SASS inspection, such as checking for LDG.E.128.SYS, and consider both vector loads and stores; done means the test reliably verifies the expected vector SASS instructions rather than only PTX.

Written by the indexing model from the issue text.

Assessment

Tech stack
cpp
Domain
performance, testing-qa
Issue type
Feature
Difficulty
4/5
Estimated time
3-5 days
Activity status
Stale
Clarity
Needs clarification
Newbie friendliness
35/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.