NVIDIA / NVIDIA/cutlass

[BUG] Conv2D Frop Gives Wrong Output

Open
#1,690 2 comments 0 reactions 0 assignees View on GitHub

Nobody has claimed this yet.

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

Description

Describe the bug
Conv2D Frop gives wrong output when calculating a small conv in FP32

Steps/Code to reproduce bug
Simply run

import torch
import numpy as np
import cutlass

inputs = torch.from_numpy(np.array([[
    [[1 ,1], [1, 1]],
    [[2 ,2], [2, 2]],
    [[3 ,3], [3, 3]],
    [[0 ,0], [0, 0]],
]]))


inputs = inputs.to(torch.float32).to(memory_format=torch.channels_last)

weights = torch.from_numpy(np.array([[[[1]], [[1]], [[1]], [[1]]]]))
weights = weights.to(torch.float32).to(memory_format=torch.channels_last)

H, C, W, H = inputs.shape
K, C, R, S = weights.shape

N, P, Q, K = cutlass.Conv2d.output_size((N, H, W, C), (K, R, S, C), padding, stride, dilation)


N, P, Q, K = cutlass.Conv2d.output_size((N, H, W, C), (K, R, S, C), padding, stride, dilation)


dtype = torch.float32

plan = cutlass.Conv2dFprop(element_input=dtype, element_weight=dtype, element_C=dtype, element_output=dtype, element_accumulator=dtype)

tensor_C = torch.empty(size=(N, K, P, Q), dtype=dtype, device="cuda").to(memory_format=torch.channels_last)
output = torch.zeros_like(torch.empty(size=(N, K, P, Q), dtype=dtype, device="cuda").to(memory_format=torch.channels_last))
plan.run(inputs, weights, tensor_C, output, stride, padding, dilation, alpha, beta, print_module=print_module)

print(output)

> 
tensor([[[[4., 6.],
          [6., 6.]]]], device='cuda:0')

Expected behavior

Output should be:

tensor([[[[6., 6.],
          [6., 6.]]]], device='cuda:0')

Environment details (please complete the following information):

  • Environment location: Local

Additional context
Add any other context about the problem here.

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

Begin at the cutlass.Conv2dFprop entry point and reduce the supplied FP32, channels-last reproduction to a runnable case, resolving the undefined parameters shown in the snippet. Compare the CUDA output with the stated all-6 tensor; done means the reproduction returns that expected output.

Written by the indexing model from the issue text.

Assessment

Tech stack
cpp, numpy, python, pytorch
Domain
machine-learning, performance
Issue type
Bug
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.