Skip to content

Commit

Permalink
Merge pull request #303 from emankov/HIPIFY
Browse files Browse the repository at this point in the history 
[HIPIFY][perl][#299] Step 1: Implement N-level nesting support in Kernel Launch stmnt
  • Loading branch information
@emankov
emankov authored Mar 19, 2021
2 parents 1310672 + 2f126eb commit 4085755
Show file tree
Hide file tree
Showing 3 changed files with 56 additions and 46 deletions.
40 changes: 20 additions & 20 deletions bin/hipify-perl
View file
Original file line number Diff line number Diff line change
Expand Up @@ -2991,32 +2991,32 @@ sub transformKernelLaunch {
no warnings qw/uninitialized/;
my $k = 0;

# kern<...><<<Dg, Db>>>() 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<...><<<Dg, Db>>>(...) 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<<<Dg, Db>>>() syntax
$k += s/([:|\w]+)\s*<<<\s*([^,]+|[\w:]*\([\w|,|:]+\))\s*,\s*([^,]+|[\w:]*\([\w|,|:]+\))\s*>>>\s*\(\s*\)/hipLaunchKernelGGL($1, $2, $3, 0, 0)/g;
# kern<<<Dg, Db>>>(...) syntax
$k += s/([:|\w]+)\s*<<<\s*([^,]+|[\w:]*\([\w|,|:]+\))\s*,\s*([^,]+|[\w:]*\([\w|,|:]+\))\s*>>>\s*\(/hipLaunchKernelGGL($1, $2, $3, 0, 0, /g;
# kern<...><<<Dg, Db, Ns, S>>>() 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<...><<<Dg, Db, Ns, S>>>(...) 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<<<Dg, Db, Ns, S>>>() 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<<<Dg, Db, Ns, S>>>(...) 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<...><<<Dg, Db, Ns>>>() 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<...><<<Dg, Db, Ns>>>(...) 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<<<Dg, Db, Ns>>>() 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<<<Dg, Db, Ns>>>(...) 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<...><<<Dg, Db, Ns, S>>>() 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<...><<<Dg, Db, Ns, S>>>(...) 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<<<Dg, Db, Ns, S>>>() 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<<<Dg, Db, Ns, S>>>(...) 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<...><<<Dg, Db>>>() 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<...><<<Dg, Db>>>(...) 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<<<Dg, Db>>>() 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<<<Dg, Db>>>(...) 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;
Expand Down
41 changes: 21 additions & 20 deletions src/CUDA2HIP_Perl.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -310,32 +310,33 @@ namespace perl {

string s_k = "$k += s/([:|\\w]+)\\s*";

*streamPtr.get() << tab << "# kern<...><<<Dg, Db>>>() 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<...><<<Dg, Db>>>(...) 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<<<Dg, Db>>>() 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<<<Dg, Db>>>(...) 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<...><<<Dg, Db, Ns, S>>>() 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<...><<<Dg, Db, Ns, S>>>(...) 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<<<Dg, Db, Ns, S>>>() 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<<<Dg, Db, Ns, S>>>(...) 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<...><<<Dg, Db, Ns>>>() 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<...><<<Dg, Db, Ns>>>(...) 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<<<Dg, Db, Ns>>>() 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<<<Dg, Db, Ns>>>(...) 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<...><<<Dg, Db, Ns, S>>>() 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<...><<<Dg, Db, Ns, S>>>(...) 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<<<Dg, Db, Ns, S>>>() 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<<<Dg, Db, Ns, S>>>(...) 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<...><<<Dg, Db>>>() 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<...><<<Dg, Db>>>(...) 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<<<Dg, Db>>>() 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<<<Dg, Db>>>(...) 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;
Expand Down
21 changes: 15 additions & 6 deletions tests/unit_tests/kernel_launch/kernel_launch_syntax.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -183,12 +183,21 @@ int main(int argc, char* argv[]) {
// CHECK: hipLaunchKernelGGL(nonempty, dim3(1), dim3(kDataLen), N, stream, x, y, z);
nonempty<<<dim3(1), dim3(kDataLen), N, stream>>> (x, y, z);

// CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy_2<float,double>), 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<float,double><<<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<float,double>), 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<float,double><<<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<float,double>), 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<float,double><<<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<float,double>), 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<float,double><<<dim3(x,y,z), 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<float,double>), 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<float,double><<<dim3(x,y,z), 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<float,double>), 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<float,double><<<dim3(x,y,z), std::min(kDataLen*2+10,x)>>>(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<<<dim3(x,y,z), dim3(x,y,std::min(y,z))>>>(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<<<dim3(x,y,z), dim3(x,y,std::min(std::max(x,y),z))>>>(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<<<dim3(x,y,z), dim3(x,y,std::min(std::max(x,int(N)),z))>>>(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<<<dim3(x,y,z), dim3(x,y,std::min(std::max(x,int(N+N -x/y + y*1)),z))>>>(x, y, z);

// Copy output data to host.
// CHECK: hipDeviceSynchronize();
Expand Down

0 comments on commit 4085755

Please sign in to comment.