diff --git a/bin/hipify-perl b/bin/hipify-perl index e6cae6c6..3899e16c 100755 --- a/bin/hipify-perl +++ b/bin/hipify-perl @@ -2991,32 +2991,32 @@ sub transformKernelLaunch { no warnings qw/uninitialized/; my $k = 0; - # kern<...><<>>() syntax - $k += s/([:|\w]+)\s*<(.+)>\s*<<<\s*([^,]+|[\w:]*\([\w|,|:]+\))\s*,\s*([^,]+|[\w:]*\([\w|,|:]+\))\s*>>>\s*\(\s*\)/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), $3, $4, 0, 0)/g; - # kern<...><<>>(...) syntax - $k += s/([:|\w]+)\s*<(.+)>\s*<<<\s*([^,]+|[\w:]*\([\w|,|:]+\))\s*,\s*([^,]+|[\w:]*\([\w|,|:]+\))\s*>>>\s*\(/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), $3, $4, 0, 0, /g; - # kern<<>>() syntax - $k += s/([:|\w]+)\s*<<<\s*([^,]+|[\w:]*\([\w|,|:]+\))\s*,\s*([^,]+|[\w:]*\([\w|,|:]+\))\s*>>>\s*\(\s*\)/hipLaunchKernelGGL($1, $2, $3, 0, 0)/g; - # kern<<>>(...) syntax - $k += s/([:|\w]+)\s*<<<\s*([^,]+|[\w:]*\([\w|,|:]+\))\s*,\s*([^,]+|[\w:]*\([\w|,|:]+\))\s*>>>\s*\(/hipLaunchKernelGGL($1, $2, $3, 0, 0, /g; + # kern<...><<>>() syntax + $k += s/([:|\w]+)\s*<(.+)>\s*<<<\s*([^,\(\)]+|[\w\s:]*\([\w|\s|,|:|\+|\*|\-|\/|(?R)]+\))\s*,\s*([^,\(\)]+|[\w\s:]*\([\w|\s|,|:|\+|\*|\-|\/|(?R)]+\))\s*,\s*([^,\(\)]+|[\w\s:]*\([\w|\s|,|:|\+|\*|\-|\/|(?R)]+\))\s*,\s*([^,\(\)]+|[\w\s:]*\([\w|\s|,|:|\+|\*|\-|\/|(?R)]+\))\s*>>>\s*\(\s*\)/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), $3, $4, $5, $6)/g; + # kern<...><<>>(...) syntax + $k += s/([:|\w]+)\s*<(.+)>\s*<<<\s*([^,\(\)]+|[\w\s:]*\([\w|\s|,|:|\+|\*|\-|\/|(?R)]+\))\s*,\s*([^,\(\)]+|[\w\s:]*\([\w|\s|,|:|\+|\*|\-|\/|(?R)]+\))\s*,\s*([^,\(\)]+|[\w\s:]*\([\w|\s|,|:|\+|\*|\-|\/|(?R)]+\))\s*,\s*([^,\(\)]+|[\w\s:]*\([\w|\s|,|:|\+|\*|\-|\/|(?R)]+\))\s*>>>\s*\(/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), $3, $4, $5, $6, /g; + # kern<<>>() syntax + $k += s/([:|\w]+)\s*<<<\s*([^,\(\)]+|[\w\s:]*\([\w|\s|,|:|\+|\*|\-|\/|(?R)]+\))\s*,\s*([^,\(\)]+|[\w\s:]*\([\w|\s|,|:|\+|\*|\-|\/|(?R)]+\))\s*,\s*([^,\(\)]+|[\w\s:]*\([\w|\s|,|:|\+|\*|\-|\/|(?R)]+\))\s*,\s*([^,\(\)]+|[\w\s:]*\([\w|\s|,|:|\+|\*|\-|\/|(?R)]+\))\s*>>>\s*\(\s*\)/hipLaunchKernelGGL($1, $2, $3, $4, $5)/g; + # kern<<>>(...) syntax + $k += s/([:|\w]+)\s*<<<\s*([^,\(\)]+|[\w\s:]*\([\w|\s|,|:|\+|\*|\-|\/|(?R)]+\))\s*,\s*([^,\(\)]+|[\w\s:]*\([\w|\s|,|:|\+|\*|\-|\/|(?R)]+\))\s*,\s*([^,\(\)]+|[\w\s:]*\([\w|\s|,|:|\+|\*|\-|\/|(?R)]+\))\s*,\s*([^,\(\)]+|[\w\s:]*\([\w|\s|,|:|\+|\*|\-|\/|(?R)]+\))\s*>>>\s*\(/hipLaunchKernelGGL($1, $2, $3, $4, $5, /g; # kern<...><<>>() syntax - $k += s/([:|\w]+)\s*<(.+)>\s*<<<\s*([^,]+|[\w:]*\([\w|,|:]+\))\s*,\s*([^,]+|[\w:]*\([\w|,|:]+\))\s*,\s*([^,]+|[\w:]*\([\w|,|:]+\))\s*>>>\s*\(\s*\)/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), $3, $4, $5, 0)/g; + $k += s/([:|\w]+)\s*<(.+)>\s*<<<\s*([^,\(\)]+|[\w\s:]*\([\w|\s|,|:|\+|\*|\-|\/|(?R)]+\))\s*,\s*([^,\(\)]+|[\w\s:]*\([\w|\s|,|:|\+|\*|\-|\/|(?R)]+\))\s*,\s*([^,\(\)]+|[\w\s:]*\([\w|\s|,|:|\+|\*|\-|\/|(?R)]+\))\s*>>>\s*\(\s*\)/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), $3, $4, $5, 0)/g; # kern<...><<>>(...) syntax - $k += s/([:|\w]+)\s*<(.+)>\s*<<<\s*([^,]+|[\w:]*\([\w|,|:]+\))\s*,\s*([^,]+|[\w:]*\([\w|,|:]+\))\s*,\s*([^,]+|[\w:]*\([\w|,|:]+\))\s*>>>\s*\(/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), $3, $4, $5, 0, /g; + $k += s/([:|\w]+)\s*<(.+)>\s*<<<\s*([^,\(\)]+|[\w\s:]*\([\w|\s|,|:|\+|\*|\-|\/|(?R)]+\))\s*,\s*([^,\(\)]+|[\w\s:]*\([\w|\s|,|:|\+|\*|\-|\/|(?R)]+\))\s*,\s*([^,\(\)]+|[\w\s:]*\([\w|\s|,|:|\+|\*|\-|\/|(?R)]+\))\s*>>>\s*\(/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), $3, $4, $5, 0, /g; # kern<<>>() syntax - $k += s/([:|\w]+)\s*<<<\s*([^,]+|[\w:]*\([\w|,|:]+\))\s*,\s*([^,]+|[\w:]*\([\w|,|:]+\))\s*,\s*([^,]+|[\w:]*\([\w|,|:]+\))\s*>>>\s*\(\s*\)/hipLaunchKernelGGL($1, $2, $3, $4, 0)/g; + $k += s/([:|\w]+)\s*<<<\s*([^,\(\)]+|[\w\s:]*\([\w|\s|,|:|\+|\*|\-|\/|(?R)]+\))\s*,\s*([^,\(\)]+|[\w\s:]*\([\w|\s|,|:|\+|\*|\-|\/|(?R)]+\))\s*,\s*([^,\(\)]+|[\w\s:]*\([\w|\s|,|:|\+|\*|\-|\/|(?R)]+\))\s*>>>\s*\(\s*\)/hipLaunchKernelGGL($1, $2, $3, $4, 0)/g; # kern<<>>(...) syntax - $k += s/([:|\w]+)\s*<<<\s*([^,]+|[\w:]*\([\w|,|:]+\))\s*,\s*([^,]+|[\w:]*\([\w|,|:]+\))\s*,\s*([^,]+|[\w:]*\([\w|,|:]+\))\s*>>>\s*\(/hipLaunchKernelGGL($1, $2, $3, $4, 0, /g; + $k += s/([:|\w]+)\s*<<<\s*([^,\(\)]+|[\w\s:]*\([\w|\s|,|:|\+|\*|\-|\/|(?R)]+\))\s*,\s*([^,\(\)]+|[\w\s:]*\([\w|\s|,|:|\+|\*|\-|\/|(?R)]+\))\s*,\s*([^,\(\)]+|[\w\s:]*\([\w|\s|,|:|\+|\*|\-|\/|(?R)]+\))\s*>>>\s*\(/hipLaunchKernelGGL($1, $2, $3, $4, 0, /g; - # kern<...><<>>() syntax - $k += s/([:|\w]+)\s*<(.+)>\s*<<<\s*([^,]+|[\w:]*\([\w|,|:]+\))\s*,\s*([^,]+|[\w:]*\([\w|,|:]+\))\s*,\s*([^,]+|[\w:]*\([\w|,|:]+\))\s*,\s*([^,]+|[\w:]*\([\w|,|:]+\))\s*>>>\s*\(\s*\)/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), $3, $4, $5, $6)/g; - # kern<...><<>>(...) syntax - $k += s/([:|\w]+)\s*<(.+)>\s*<<<\s*([^,]+|[\w:]*\([\w|,|:]+\))\s*,\s*([^,]+|[\w:]*\([\w|,|:]+\))\s*,\s*([^,]+|[\w:]*\([\w|,|:]+\))\s*,\s*([^,]+|[\w:]*\([\w|,|:]+\))\s*>>>\s*\(/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), $3, $4, $5, $6, /g; - # kern<<>>() syntax - $k += s/([:|\w]+)\s*<<<\s*([^,]+|[\w:]*\([\w|,|:]+\))\s*,\s*([^,]+|[\w:]*\([\w|,|:]+\))\s*,\s*([^,]+|[\w:]*\([\w|,|:]+\))\s*,\s*([^,]+|[\w:]*\([\w|,|:]+\))\s*>>>\s*\(\s*\)/hipLaunchKernelGGL($1, $2, $3, $4, $5)/g; - # kern<<>>(...) syntax - $k += s/([:|\w]+)\s*<<<\s*([^,]+|[\w:]*\([\w|,|:]+\))\s*,\s*([^,]+|[\w:]*\([\w|,|:]+\))\s*,\s*([^,]+|[\w:]*\([\w|,|:]+\))\s*,\s*([^,]+|[\w:]*\([\w|,|:]+\))\s*>>>\s*\(/hipLaunchKernelGGL($1, $2, $3, $4, $5, /g; + # kern<...><<>>() syntax + $k += s/([:|\w]+)\s*<(.+)>\s*<<<\s*([^,\(\)]+|[\w\s:]*\([\w|\s|,|:|\+|\*|\-|\/|(?R)]+\))\s*,\s*([^,\(\)]+|[\w\s:]*\([\w|\s|,|:|\+|\*|\-|\/|(?R)]+\))\s*>>>\s*\(\s*\)/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), $3, $4, 0, 0)/g; + # kern<...><<>>(...) syntax + $k += s/([:|\w]+)\s*<(.+)>\s*<<<\s*([^,\(\)]+|[\w\s:]*\([\w|\s|,|:|\+|\*|\-|\/|(?R)]+\))\s*,\s*([^,\(\)]+|[\w\s:]*\([\w|\s|,|:|\+|\*|\-|\/|(?R)]+\))\s*>>>\s*\(/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), $3, $4, 0, 0, /g; + # kern<<>>() syntax + $k += s/([:|\w]+)\s*<<<\s*([^,\(\)]+|[\w\s:]*\([\w|\s|,|:|\+|\*|\-|\/|(?R)]+\))\s*,\s*([^,\(\)]+|[\w\s:]*\([\w|\s|,|:|\+|\*|\-|\/|(?R)]+\))\s*>>>\s*\(\s*\)/hipLaunchKernelGGL($1, $2, $3, 0, 0)/g; + # kern<<>>(...) syntax + $k += s/([:|\w]+)\s*<<<\s*([^,\(\)]+|[\w\s:]*\([\w|\s|,|:|\+|\*|\-|\/|(?R)]+\))\s*,\s*([^,\(\)]+|[\w\s:]*\([\w|\s|,|:|\+|\*|\-|\/|(?R)]+\))\s*>>>\s*\(/hipLaunchKernelGGL($1, $2, $3, 0, 0, /g; if ($k) { $ft{'kernel_launch'} += $k; diff --git a/src/CUDA2HIP_Perl.cpp b/src/CUDA2HIP_Perl.cpp index 5a44292d..f316a4d4 100644 --- a/src/CUDA2HIP_Perl.cpp +++ b/src/CUDA2HIP_Perl.cpp @@ -310,32 +310,33 @@ namespace perl { string s_k = "$k += s/([:|\\w]+)\\s*"; - *streamPtr.get() << tab << "# kern<...><<>>() syntax" << endl; - *streamPtr.get() << tab << s_k << "<(.+)>\\s*<<<\\s*([^,]+|[\\w:]*\\([\\w|,|:]+\\))\\s*,\\s*([^,]+|[\\w:]*\\([\\w|,|:]+\\))\\s*>>>\\s*\\(\\s*\\)/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), $3, $4, 0, 0)/g;" << endl; - *streamPtr.get() << tab << "# kern<...><<>>(...) syntax" << endl; - *streamPtr.get() << tab << s_k << "<(.+)>\\s*<<<\\s*([^,]+|[\\w:]*\\([\\w|,|:]+\\))\\s*,\\s*([^,]+|[\\w:]*\\([\\w|,|:]+\\))\\s*>>>\\s*\\(/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), $3, $4, 0, 0, /g;" << endl; - *streamPtr.get() << tab << "# kern<<>>() syntax" << endl; - *streamPtr.get() << tab << s_k << "<<<\\s*([^,]+|[\\w:]*\\([\\w|,|:]+\\))\\s*,\\s*([^,]+|[\\w:]*\\([\\w|,|:]+\\))\\s*>>>\\s*\\(\\s*\\)/hipLaunchKernelGGL($1, $2, $3, 0, 0)/g;" << endl; - *streamPtr.get() << tab << "# kern<<>>(...) syntax" << endl; - *streamPtr.get() << tab << s_k << "<<<\\s*([^,]+|[\\w:]*\\([\\w|,|:]+\\))\\s*,\\s*([^,]+|[\\w:]*\\([\\w|,|:]+\\))\\s*>>>\\s*\\(/hipLaunchKernelGGL($1, $2, $3, 0, 0, /g;" << endl_2; + + *streamPtr.get() << tab << "# kern<...><<>>() syntax" << endl; + *streamPtr.get() << tab << s_k << "<(.+)>\\s*<<<\\s*([^,\\(\\)]+|[\\w\\s:]*\\([\\w|\\s|,|:|\\+|\\*|\\-|\\/|(?R)]+\\))\\s*,\\s*([^,\\(\\)]+|[\\w\\s:]*\\([\\w|\\s|,|:|\\+|\\*|\\-|\\/|(?R)]+\\))\\s*,\\s*([^,\\(\\)]+|[\\w\\s:]*\\([\\w|\\s|,|:|\\+|\\*|\\-|\\/|(?R)]+\\))\\s*,\\s*([^,\\(\\)]+|[\\w\\s:]*\\([\\w|\\s|,|:|\\+|\\*|\\-|\\/|(?R)]+\\))\\s*>>>\\s*\\(\\s*\\)/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), $3, $4, $5, $6)/g;" << endl; + *streamPtr.get() << tab << "# kern<...><<>>(...) syntax" << endl; + *streamPtr.get() << tab << s_k << "<(.+)>\\s*<<<\\s*([^,\\(\\)]+|[\\w\\s:]*\\([\\w|\\s|,|:|\\+|\\*|\\-|\\/|(?R)]+\\))\\s*,\\s*([^,\\(\\)]+|[\\w\\s:]*\\([\\w|\\s|,|:|\\+|\\*|\\-|\\/|(?R)]+\\))\\s*,\\s*([^,\\(\\)]+|[\\w\\s:]*\\([\\w|\\s|,|:|\\+|\\*|\\-|\\/|(?R)]+\\))\\s*,\\s*([^,\\(\\)]+|[\\w\\s:]*\\([\\w|\\s|,|:|\\+|\\*|\\-|\\/|(?R)]+\\))\\s*>>>\\s*\\(/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), $3, $4, $5, $6, /g;" << endl; + *streamPtr.get() << tab << "# kern<<>>() syntax" << endl; + *streamPtr.get() << tab << s_k << "<<<\\s*([^,\\(\\)]+|[\\w\\s:]*\\([\\w|\\s|,|:|\\+|\\*|\\-|\\/|(?R)]+\\))\\s*,\\s*([^,\\(\\)]+|[\\w\\s:]*\\([\\w|\\s|,|:|\\+|\\*|\\-|\\/|(?R)]+\\))\\s*,\\s*([^,\\(\\)]+|[\\w\\s:]*\\([\\w|\\s|,|:|\\+|\\*|\\-|\\/|(?R)]+\\))\\s*,\\s*([^,\\(\\)]+|[\\w\\s:]*\\([\\w|\\s|,|:|\\+|\\*|\\-|\\/|(?R)]+\\))\\s*>>>\\s*\\(\\s*\\)/hipLaunchKernelGGL($1, $2, $3, $4, $5)/g;" << endl; + *streamPtr.get() << tab << "# kern<<>>(...) syntax" << endl; + *streamPtr.get() << tab << s_k << "<<<\\s*([^,\\(\\)]+|[\\w\\s:]*\\([\\w|\\s|,|:|\\+|\\*|\\-|\\/|(?R)]+\\))\\s*,\\s*([^,\\(\\)]+|[\\w\\s:]*\\([\\w|\\s|,|:|\\+|\\*|\\-|\\/|(?R)]+\\))\\s*,\\s*([^,\\(\\)]+|[\\w\\s:]*\\([\\w|\\s|,|:|\\+|\\*|\\-|\\/|(?R)]+\\))\\s*,\\s*([^,\\(\\)]+|[\\w\\s:]*\\([\\w|\\s|,|:|\\+|\\*|\\-|\\/|(?R)]+\\))\\s*>>>\\s*\\(/hipLaunchKernelGGL($1, $2, $3, $4, $5, /g;" << endl_2; *streamPtr.get() << tab << "# kern<...><<>>() syntax" << endl; - *streamPtr.get() << tab << s_k << "<(.+)>\\s*<<<\\s*([^,]+|[\\w:]*\\([\\w|,|:]+\\))\\s*,\\s*([^,]+|[\\w:]*\\([\\w|,|:]+\\))\\s*,\\s*([^,]+|[\\w:]*\\([\\w|,|:]+\\))\\s*>>>\\s*\\(\\s*\\)/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), $3, $4, $5, 0)/g;" << endl; + *streamPtr.get() << tab << s_k << "<(.+)>\\s*<<<\\s*([^,\\(\\)]+|[\\w\\s:]*\\([\\w|\\s|,|:|\\+|\\*|\\-|\\/|(?R)]+\\))\\s*,\\s*([^,\\(\\)]+|[\\w\\s:]*\\([\\w|\\s|,|:|\\+|\\*|\\-|\\/|(?R)]+\\))\\s*,\\s*([^,\\(\\)]+|[\\w\\s:]*\\([\\w|\\s|,|:|\\+|\\*|\\-|\\/|(?R)]+\\))\\s*>>>\\s*\\(\\s*\\)/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), $3, $4, $5, 0)/g;" << endl; *streamPtr.get() << tab << "# kern<...><<>>(...) syntax" << endl; - *streamPtr.get() << tab << s_k << "<(.+)>\\s*<<<\\s*([^,]+|[\\w:]*\\([\\w|,|:]+\\))\\s*,\\s*([^,]+|[\\w:]*\\([\\w|,|:]+\\))\\s*,\\s*([^,]+|[\\w:]*\\([\\w|,|:]+\\))\\s*>>>\\s*\\(/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), $3, $4, $5, 0, /g;" << endl; + *streamPtr.get() << tab << s_k << "<(.+)>\\s*<<<\\s*([^,\\(\\)]+|[\\w\\s:]*\\([\\w|\\s|,|:|\\+|\\*|\\-|\\/|(?R)]+\\))\\s*,\\s*([^,\\(\\)]+|[\\w\\s:]*\\([\\w|\\s|,|:|\\+|\\*|\\-|\\/|(?R)]+\\))\\s*,\\s*([^,\\(\\)]+|[\\w\\s:]*\\([\\w|\\s|,|:|\\+|\\*|\\-|\\/|(?R)]+\\))\\s*>>>\\s*\\(/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), $3, $4, $5, 0, /g;" << endl; *streamPtr.get() << tab << "# kern<<>>() syntax" << endl; - *streamPtr.get() << tab << s_k << "<<<\\s*([^,]+|[\\w:]*\\([\\w|,|:]+\\))\\s*,\\s*([^,]+|[\\w:]*\\([\\w|,|:]+\\))\\s*,\\s*([^,]+|[\\w:]*\\([\\w|,|:]+\\))\\s*>>>\\s*\\(\\s*\\)/hipLaunchKernelGGL($1, $2, $3, $4, 0)/g;" << endl; + *streamPtr.get() << tab << s_k << "<<<\\s*([^,\\(\\)]+|[\\w\\s:]*\\([\\w|\\s|,|:|\\+|\\*|\\-|\\/|(?R)]+\\))\\s*,\\s*([^,\\(\\)]+|[\\w\\s:]*\\([\\w|\\s|,|:|\\+|\\*|\\-|\\/|(?R)]+\\))\\s*,\\s*([^,\\(\\)]+|[\\w\\s:]*\\([\\w|\\s|,|:|\\+|\\*|\\-|\\/|(?R)]+\\))\\s*>>>\\s*\\(\\s*\\)/hipLaunchKernelGGL($1, $2, $3, $4, 0)/g;" << endl; *streamPtr.get() << tab << "# kern<<>>(...) syntax" << endl; - *streamPtr.get() << tab << s_k << "<<<\\s*([^,]+|[\\w:]*\\([\\w|,|:]+\\))\\s*,\\s*([^,]+|[\\w:]*\\([\\w|,|:]+\\))\\s*,\\s*([^,]+|[\\w:]*\\([\\w|,|:]+\\))\\s*>>>\\s*\\(/hipLaunchKernelGGL($1, $2, $3, $4, 0, /g;" << endl_2; + *streamPtr.get() << tab << s_k << "<<<\\s*([^,\\(\\)]+|[\\w\\s:]*\\([\\w|\\s|,|:|\\+|\\*|\\-|\\/|(?R)]+\\))\\s*,\\s*([^,\\(\\)]+|[\\w\\s:]*\\([\\w|\\s|,|:|\\+|\\*|\\-|\\/|(?R)]+\\))\\s*,\\s*([^,\\(\\)]+|[\\w\\s:]*\\([\\w|\\s|,|:|\\+|\\*|\\-|\\/|(?R)]+\\))\\s*>>>\\s*\\(/hipLaunchKernelGGL($1, $2, $3, $4, 0, /g;" << endl_2; - *streamPtr.get() << tab << "# kern<...><<>>() syntax" << endl; - *streamPtr.get() << tab << s_k << "<(.+)>\\s*<<<\\s*([^,]+|[\\w:]*\\([\\w|,|:]+\\))\\s*,\\s*([^,]+|[\\w:]*\\([\\w|,|:]+\\))\\s*,\\s*([^,]+|[\\w:]*\\([\\w|,|:]+\\))\\s*,\\s*([^,]+|[\\w:]*\\([\\w|,|:]+\\))\\s*>>>\\s*\\(\\s*\\)/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), $3, $4, $5, $6)/g;" << endl; - *streamPtr.get() << tab << "# kern<...><<>>(...) syntax" << endl; - *streamPtr.get() << tab << s_k << "<(.+)>\\s*<<<\\s*([^,]+|[\\w:]*\\([\\w|,|:]+\\))\\s*,\\s*([^,]+|[\\w:]*\\([\\w|,|:]+\\))\\s*,\\s*([^,]+|[\\w:]*\\([\\w|,|:]+\\))\\s*,\\s*([^,]+|[\\w:]*\\([\\w|,|:]+\\))\\s*>>>\\s*\\(/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), $3, $4, $5, $6, /g;" << endl; - *streamPtr.get() << tab << "# kern<<>>() syntax" << endl; - *streamPtr.get() << tab << s_k << "<<<\\s*([^,]+|[\\w:]*\\([\\w|,|:]+\\))\\s*,\\s*([^,]+|[\\w:]*\\([\\w|,|:]+\\))\\s*,\\s*([^,]+|[\\w:]*\\([\\w|,|:]+\\))\\s*,\\s*([^,]+|[\\w:]*\\([\\w|,|:]+\\))\\s*>>>\\s*\\(\\s*\\)/hipLaunchKernelGGL($1, $2, $3, $4, $5)/g;" << endl; - *streamPtr.get() << tab << "# kern<<>>(...) syntax" << endl; - *streamPtr.get() << tab << s_k << "<<<\\s*([^,]+|[\\w:]*\\([\\w|,|:]+\\))\\s*,\\s*([^,]+|[\\w:]*\\([\\w|,|:]+\\))\\s*,\\s*([^,]+|[\\w:]*\\([\\w|,|:]+\\))\\s*,\\s*([^,]+|[\\w:]*\\([\\w|,|:]+\\))\\s*>>>\\s*\\(/hipLaunchKernelGGL($1, $2, $3, $4, $5, /g;" << endl_2; + *streamPtr.get() << tab << "# kern<...><<>>() syntax" << endl; + *streamPtr.get() << tab << s_k << "<(.+)>\\s*<<<\\s*([^,\\(\\)]+|[\\w\\s:]*\\([\\w|\\s|,|:|\\+|\\*|\\-|\\/|(?R)]+\\))\\s*,\\s*([^,\\(\\)]+|[\\w\\s:]*\\([\\w|\\s|,|:|\\+|\\*|\\-|\\/|(?R)]+\\))\\s*>>>\\s*\\(\\s*\\)/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), $3, $4, 0, 0)/g;" << endl; + *streamPtr.get() << tab << "# kern<...><<>>(...) syntax" << endl; + *streamPtr.get() << tab << s_k << "<(.+)>\\s*<<<\\s*([^,\\(\\)]+|[\\w\\s:]*\\([\\w|\\s|,|:|\\+|\\*|\\-|\\/|(?R)]+\\))\\s*,\\s*([^,\\(\\)]+|[\\w\\s:]*\\([\\w|\\s|,|:|\\+|\\*|\\-|\\/|(?R)]+\\))\\s*>>>\\s*\\(/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), $3, $4, 0, 0, /g;" << endl; + *streamPtr.get() << tab << "# kern<<>>() syntax" << endl; + *streamPtr.get() << tab << s_k << "<<<\\s*([^,\\(\\)]+|[\\w\\s:]*\\([\\w|\\s|,|:|\\+|\\*|\\-|\\/|(?R)]+\\))\\s*,\\s*([^,\\(\\)]+|[\\w\\s:]*\\([\\w|\\s|,|:|\\+|\\*|\\-|\\/|(?R)]+\\))\\s*>>>\\s*\\(\\s*\\)/hipLaunchKernelGGL($1, $2, $3, 0, 0)/g;" << endl; + *streamPtr.get() << tab << "# kern<<>>(...) syntax" << endl; + *streamPtr.get() << tab << s_k << "<<<\\s*([^,\\(\\)]+|[\\w\\s:]*\\([\\w|\\s|,|:|\\+|\\*|\\-|\\/|(?R)]+\\))\\s*,\\s*([^,\\(\\)]+|[\\w\\s:]*\\([\\w|\\s|,|:|\\+|\\*|\\-|\\/|(?R)]+\\))\\s*>>>\\s*\\(/hipLaunchKernelGGL($1, $2, $3, 0, 0, /g;" << endl_2; *streamPtr.get() << tab << "if ($k) {" << endl; *streamPtr.get() << tab_2 << "$ft{'kernel_launch'} += $k;" << endl; diff --git a/tests/unit_tests/kernel_launch/kernel_launch_syntax.cu b/tests/unit_tests/kernel_launch/kernel_launch_syntax.cu index 7d1dca4a..f048423d 100644 --- a/tests/unit_tests/kernel_launch/kernel_launch_syntax.cu +++ b/tests/unit_tests/kernel_launch/kernel_launch_syntax.cu @@ -183,12 +183,21 @@ int main(int argc, char* argv[]) { // CHECK: hipLaunchKernelGGL(nonempty, dim3(1), dim3(kDataLen), N, stream, x, y, z); nonempty<<>> (x, y, z); - // CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy_2), dim3(1), dim3(std::min(kDataLen*2+10,x)), std::min(x,y), stream, a, std::min(d_x,d_y), std::max(d_x,d_y)); - axpy_2<<<1, std::min(kDataLen*2+10,x), std::min(x,y), stream>>>(a, std::min(d_x,d_y), std::max(d_x,d_y)); - // CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy_2), dim3(1), dim3(std::min(kDataLen*2+10,x)), std::min(x,y), 0, a, std::min(d_x,d_y), std::max(d_x,d_y)); - axpy_2<<<1, std::min(kDataLen*2+10,x), std::min(x,y)>>>(a, std::min(d_x,d_y), std::max(d_x,d_y)); - // CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy_2), dim3(1), dim3(std::min(kDataLen*2+10,x)), 0, 0, a, std::min(d_x,d_y), std::max(d_x,d_y)); - axpy_2<<<1, std::min(kDataLen*2+10,x)>>>(a, std::min(d_x,d_y), std::max(d_x,d_y)); + // CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy_2), dim3(x,y,z), dim3(std::min(kDataLen*2+10,x)), std::min(x,y), stream, a, std::min(d_x,d_y), std::max(d_x,d_y)); + axpy_2<<>>(a, std::min(d_x,d_y), std::max(d_x,d_y)); + // CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy_2), dim3(x,y,z), dim3(std::min(kDataLen*2+10,x)), std::min(x,y), 0, a, std::min(d_x,d_y), std::max(d_x,d_y)); + axpy_2<<>>(a, std::min(d_x,d_y), std::max(d_x,d_y)); + // CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy_2), dim3(x,y,z), dim3(std::min(kDataLen*2+10,x)), 0, 0, a, std::min(d_x,d_y), std::max(d_x,d_y)); + axpy_2<<>>(a, std::min(d_x,d_y), std::max(d_x,d_y)); + + // CHECK: hipLaunchKernelGGL(nonempty, dim3(x,y,z), dim3(x,y,std::min(y,z)), 0, 0, x, y, z); + nonempty<<>>(x, y, z); + // CHECK: hipLaunchKernelGGL(nonempty, dim3(x,y,z), dim3(x,y,std::min(std::max(x,y),z)), 0, 0, x, y, z); + nonempty<<>>(x, y, z); + // CHECK: hipLaunchKernelGGL(nonempty, dim3(x,y,z), dim3(x,y,std::min(std::max(x,int(N)),z)), 0, 0, x, y, z); + nonempty<<>>(x, y, z); + // CHECK: hipLaunchKernelGGL(nonempty, dim3(x,y,z), dim3(x,y,std::min(std::max(x,int(N+N -x/y + y*1)),z)), 0, 0, x, y, z); + nonempty<<>>(x, y, z); // Copy output data to host. // CHECK: hipDeviceSynchronize();