Skip to content

Add TMA example for Hopper H100 #214

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
wants to merge 3 commits into
base: master
Choose a base branch
from
Open

Add TMA example for Hopper H100 #214

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
5925483
Add TMA example
ahendriksen Jun 30, 2023
a3b5b81
Use CTK-provided type for cuTensorMapEncodeTiled
ahendriksen Aug 7, 2023
fc06e3d
Make expected failures more obvious
ahendriksen Aug 7, 2023
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
49 changes: 49 additions & 0 deletions Samples/3_CUDA_Features/globalToShmemTMACopy/README.md
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,49 @@
# globalToShmemTMACopy - Global Memory to Shared Memory TMA Copy

## Description

This sample shows how to use the CUDA driver API and inline PTX assembly to copy
a 2D tile of a tensor into shared memory. It also demonstrates arrive-wait
barrier for synchronization.

## Key Concepts

CUDA Runtime API, CUDA Driver API, PTX ISA, CPP11 CUDA

## Supported SM Architectures

This sample requires compute capability 9.0 or higher.

[SM 9.0 ](https://developer.nvidia.com/cuda-gpus)

## Supported OSes

Linux, Windows, QNX

## Supported CPU Architecture

x86_64, ppc64le, armv7l, aarch64

## CUDA APIs involved

### [CUDA Runtime API](http://docs.nvidia.com/cuda/cuda-runtime-api/index.html)
cudaMalloc, cudaMemcpy, cudaFree, cudaDeviceSynchronize

### [CUDA Driver API](http://docs.nvidia.com/cuda/cuda-driver-api/index.html)
cudaMalloc, cudaMemcpy, cudaFree, cudaDeviceSynchronize

### [CUDA PTX ISA](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html)

## Dependencies needed to build/run
[CPP11](../../../README.md#cpp11)

## Prerequisites

Download and install the [CUDA Toolkit 12.2](https://developer.nvidia.com/cuda-downloads) for your corresponding platform.
Make sure the dependencies mentioned in [Dependencies]() section above are installed.

## Build and Run


## References (for more details)

Loading