Skip to content
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

Jump to bottom

Add int4 & uint4 types to MigraphX #3378

Draft
wants to merge 4 commits into
base: develop
Choose a base branch
from
Draft

Add int4 & uint4 types to MigraphX #3378

Show file tree
Hide file tree
Changes from 3 commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
45bc16b
add int4 & uint4 types
lakhinderwalia Aug 15, 2024
3d4ff7a
misc. fixes for int4 types
lakhinderwalia Aug 15, 2024
8f67ab4
fixes for broken test and warnings
lakhinderwalia Aug 15, 2024
45e0397
Review comments on int4/uint4 type shape
lakhinderwalia Aug 16, 2024
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
3 changes: 3 additions & 0 deletions src/api/api.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -115,6 +115,9 @@ migraphx_shape_datatype_t to_shape_type(shape::type_t t)
case shape::x: return migraphx_shape_##x;
MIGRAPHX_SHAPE_VISIT_TYPES(MIGRAPHX_DETAIL_SHAPE_CASE_CONVERT)
#undef MIGRAPHX_DETAIL_SHAPE_CASE_CONVERT
case shape::uint4_type:
case shape::int4_type:
break;
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This wrong, the enum values should be exposed in the API and not throw an error.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

OK.

}
MIGRAPHX_THROW(migraphx_status_bad_param, "Unknown type");
}
Expand Down
6 changes: 5 additions & 1 deletion src/include/migraphx/shape.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -68,7 +68,9 @@ struct MIGRAPHX_EXPORT shape
#define MIGRAPHX_SHAPE_GENERATE_ENUM_TYPES(x, t) x,
enum type_t
{
MIGRAPHX_SHAPE_VISIT_TYPES(MIGRAPHX_SHAPE_GENERATE_ENUM_TYPES) tuple_type
MIGRAPHX_SHAPE_VISIT_TYPES(MIGRAPHX_SHAPE_GENERATE_ENUM_TYPES) tuple_type,
uint4_type,
int4_type
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

We should have a macro to list these types.

};
#undef MIGRAPHX_SHAPE_GENERATE_ENUM_TYPES

Expand Down Expand Up @@ -381,6 +383,8 @@ struct MIGRAPHX_EXPORT shape
{
switch(t)
{
case uint4_type:
case int4_type:
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The tuple visitor should not be called for these types. It should throw an error.

case tuple_type: {
tv();
return;
Expand Down
22 changes: 15 additions & 7 deletions src/shape.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -48,20 +48,20 @@

shape_impl(shape::type_t t) : m_type(t), m_lens({1}), m_strides({0}), m_standard(true)
{
assert(t != shape::tuple_type);
assert(t != shape::tuple_type && t != shape::uint4_type && t != shape::int4_type);
}

shape_impl(shape::type_t t, std::vector<std::size_t> l)
: m_type(t), m_lens(std::move(l)), m_standard(true)
{
assert(t != shape::tuple_type);
assert(t != shape::tuple_type && t != shape::uint4_type && t != shape::int4_type);
Copy link
Collaborator

@pfultz2 pfultz2 Aug 15, 2024 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This assert shouldn't be changed here. It should be valid to construct a shape of int4_type otherwise whats the point of adding this type.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Thanks. This one is definitely required.

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

What do you mean?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I meant you are correct. This constructor code, as you suggest, is definitely required for correct functioning.
(I had just quickly hacked those places to make sure the CI would pass the compile stage -- as some switch statements etc. were failing.) Thanks.

this->calculate_strides();
}

shape_impl(shape::type_t t, std::vector<std::size_t> l, std::vector<std::size_t> s)
: m_type(t), m_lens(std::move(l)), m_strides(std::move(s))
{
assert(t != shape::tuple_type);
assert(t != shape::tuple_type && t != shape::uint4_type && t != shape::int4_type);
assert(m_lens.size() == m_strides.size());
m_standard = this->elements() == this->element_space() and not skips() and
std::is_sorted(m_strides.rbegin(), m_strides.rend());
Expand Down Expand Up @@ -122,7 +122,7 @@
{
if(not m_dyn_dims.empty())
{
auto maxes = max_lens();
auto maxes = max_lens();
std::size_t max_val = std::numeric_limits<std::size_t>::max();

return std::accumulate(
Expand Down Expand Up @@ -224,7 +224,9 @@
{
static const std::vector<shape::type_t> result = {
#define MIGRAPHX_GENERATE_TYPE_VECTOR(x, t) x,
MIGRAPHX_SHAPE_VISIT_TYPES(MIGRAPHX_GENERATE_TYPE_VECTOR) tuple_type};
MIGRAPHX_SHAPE_VISIT_TYPES(MIGRAPHX_GENERATE_TYPE_VECTOR) tuple_type,
int4_type,
uint4_type};

Check warning on line 229 in src/shape.cpp

View check run for this annotation

Codecov / codecov/patch

src/shape.cpp#L229 

Added line #L229 was not covered by tests
return result;
}

Expand All @@ -233,6 +235,8 @@
switch(t)
{
case tuple_type: return "tuple_type";
case int4_type: return "int4_type";
case uint4_type: return "uint4_type";

Check warning on line 239 in src/shape.cpp

View check run for this annotation

Codecov / codecov/patch

src/shape.cpp#L238-L239 

Added lines #L238 - L239 were not covered by tests
#define MIGRAPHX_SHAPE_GENERATE_TYPE_NAME_CASE(x, t) \
case x: return #x;
MIGRAPHX_SHAPE_VISIT_TYPES(MIGRAPHX_SHAPE_GENERATE_TYPE_NAME_CASE)
Expand All @@ -246,6 +250,8 @@
switch(t)
{
case tuple_type: MIGRAPHX_THROW("No C++ type for tuple");
case int4_type: MIGRAPHX_THROW("No C++ type for int4_type");
case uint4_type: MIGRAPHX_THROW("No C++ type for uint4_type");

Check warning on line 254 in src/shape.cpp

View check run for this annotation

Codecov / codecov/patch

src/shape.cpp#L253-L254 

Added lines #L253 - L254 were not covered by tests
#define MIGRAPHX_SHAPE_GENERATE_CPP_TYPE_CASE(x, t) \
case x: return #t;
MIGRAPHX_SHAPE_VISIT_TYPES(MIGRAPHX_SHAPE_GENERATE_CPP_TYPE_CASE)
Expand Down Expand Up @@ -728,7 +734,9 @@
#define MIGRAPHX_SHAPE_GENERATE_TYPE_STRING_MAP(x, t) {#x, x}, {#t, x},
MIGRAPHX_SHAPE_VISIT_TYPES(MIGRAPHX_SHAPE_GENERATE_TYPE_STRING_MAP){"tuple_type",
tuple_type},
{"tuple", tuple_type}};
{"tuple", tuple_type},
{"int4_type", int4_type},
{"uint4_type", uint4_type}};
return m.at(s);
}

Expand All @@ -744,7 +752,7 @@
auto subs = flatten(s.sub_shapes());
result.insert(result.end(), subs.begin(), subs.end());
}
else
else if(s.type() != shape::uint4_type && s.type() != shape::int4_type)

Check warning on line 755 in src/shape.cpp

View check run for this annotation

Codecov / codecov/patch

src/shape.cpp#L755 

Added line #L755 was not covered by tests
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This shouldn't be changed, we should still be able to flatten an tuples of int4_type.

{
result.push_back(s);
}
Expand Down
2 changes: 2 additions & 0 deletions src/targets/gpu/gemm_impl.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -69,6 +69,8 @@ rocblas_datatype get_type(shape::type_t type)
case shape::uint16_type:
case shape::int16_type:
case shape::int64_type:
case shape::int4_type:
case shape::uint4_type:
case shape::uint64_type: MIGRAPHX_THROW("ROCBLAS_GEMM: data type not supported!");
}

Expand Down
10 changes: 8 additions & 2 deletions test/gpu/jit.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -239,7 +239,7 @@ TEST_CASE(code_object_hip)

std::vector<migraphx::shape> expected_inputs = {input, input};
auto co = migraphx::make_op("gpu::code_object",
{{"code_object", migraphx::value::binary{binaries.front()}},
{{"code_object", migraphx::value::binary{binaries.front()}},
{"symbol_name", "add_2"},
{"global", input.elements()},
{"local", 1024},
Expand Down Expand Up @@ -350,7 +350,11 @@ TEST_CASE(compile_math)
auto vec_sizes = {2, 4, 6};
for(auto&& t : migraphx::shape::types())
{
if(contains({migraphx::shape::bool_type, migraphx::shape::tuple_type}, t))
if(contains({migraphx::shape::bool_type,
migraphx::shape::tuple_type,
migraphx::shape::int4_type,
migraphx::shape::uint4_type},
t))
continue;
auto name = migraphx::shape::cpp_type(t);
if(t == migraphx::shape::half_type)
Expand Down Expand Up @@ -403,6 +407,8 @@ TEST_CASE(assert_type_min_max)
for(auto&& t : migraphx::shape::types())
{
if(contains({migraphx::shape::bool_type,
migraphx::shape::uint4_type,
migraphx::shape::int4_type,
migraphx::shape::fp8e4m3fnuz_type,
migraphx::shape::tuple_type},
t))
Expand Down
Loading