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
+38 −5
Draft

Add int4 & uint4 types to MigraphX #3378

Show file tree
Hide file tree
Changes from all 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
11 changes: 11 additions & 0 deletions src/api/api.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -98,6 +98,10 @@ shape::type_t to_shape_type(migraphx_shape_datatype_t t)
switch(t)
{
case migraphx_shape_tuple_type: return shape::tuple_type;

// case migraphx_shape_uint4_type: return shape::uint4_type;
// case migraphx_shape_int4_type: return shape::int4_type;

#define MIGRAPHX_DETAIL_SHAPE_CASE_CONVERT(x, y) \
case migraphx_shape_##x: return shape::x;
MIGRAPHX_SHAPE_VISIT_TYPES(MIGRAPHX_DETAIL_SHAPE_CASE_CONVERT)
Expand All @@ -111,10 +115,17 @@ migraphx_shape_datatype_t to_shape_type(shape::type_t t)
switch(t)
{
case shape::tuple_type: return migraphx_shape_tuple_type;

// case shape::uint4_type: return migraphx_shape_uint4_type;
// case shape::int4_type: return migraphx_shape_int4_type;

#define MIGRAPHX_DETAIL_SHAPE_CASE_CONVERT(x, y) \
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
16 changes: 13 additions & 3 deletions src/shape.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
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,11 @@
#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},
{"int4", int4_type},
{"uint4_type", uint4_type},
{"uint4", uint4_type}};
return m.at(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
8 changes: 7 additions & 1 deletion test/gpu/jit.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
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