Merge pull request #1387 from emankov/master

[HIPIFY][perl][#259] Fix

[ROCm/clr commit: 9af6190ec9]
This commit is contained in:
Evgeny Mankov
2019-09-04 16:11:16 +03:00
committato da GitHub
2 ha cambiato i file con 25 aggiunte e 13 eliminazioni
+14 -13
Vedi File
@@ -403,24 +403,25 @@ while (@ARGV) {
#--------
# CUDA Launch Syntax
# Note these only work if launch is on a single line.
# Handle the <<numBlocks, blockDim>>> syntax:
{
# match uses ? for <.*> which will be unitialized if this is not present in launch syntax.
no warnings qw/uninitialized/;
my $k = 0;
my $kernelName;
# Handle the <<numBlocks, blockDim, sharedSize, stream>>> syntax:
$k += s/(\w+)\s*(<.*>)?\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>([\s*\\]*)\(/hipLaunchKernelGGL(($1$2), dim3($3), dim3($4), $5, $6, /g;
$kernelName = $1 if $k;
# Handle the <<numBlocks, blockDim, sharedSize>>> syntax:
$k += s/(\w+)\s*(<.*>)?\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>([\s*\\]*)\(/hipLaunchKernelGGL(($1$2), dim3($3), dim3($4), $5, 0, /g;
$kernelName = $1 if $k;
# Handle the <<numBlocks, blockDim>>> syntax:
$k += s/(\w+)\s*(<.*>)?\s*<<<\s*(.+)\s*,\s*(.+)\s*>>>([\s\\]*)\(/hipLaunchKernelGGL(($1$2), dim3($3), dim3($4), 0, 0, /g;
$kernelName = $1 if $k;
$ft{'kern'} += $k;
# Handle the <<numBlocks, blockDim, sharedSize, stream>>> syntax with empty args:
$k += s/(\w+)\s*(<.*>)?\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\((\s*)\)/hipLaunchKernelGGL($1$2, dim3($3), dim3($4), $5, $6)/g;
# Handle the <<numBlocks, blockDim, sharedSize, stream>>> syntax with non-empty args:
$k += s/(\w+)\s*(<.*>)?\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\(/hipLaunchKernelGGL($1$2, dim3($3), dim3($4), $5, $6, /g;
# Handle the <<numBlocks, blockDim, sharedSize>>> syntax with empty args:
$k += s/(\w+)\s*(<.*>)?\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\((\s*)\)/hipLaunchKernelGGL($1$2, dim3($3), dim3($4), $5, 0)/g;
# Handle the <<numBlocks, blockDim, sharedSize>>> syntax with non-empty args:
$k += s/(\w+)\s*(<.*>)?\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\(/hipLaunchKernelGGL($1$2, dim3($3), dim3($4), $5, 0, /g;
# Handle the <<numBlocks, blockDim>>> syntax with empty args:
$k += s/(\w+)\s*(<.*>)?\s*<<<\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\((\s*)\)/hipLaunchKernelGGL($1$2, dim3($3), dim3($4), 0, 0)/g;
# Handle the <<numBlocks, blockDim>>> syntax with non-empty args:
$k += s/(\w+)\s*(<.*>)?\s*<<<\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\(/hipLaunchKernelGGL($1$2, dim3($3), dim3($4), 0, 0, /g;
if ($k) {
$Tkernels{$kernelName} ++;
$ft{'kern'} += $k;
$Tkernels{$1} ++;
}
}
unless ($no_translate_textures) {
@@ -19,6 +19,8 @@ __global__ void axpy(T a, T *x, T *y) {
y[threadIdx.x] = a * x[threadIdx.x];
}
__global__ void empty() {
}
int main(int argc, char* argv[]) {
const int kDataLen = 4;
@@ -64,6 +66,15 @@ int main(int argc, char* argv[]) {
// CHECK: hipLaunchKernelGGL(axpy<float>, dim3(1), dim3(kDataLen), 0, 0, ARG_LIST_AS_MACRO);
KERNEL_CALL_AS_MACRO(ARG_LIST_AS_MACRO);
// CHECK: hipLaunchKernelGGL(empty, dim3(1), dim3(kDataLen), 0, 0);
empty<<<1, kDataLen>>> ( );
// CHECK: hipLaunchKernelGGL(empty, dim3(1), dim3(kDataLen), 0, 0);
empty<<<1, kDataLen, 0>>>();
// CHECK: hipLaunchKernelGGL(empty, dim3(1), dim3(kDataLen), 0, 0);
empty<<<1, kDataLen, 0, 0>>>();
// CHECK: COMPLETE_LAUNCH;
COMPLETE_LAUNCH;