Skip to content

Commit 84c098c

Browse files
emankovemankov
authored
Merge pull request #2134 from emankov/amd-develop
[HIPIFY][Device][fix] The missing vector functions and types - Part 2
2 parents 682be66 + e26f2d0 commit 84c098c

File tree

5 files changed

+280
-0
lines changed

5 files changed

+280
-0
lines changed

‎bin/hipify-perl‎

Lines changed: 32 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -6960,6 +6960,10 @@ sub simpleSubstitutions {
69606960
subst("make_char3", "make_char3", "device_function");
69616961
subst("make_char4", "make_char4", "device_function");
69626962
subst("make_half2", "make_half2", "device_function");
6963+
subst("make_int1", "make_int1", "device_function");
6964+
subst("make_int2", "make_int2", "device_function");
6965+
subst("make_int3", "make_int3", "device_function");
6966+
subst("make_int4", "make_int4", "device_function");
69636967
subst("make_short1", "make_short1", "device_function");
69646968
subst("make_short2", "make_short2", "device_function");
69656969
subst("make_short3", "make_short3", "device_function");
@@ -6968,6 +6972,10 @@ sub simpleSubstitutions {
69686972
subst("make_uchar2", "make_uchar2", "device_function");
69696973
subst("make_uchar3", "make_uchar3", "device_function");
69706974
subst("make_uchar4", "make_uchar4", "device_function");
6975+
subst("make_uint1", "make_uint1", "device_function");
6976+
subst("make_uint2", "make_uint2", "device_function");
6977+
subst("make_uint3", "make_uint3", "device_function");
6978+
subst("make_uint4", "make_uint4", "device_function");
69716979
subst("make_ushort1", "make_ushort1", "device_function");
69726980
subst("make_ushort2", "make_ushort2", "device_function");
69736981
subst("make_ushort3", "make_ushort3", "device_function");
@@ -7083,6 +7091,10 @@ sub simpleSubstitutions {
70837091
subst("char3", "char3", "device_type");
70847092
subst("char4", "char4", "device_type");
70857093
subst("cudaRoundMode", "hipRoundMode", "device_type");
7094+
subst("int1", "int1", "device_type");
7095+
subst("int2", "int2", "device_type");
7096+
subst("int3", "int3", "device_type");
7097+
subst("int4", "int4", "device_type");
70867098
subst("nv_bfloat16", "hip_bfloat16", "device_type");
70877099
subst("short1", "short1", "device_type");
70887100
subst("short2", "short2", "device_type");
@@ -7092,6 +7104,10 @@ sub simpleSubstitutions {
70927104
subst("uchar2", "uchar2", "device_type");
70937105
subst("uchar3", "uchar3", "device_type");
70947106
subst("uchar4", "uchar4", "device_type");
7107+
subst("uint1", "uint1", "device_type");
7108+
subst("uint2", "uint2", "device_type");
7109+
subst("uint3", "uint3", "device_type");
7110+
subst("uint4", "uint4", "device_type");
70957111
subst("ushort1", "ushort1", "device_type");
70967112
subst("ushort2", "ushort2", "device_type");
70977113
subst("ushort3", "ushort3", "device_type");
@@ -9645,6 +9661,10 @@ sub transformHostFunctions {
96459661
"make_ushort3",
96469662
"make_ushort2",
96479663
"make_ushort1",
9664+
"make_uint4",
9665+
"make_uint3",
9666+
"make_uint2",
9667+
"make_uint1",
96489668
"make_uchar4",
96499669
"make_uchar3",
96509670
"make_uchar2",
@@ -9653,6 +9673,10 @@ sub transformHostFunctions {
96539673
"make_short3",
96549674
"make_short2",
96559675
"make_short1",
9676+
"make_int4",
9677+
"make_int3",
9678+
"make_int2",
9679+
"make_int1",
96569680
"make_half2",
96579681
"make_char4",
96589682
"make_char3",
@@ -10507,6 +10531,10 @@ sub warnUnsupportedDeviceFunctions {
1050710531
"ushort3",
1050810532
"ushort2",
1050910533
"ushort1",
10534+
"uint4",
10535+
"uint3",
10536+
"uint2",
10537+
"uint1",
1051010538
"uchar4",
1051110539
"uchar3",
1051210540
"uchar2",
@@ -10516,6 +10544,10 @@ sub warnUnsupportedDeviceFunctions {
1051610544
"short2",
1051710545
"short1",
1051810546
"nv_bfloat16",
10547+
"int4",
10548+
"int3",
10549+
"int2",
10550+
"int1",
1051910551
"cudaRoundZero",
1052010552
"cudaRoundPosInf",
1052110553
"cudaRoundNearest",

‎docs/reference/tables/CUDA_Device_API_supported_by_HIP.md‎

Lines changed: 16 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -815,6 +815,10 @@
815815
|`make_char3`| | | | |`make_char3`|1.6.0| | | | |
816816
|`make_char4`| | | | |`make_char4`|1.6.0| | | | |
817817
|`make_half2`|12.2| | | |`make_half2`|4.5.0| | | | |
818+
|`make_int1`| | | | |`make_int1`|1.6.0| | | | |
819+
|`make_int2`| | | | |`make_int2`|1.6.0| | | | |
820+
|`make_int3`| | | | |`make_int3`|1.6.0| | | | |
821+
|`make_int4`| | | | |`make_int4`|1.6.0| | | | |
818822
|`make_short1`| | | | |`make_short1`|1.6.0| | | | |
819823
|`make_short2`| | | | |`make_short2`|1.6.0| | | | |
820824
|`make_short3`| | | | |`make_short3`|1.6.0| | | | |
@@ -823,6 +827,10 @@
823827
|`make_uchar2`| | | | |`make_uchar2`|1.6.0| | | | |
824828
|`make_uchar3`| | | | |`make_uchar3`|1.6.0| | | | |
825829
|`make_uchar4`| | | | |`make_uchar4`|1.6.0| | | | |
830+
|`make_uint1`| | | | |`make_uint1`|1.6.0| | | | |
831+
|`make_uint2`| | | | |`make_uint2`|1.6.0| | | | |
832+
|`make_uint3`| | | | |`make_uint3`|1.6.0| | | | |
833+
|`make_uint4`| | | | |`make_uint4`|1.6.0| | | | |
826834
|`make_ushort1`| | | | |`make_ushort1`|1.6.0| | | | |
827835
|`make_ushort2`| | | | |`make_ushort2`|1.6.0| | | | |
828836
|`make_ushort3`| | | | |`make_ushort3`|1.6.0| | | | |
@@ -978,6 +986,10 @@
978986
|`cudaRoundNearest`| | | | |`hipRoundNearest`|7.0.0| | | | |
979987
|`cudaRoundPosInf`| | | | |`hipRoundPosInf`|7.0.0| | | | |
980988
|`cudaRoundZero`| | | | |`hipRoundZero`|7.0.0| | | | |
989+
|`int1`| | | | |`int1`|1.6.0| | | | |
990+
|`int2`| | | | |`int2`|1.6.0| | | | |
991+
|`int3`| | | | |`int3`|1.6.0| | | | |
992+
|`int4`| | | | |`int4`|1.6.0| | | | |
981993
|`nv_bfloat16`|11.0| | | |`hip_bfloat16`|3.5.0| | | | |
982994
|`nv_bfloat162`|11.0| | | | | | | | | |
983995
|`short1`| | | | |`short1`|1.6.0| | | | |
@@ -988,6 +1000,10 @@
9881000
|`uchar2`| | | | |`uchar2`|1.6.0| | | | |
9891001
|`uchar3`| | | | |`uchar3`|1.6.0| | | | |
9901002
|`uchar4`| | | | |`uchar4`|1.6.0| | | | |
1003+
|`uint1`| | | | |`uint1`|1.6.0| | | | |
1004+
|`uint2`| | | | |`uint2`|1.6.0| | | | |
1005+
|`uint3`| | | | |`uint3`|1.6.0| | | | |
1006+
|`uint4`| | | | |`uint4`|1.6.0| | | | |
9911007
|`ushort1`| | | | |`ushort1`|1.6.0| | | | |
9921008
|`ushort2`| | | | |`ushort2`|1.6.0| | | | |
9931009
|`ushort3`| | | | |`ushort3`|1.6.0| | | | |

‎src/CUDA2HIP_Device_functions.cpp‎

Lines changed: 16 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -939,6 +939,14 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DEVICE_FUNCTION_MAP {
939939
{"make_ushort2", {"make_ushort2", "", CONV_DEVICE_FUNC, API_RUNTIME, 1}},
940940
{"make_ushort3", {"make_ushort3", "", CONV_DEVICE_FUNC, API_RUNTIME, 1}},
941941
{"make_ushort4", {"make_ushort4", "", CONV_DEVICE_FUNC, API_RUNTIME, 1}},
942+
{"make_int1", {"make_int1", "", CONV_DEVICE_FUNC, API_RUNTIME, 1}},
943+
{"make_int2", {"make_int2", "", CONV_DEVICE_FUNC, API_RUNTIME, 1}},
944+
{"make_int3", {"make_int3", "", CONV_DEVICE_FUNC, API_RUNTIME, 1}},
945+
{"make_int4", {"make_int4", "", CONV_DEVICE_FUNC, API_RUNTIME, 1}},
946+
{"make_uint1", {"make_uint1", "", CONV_DEVICE_FUNC, API_RUNTIME, 1}},
947+
{"make_uint2", {"make_uint2", "", CONV_DEVICE_FUNC, API_RUNTIME, 1}},
948+
{"make_uint3", {"make_uint3", "", CONV_DEVICE_FUNC, API_RUNTIME, 1}},
949+
{"make_uint4", {"make_uint4", "", CONV_DEVICE_FUNC, API_RUNTIME, 1}},
942950
};
943951

944952
const std::map<llvm::StringRef, cudaAPIversions> CUDA_DEVICE_FUNCTION_VER_MAP {
@@ -1737,6 +1745,14 @@ const std::map<llvm::StringRef, hipAPIversions> HIP_DEVICE_FUNCTION_VER_MAP {
17371745
{"make_ushort2", {HIP_1060, HIP_0, HIP_0 }},
17381746
{"make_ushort3", {HIP_1060, HIP_0, HIP_0 }},
17391747
{"make_ushort4", {HIP_1060, HIP_0, HIP_0 }},
1748+
{"make_int1", {HIP_1060, HIP_0, HIP_0 }},
1749+
{"make_int2", {HIP_1060, HIP_0, HIP_0 }},
1750+
{"make_int3", {HIP_1060, HIP_0, HIP_0 }},
1751+
{"make_int4", {HIP_1060, HIP_0, HIP_0 }},
1752+
{"make_uint1", {HIP_1060, HIP_0, HIP_0 }},
1753+
{"make_uint2", {HIP_1060, HIP_0, HIP_0 }},
1754+
{"make_uint3", {HIP_1060, HIP_0, HIP_0 }},
1755+
{"make_uint4", {HIP_1060, HIP_0, HIP_0 }},
17401756
};
17411757

17421758
const std::map<llvm::StringRef, cudaAPIChangedVersions> CUDA_DEVICE_FUNCTION_CHANGED_VER_MAP {

‎src/CUDA2HIP_Device_types.cpp‎

Lines changed: 16 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -108,6 +108,14 @@ const std::map<llvm::StringRef, hipCounter> CUDA_DEVICE_TYPE_NAME_MAP {
108108
{"ushort2", {"ushort2", "", CONV_DEVICE_TYPE, API_RUNTIME, 2}},
109109
{"ushort3", {"ushort3", "", CONV_DEVICE_TYPE, API_RUNTIME, 2}},
110110
{"ushort4", {"ushort4", "", CONV_DEVICE_TYPE, API_RUNTIME, 2}},
111+
{"int1", {"int1", "", CONV_DEVICE_TYPE, API_RUNTIME, 2}},
112+
{"int2", {"int2", "", CONV_DEVICE_TYPE, API_RUNTIME, 2}},
113+
{"int3", {"int3", "", CONV_DEVICE_TYPE, API_RUNTIME, 2}},
114+
{"int4", {"int4", "", CONV_DEVICE_TYPE, API_RUNTIME, 2}},
115+
{"uint1", {"uint1", "", CONV_DEVICE_TYPE, API_RUNTIME, 2}},
116+
{"uint2", {"uint2", "", CONV_DEVICE_TYPE, API_RUNTIME, 2}},
117+
{"uint3", {"uint3", "", CONV_DEVICE_TYPE, API_RUNTIME, 2}},
118+
{"uint4", {"uint4", "", CONV_DEVICE_TYPE, API_RUNTIME, 2}},
111119
};
112120

113121
const std::map<llvm::StringRef, cudaAPIversions> CUDA_DEVICE_TYPE_NAME_VER_MAP {
@@ -225,6 +233,14 @@ const std::map<llvm::StringRef, hipAPIversions> HIP_DEVICE_TYPE_NAME_VER_MAP {
225233
{"ushort2", {HIP_1060, HIP_0, HIP_0 }},
226234
{"ushort3", {HIP_1060, HIP_0, HIP_0 }},
227235
{"ushort4", {HIP_1060, HIP_0, HIP_0 }},
236+
{"int1", {HIP_1060, HIP_0, HIP_0 }},
237+
{"int2", {HIP_1060, HIP_0, HIP_0 }},
238+
{"int3", {HIP_1060, HIP_0, HIP_0 }},
239+
{"int4", {HIP_1060, HIP_0, HIP_0 }},
240+
{"uint1", {HIP_1060, HIP_0, HIP_0 }},
241+
{"uint2", {HIP_1060, HIP_0, HIP_0 }},
242+
{"uint3", {HIP_1060, HIP_0, HIP_0 }},
243+
{"uint4", {HIP_1060, HIP_0, HIP_0 }},
228244

229245
{"rocblas_half", {HIP_1050, HIP_0, HIP_0 }},
230246
{"rocblas_bfloat16", {HIP_3050, HIP_0, HIP_0 }},

0 commit comments

Comments
 (0)