NVIDIA / NVIDIA/cutlass

[QST] How do I import the `cutlass::gemm::kernel::GemmUniversal`

Open
#1,814 10 comments 0 reactions 0 assignees View on GitHub

Nobody has claimed this yet.

? - Needs Triage inactive-30d inactive-90d question
Dominant language
C++
Stars
10.5k
Forks
2.1k
Avg merge
3d 11h
Merged PRs (30d)
7

Description

Following up on #1291, I am finally making the time to properly understand Cutlass. I can build and run the examples in the repo, but I am having a lot of difficulty understanding the type errors I am getting when I try to import the GemmUniversal classes. For example...

#include <iostream>

#include "cutlass/cutlass.h"
#include "cute/tensor.hpp"
#include "cute/atom/mma_atom.hpp"

#include "cutlass/numeric_types.h"

#include "cutlass/gemm/device/gemm_universal_adapter.h"
#include "default_gemm_configuration.hpp"

using namespace cute;

/////////////////////////////////////////////////////////////////////////////////////////////////

using Config = cutlass::gemm::device::DefaultGemmConfigurationToCutlass3Types<
	cutlass::arch::OpClassSimt, cutlass::arch::Sm50,
	double, cutlass::layout::ColumnMajor,
	double, cutlass::layout::ColumnMajor,
	double, cutlass::layout::ColumnMajor,
	double>;

using GemmKernel = cutlass::gemm::kernel::GemmUniversal<
	Shape<int, int, int, int>,
	Config::CollectiveMainloop,
	Config::CollectiveEpilogue
>;

using Gemm = cutlass::gemm::device::GemmUniversalAdapter<GemmKernel>;

I've isolated this piece of code from one of the tests and am trying to bring it into another project without success. Even when I seemingly have all the necessary includes, I see something like this...

Image

Probably I shouldn't be trying to import test harness code, but the trouble is that there are no clean examples how to use the universal kernel.

https://github.com/NVIDIA/cutlass/blob/main/media/docs/quickstart.md#launching-a-gemm-kernel-using-cutlass-30-or-newer

The quickstart one is clean but doesn't typecheck for me.

There are some examples like the 48_hopper_warp_specialized_gemm which do have use examples of it, but when I get rid of the device arch checking macros, I just get a ton of type errors. Those kinds of examples aren't usable.

I am going to try recompiling the Cutlass library for 9.0 devices just to see if something changes.

But otherwise, could I get some help to start me off with Cutlass? Thanks.

Contributor guide

No contributing guide indexed for this repository

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

Start with media/docs/quickstart.md and compare its universal-kernel example with the 48_hopper_warp_specialized_gemm examples and the included test code. Reproduce the reported type errors in the shown GemmUniversal setup, then clarify the required includes and usage in a standalone example that typechecks.

Written by the indexing model from the issue text.

Assessment

Tech stack
cpp
Domain
hpc
Issue type
Documentation
Difficulty
4/5
Estimated time
3-5 days
Activity status
Stale
Clarity
Needs clarification
Newbie friendliness
25/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.