Nota tentang kuliah Mod GPU 1

Profiler

Prestasi komputer adalah mengenai pertukaran masa dan memori. Memandangkan pengiraan peranti jauh lebih mahal, kebanyakan masa, masa adalah keutamaan untuk diambil berat.

Kenapa guna profiler?

1. CUDA tidak segerak jadi tidak boleh menggunakan modul masa Python
2. Profil adalah jauh lebih berkuasa

Alatan

Terdapat tiga pemprofil:

• profil autograd: berangka
• Pytorch profiler: visual
• Pengiraan Nsight NVIDIA

Profil autograd menggunakan torch.cuda.Event() untuk mengukur prestasi.

Profil PyTorch menggunakan profil kaedah() daripada pengurus konteks Profiler obor.profiler untuk menganalisis prestasi.
Anda boleh mengeksport hasil sebagai fail .json dan memuat naiknya ke chrome://tracing/ untuk menggambarkannya.

Demo

Kursus ini menyediakan program mudah untuk menunjukkan cara menggunakan profiler autograd untuk menganalisis prestasi tiga cara untuk melakukan operasi segi empat sama:

• oleh obor.square()
• oleh ** pengendali
• oleh * pengendali

def time_pytorch_function(func, input):
    # CUDA IS ASYNC so can't use python time module
    start = torch.cuda.Event(enable_timing=True)
    end = torch.cuda.Event(enable_timing=True)

    # Warmup
    for _ in range(5):
        func(input)

    start.record()
    func(input)
    end.record()
    torch.cuda.synchronize()
    return start.elapsed_time(end)

time_pytorch_function(torch.square, b)
time_pytorch_function(square_2, b)
time_pytorch_function(square_3, b)

Hasil di bawah dilakukan pada GPU NVIDIA T4.

Profiling torch.square:
Self CPU time total: 10.577ms
Self CUDA time total: 3.266ms

Profiling a * a:
Self CPU time total: 5.417ms
Self CUDA time total: 3.276ms

Profiling a ** 2:
Self CPU time total: 6.183ms
Self CUDA time total: 3.274ms

Ternyata:

• Pengendalian CUDA lebih pantas daripada CPU.
• Pengendali * sedang melakukan operasi aten::multiply dan bukannya aten::pow, dan yang pertama adalah lebih pantas. Ini mungkin kerana penggandaan itu digunakan lebih daripada pow dan ramai pembangun menghabiskan masa untuk mengoptimumkannya.
• Perbezaan prestasi pada CUDA adalah minimum. torch.square ialah operasi paling perlahan memandangkan masa CPU
• aten::square ialah panggilan kepada aten::pow
• Ketiga-tiga kaedah melancarkan kernel cuda yang dipanggil native::vectorized_elementwise_kernel<4, di...

---

Mengintegrasikan kernel CUDA dalam PyTorch

Terdapat beberapa cara untuk melakukannya:

• gunakan load_inline daripada torch.utils.cpp_extendsion
• gunakan Numba iaitu pengkompil yang menyusun fungsi Python yang dihias ke dalam kod mesin yang berjalan pada kedua-dua CPU dan GPU
• gunakan Triton

Kita boleh menggunakan load_inline daripada torch.utils.cpp_extendsion untuk memuatkan kernel CUDA sebagai sambungan PyTorch oleh load_inline(nama, cpp_sources, cuda_sources, functions, with_cuda, build_directory).

from torch.utils.cpp_extension import load_inline

square_matrix_extension = load_inline(
    name='square_matrix_extension',
    cpp_sources=cpp_source,
    cuda_sources=cuda_source,
    functions=['square_matrix'],
    with_cuda=True,
    extra_cuda_cflags=["-O2"],
    build_directory='./load_inline_cuda',
    # extra_cuda_cflags=['--expt-relaxed-constexpr']
)

a = torch.tensor([[1., 2., 3.], [4., 5., 6.]], device='cuda')
print(square_matrix_extension.square_matrix(a))

---

Hands-on

Gunakan profiler autograd pada operasi min

Apabila menggunakan autograd profiler, ingat:

1. Panaskan GPU sebelum merakam supaya GPU memasuki keadaan mantap
2. Purata berbilang larian untuk hasil yang lebih dipercayai

def time_pytorch_function(func, input):
    # CUDA IS ASYNC so can't use python time module
    start = torch.cuda.Event(enable_timing=True)
    end = torch.cuda.Event(enable_timing=True)

    # Warmup
    for _ in range(5):
        func(input)

    start.record()
    func(input)
    end.record()
    torch.cuda.synchronize()
    return start.elapsed_time(end)

time_pytorch_function(torch.square, b)
time_pytorch_function(square_2, b)
time_pytorch_function(square_3, b)

Gunakan pemprofil Pytorch pada operasi min

Profiling torch.square:
Self CPU time total: 10.577ms
Self CUDA time total: 3.266ms

Profiling a * a:
Self CPU time total: 5.417ms
Self CUDA time total: 3.276ms

Profiling a ** 2:
Self CPU time total: 6.183ms
Self CUDA time total: 3.274ms

Melaksanakan kod triton untuk torch.mean()

from torch.utils.cpp_extension import load_inline

square_matrix_extension = load_inline(
    name='square_matrix_extension',
    cpp_sources=cpp_source,
    cuda_sources=cuda_source,
    functions=['square_matrix'],
    with_cuda=True,
    extra_cuda_cflags=["-O2"],
    build_directory='./load_inline_cuda',
    # extra_cuda_cflags=['--expt-relaxed-constexpr']
)

a = torch.tensor([[1., 2., 3.], [4., 5., 6.]], device='cuda')
print(square_matrix_extension.square_matrix(a))

---

Rujukan

• kuliah mod gpu - Github
• Acara - Pytorch
• Profil PyTorch
• Pengiraan Nsight NVIDIA
• torch.utils.cpp_extension.load_inline
• Triton

Atas ialah kandungan terperinci Nota tentang kuliah Mod GPU 1. Untuk maklumat lanjut, sila ikut artikel berkaitan lain di laman web China PHP!