From 6bb9913e8a5c9c2f32c8e52479f6d5e1ca4b138f Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Wed, 2 Oct 2019 16:01:07 +0300 Subject: [PATCH] [HIPIFY] Return to wrapping with HIP_KERNEL_NAME(...) macro of a template instantiation kernel launch [REASON] 1. hip-clang is fine with the templated kernel launch, brackets are unneeded: HIP_KERNEL_NAME(...) __VA_ARGS__ 2. HCC is not, thus: HIP_KERNEL_NAME(...) (__VA_ARGS__) [TODO] Clean-up entirely kernel name wrapping when HCC is finally obsolete. + Update perl generation, hipify-perl, and affected tests accordingly. --- bin/hipify-perl | 12 ++++++------ hipify-clang/src/CUDA2HIP_Perl.cpp | 12 ++++++------ hipify-clang/src/HipifyAction.cpp | 3 ++- tests/hipify-clang/unit_tests/device/atomics.cu | 2 +- .../libraries/cuRAND/benchmark_curand_kernel.cpp | 10 +++++----- tests/hipify-clang/unit_tests/samples/axpy.cu | 14 +++++++------- tests/hipify-clang/unit_tests/samples/square.cu | 2 +- 7 files changed, 28 insertions(+), 27 deletions(-) diff --git a/bin/hipify-perl b/bin/hipify-perl index 2a87271a60..609b0fc3f4 100755 --- a/bin/hipify-perl +++ b/bin/hipify-perl @@ -1610,32 +1610,32 @@ sub transformKernelLaunch { my $k = 0; # Handle the kern<...><<>>() 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; + $k += s/(\w+)\s*<(.+)>\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\((\s*)\)/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), dim3($3), dim3($4), $5, $6)/g; # Handle the kern<<>>() syntax with empty args: $k += s/(\w+)\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\((\s*)\)/hipLaunchKernelGGL($1, dim3($2), dim3($3), $4, $5)/g; # Handle the kern<...><<>>(...) 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; + $k += s/(\w+)\s*<(.+)>\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\(/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), dim3($3), dim3($4), $5, $6, /g; # Handle the kern<<>>(...) syntax with non-empty args: $k += s/(\w+)\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\(/hipLaunchKernelGGL($1, dim3($2), dim3($3), $4, $5, /g; # Handle the kern<...><<>>() 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; + $k += s/(\w+)\s*<(.+)>\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\((\s*)\)/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), dim3($3), dim3($4), $5, 0)/g; # Handle the kern<<>>() syntax with empty args: $k += s/(\w+)\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\((\s*)\)/hipLaunchKernelGGL($1, dim3($2), dim3($3), $4, 0)/g; # Handle the kern<...><>>(...) 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; + $k += s/(\w+)\s*<(.+)>\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\(/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), dim3($3), dim3($4), $5, 0, /g; # Handle the kern<<>>(...) syntax with non-empty args: $k += s/(\w+)\s*<<<\s*(.+)\s*,\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\(/hipLaunchKernelGGL($1, dim3($2), dim3($3), $4, 0, /g; # Handle the kern<...><<>>() 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; + $k += s/(\w+)\s*<(.+)>\s*<<<\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\((\s*)\)/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), dim3($3), dim3($4), 0, 0)/g; # Handle the kern<<>>() syntax with empty args: $k += s/(\w+)\s*<<<\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\((\s*)\)/hipLaunchKernelGGL($1, dim3($2), dim3($3), 0, 0)/g; # Handle the kern<...><<>>(...) 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; + $k += s/(\w+)\s*<(.+)>\s*<<<\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\(/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), dim3($3), dim3($4), 0, 0, /g; # Handle the kern<<>>(...) syntax with non-empty args: $k += s/(\w+)\s*<<<\s*(.+)\s*,\s*(.+)\s*>>>(\s*)\(/hipLaunchKernelGGL($1, dim3($2), dim3($3), 0, 0, /g; diff --git a/hipify-clang/src/CUDA2HIP_Perl.cpp b/hipify-clang/src/CUDA2HIP_Perl.cpp index 5d7a8790d0..f101f9b8cb 100644 --- a/hipify-clang/src/CUDA2HIP_Perl.cpp +++ b/hipify-clang/src/CUDA2HIP_Perl.cpp @@ -197,32 +197,32 @@ namespace perl { *streamPtr.get() << tab << sMy_k << std::endl << std::endl; *streamPtr.get() << tab << "# Handle the kern<...><<>>() syntax with empty args:" << std::endl; - *streamPtr.get() << tab << "$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;" << std::endl; + *streamPtr.get() << tab << "$k += s/(\\w+)\\s*<(.+)>\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\((\\s*)\\)/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), dim3($3), dim3($4), $5, $6)/g;" << std::endl; *streamPtr.get() << tab << "# Handle the kern<<>>() syntax with empty args:" << std::endl; *streamPtr.get() << tab << "$k += s/(\\w+)\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\((\\s*)\\)/hipLaunchKernelGGL($1, dim3($2), dim3($3), $4, $5)/g;" << std::endl << std::endl; *streamPtr.get() << tab << "# Handle the kern<...><<>>(...) syntax with non-empty args:" << std::endl; - *streamPtr.get() << tab << "$k += s/(\\w+)\\s*<(.+)>\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\(/hipLaunchKernelGGL(($1<$2>), dim3($3), dim3($4), $5, $6, /g;" << std::endl; + *streamPtr.get() << tab << "$k += s/(\\w+)\\s*<(.+)>\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\(/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), dim3($3), dim3($4), $5, $6, /g;" << std::endl; *streamPtr.get() << tab << "# Handle the kern<<>>(...) syntax with non-empty args:" << std::endl; *streamPtr.get() << tab << "$k += s/(\\w+)\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\(/hipLaunchKernelGGL($1, dim3($2), dim3($3), $4, $5, /g;" << std::endl << std::endl; *streamPtr.get() << tab << "# Handle the kern<...><<>>() syntax with empty args:" << std::endl; - *streamPtr.get() << tab << "$k += s/(\\w+)\\s*<(.+)>\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\((\\s*)\\)/hipLaunchKernelGGL(($1<$2>), dim3($3), dim3($4), $5, 0)/g;" << std::endl; + *streamPtr.get() << tab << "$k += s/(\\w+)\\s*<(.+)>\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\((\\s*)\\)/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), dim3($3), dim3($4), $5, 0)/g;" << std::endl; *streamPtr.get() << tab << "# Handle the kern<<>>() syntax with empty args:" << std::endl; *streamPtr.get() << tab << "$k += s/(\\w+)\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\((\\s*)\\)/hipLaunchKernelGGL($1, dim3($2), dim3($3), $4, 0)/g;" << std::endl << std::endl; *streamPtr.get() << tab << "# Handle the kern<...><>>(...) syntax with non-empty args:" << std::endl; - *streamPtr.get() << tab << "$k += s/(\\w+)\\s*<(.+)>\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\(/hipLaunchKernelGGL(($1<$2>), dim3($3), dim3($4), $5, 0, /g;" << std::endl; + *streamPtr.get() << tab << "$k += s/(\\w+)\\s*<(.+)>\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\(/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), dim3($3), dim3($4), $5, 0, /g;" << std::endl; *streamPtr.get() << tab << "# Handle the kern<<>>(...) syntax with non-empty args:" << std::endl; *streamPtr.get() << tab << "$k += s/(\\w+)\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\(/hipLaunchKernelGGL($1, dim3($2), dim3($3), $4, 0, /g;" << std::endl << std::endl; *streamPtr.get() << tab << "# Handle the kern<...><<>>() syntax with empty args:" << std::endl; - *streamPtr.get() << tab << "$k += s/(\\w+)\\s*<(.+)>\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\((\\s*)\\)/hipLaunchKernelGGL(($1<$2>), dim3($3), dim3($4), 0, 0)/g;" << std::endl; + *streamPtr.get() << tab << "$k += s/(\\w+)\\s*<(.+)>\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\((\\s*)\\)/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), dim3($3), dim3($4), 0, 0)/g;" << std::endl; *streamPtr.get() << tab << "# Handle the kern<<>>() syntax with empty args:" << std::endl; *streamPtr.get() << tab << "$k += s/(\\w+)\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\((\\s*)\\)/hipLaunchKernelGGL($1, dim3($2), dim3($3), 0, 0)/g;" << std::endl << std::endl; *streamPtr.get() << tab << "# Handle the kern<...><<>>(...) syntax with non-empty args:" << std::endl; - *streamPtr.get() << tab << "$k += s/(\\w+)\\s*<(.+)>\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\(/hipLaunchKernelGGL(($1<$2>), dim3($3), dim3($4), 0, 0, /g;" << std::endl; + *streamPtr.get() << tab << "$k += s/(\\w+)\\s*<(.+)>\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\(/hipLaunchKernelGGL(HIP_KERNEL_NAME($1<$2>), dim3($3), dim3($4), 0, 0, /g;" << std::endl; *streamPtr.get() << tab << "# Handle the kern<<>>(...) syntax with non-empty args:" << std::endl; *streamPtr.get() << tab << "$k += s/(\\w+)\\s*<<<\\s*(.+)\\s*,\\s*(.+)\\s*>>>(\\s*)\\(/hipLaunchKernelGGL($1, dim3($2), dim3($3), 0, 0, /g;" << std::endl << std::endl; diff --git a/hipify-clang/src/HipifyAction.cpp b/hipify-clang/src/HipifyAction.cpp index cb0a5eedc9..5c5663985a 100644 --- a/hipify-clang/src/HipifyAction.cpp +++ b/hipify-clang/src/HipifyAction.cpp @@ -38,6 +38,7 @@ namespace mat = clang::ast_matchers; const std::string sHIP_DYNAMIC_SHARED = "HIP_DYNAMIC_SHARED"; std::string sHIP_SYMBOL = "HIP_SYMBOL"; +std::string sHIP_KERNEL_NAME = "HIP_KERNEL_NAME"; std::string s_reinterpret_cast = "reinterpret_cast"; const std::string sHipLaunchKernelGGL = "hipLaunchKernelGGL("; const std::string sDim3 = "dim3("; @@ -348,7 +349,7 @@ bool HipifyAction::cudaLaunchKernel(const clang::ast_matchers::MatchFinder::Matc clang::LangOptions DefaultLangOptions; clang::SourceManager* SM = Result.SourceManager; OS << sHipLaunchKernelGGL; - if (caleeDecl->isTemplateInstantiation()) OS << "("; + if (caleeDecl->isTemplateInstantiation()) OS << sHIP_KERNEL_NAME << "("; OS << readSourceText(*SM, calleeExpr->getSourceRange()); if (caleeDecl->isTemplateInstantiation()) OS << ")"; OS << ", "; diff --git a/tests/hipify-clang/unit_tests/device/atomics.cu b/tests/hipify-clang/unit_tests/device/atomics.cu index d1c1fad401..1afd1ab541 100644 --- a/tests/hipify-clang/unit_tests/device/atomics.cu +++ b/tests/hipify-clang/unit_tests/device/atomics.cu @@ -257,7 +257,7 @@ void runTest() { // CHECK: hipMemcpy(dOData, hOData, memSize, hipMemcpyHostToDevice); cudaMemcpy(dOData, hOData, memSize, cudaMemcpyHostToDevice); // Execute the kernel - // CHECK: hipLaunchKernelGGL((testKernel), dim3(numBlocks), dim3(numThreads), 0, 0, dOData); + // CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(testKernel), dim3(numBlocks), dim3(numThreads), 0, 0, dOData); testKernel<<>>(dOData); // Copy result from device to host // CHECK: hipMemcpy(hOData, dOData, memSize, hipMemcpyDeviceToHost); diff --git a/tests/hipify-clang/unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp b/tests/hipify-clang/unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp index ece384f04b..bff9b77cad 100644 --- a/tests/hipify-clang/unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp +++ b/tests/hipify-clang/unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp @@ -121,7 +121,7 @@ struct runner const size_t states_size = blocks * threads; // CHECK: CUDA_CALL(hipMalloc((void **)&states, states_size * sizeof(GeneratorState))); CUDA_CALL(cudaMalloc((void **)&states, states_size * sizeof(GeneratorState))); - // CHECK: hipLaunchKernelGGL((init_kernel), dim3(blocks), dim3(threads), 0, 0, states, seed, offset); + // CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(init_kernel), dim3(blocks), dim3(threads), 0, 0, states, seed, offset); init_kernel<<>>(states, seed, offset); // CHECK: CUDA_CALL(hipPeekAtLastError()); // CHECK: CUDA_CALL(hipDeviceSynchronize()); @@ -142,7 +142,7 @@ struct runner const GenerateFunc& generate_func, const Extra extra) { - // CHECK: hipLaunchKernelGGL((generate_kernel), dim3(blocks), dim3(threads), 0, 0, states, data, size, generate_func, extra); + // CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(generate_kernel), dim3(blocks), dim3(threads), 0, 0, states, data, size, generate_func, extra); generate_kernel<<>>(states, data, size, generate_func, extra); } }; @@ -223,7 +223,7 @@ struct runner const GenerateFunc& generate_func, const Extra extra) { - // CHECK: hipLaunchKernelGGL((generate_kernel), dim3(std::min((size_t)200, blocks)), dim3(256), 0, 0, states, data, size, generate_func, extra); + // CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(generate_kernel), dim3(std::min((size_t)200, blocks)), dim3(256), 0, 0, states, data, size, generate_func, extra); generate_kernel<<>>(states, data, size, generate_func, extra); } }; @@ -304,7 +304,7 @@ struct runner CUDA_CALL(cudaMemcpy(directions, h_directions, size, cudaMemcpyHostToDevice)); const size_t blocks_x = next_power2((blocks + dimensions - 1) / dimensions); - // CHECK: hipLaunchKernelGGL((init_kernel), dim3(blocks_x, dimensions), dim3(threads), 0, 0, states, directions, offset); + // CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(init_kernel), dim3(blocks_x, dimensions), dim3(threads), 0, 0, states, directions, offset); init_kernel<<>>(states, directions, offset); // CHECK: CUDA_CALL(hipPeekAtLastError()); // CHECK: CUDA_CALL(hipDeviceSynchronize()); @@ -329,7 +329,7 @@ struct runner const Extra extra) { const size_t blocks_x = next_power2((blocks + dimensions - 1) / dimensions); - // CHECK: hipLaunchKernelGGL((generate_kernel), dim3(blocks_x, dimensions), dim3(threads), 0, 0, states, data, size / dimensions, generate_func, extra); + // CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(generate_kernel), dim3(blocks_x, dimensions), dim3(threads), 0, 0, states, data, size / dimensions, generate_func, extra); generate_kernel<<>>(states, data, size / dimensions, generate_func, extra); } }; diff --git a/tests/hipify-clang/unit_tests/samples/axpy.cu b/tests/hipify-clang/unit_tests/samples/axpy.cu index 549dc442aa..c09c372967 100644 --- a/tests/hipify-clang/unit_tests/samples/axpy.cu +++ b/tests/hipify-clang/unit_tests/samples/axpy.cu @@ -10,7 +10,7 @@ #define KERNEL_CALL_AS_MACRO axpy<<<1, kDataLen>>> #define KERNEL_NAME_MACRO axpy -// CHECK: #define COMPLETE_LAUNCH hipLaunchKernelGGL((axpy), dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y) +// CHECK: #define COMPLETE_LAUNCH hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy), dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y) #define COMPLETE_LAUNCH axpy<<<1, kDataLen>>>(a, device_x, device_y) @@ -48,22 +48,22 @@ int main(int argc, char* argv[]) { cudaMemcpy(device_x, host_x, kDataLen * sizeof(float), cudaMemcpyHostToDevice); // Launch the kernel in numerous different strange ways to exercise the prerocessor. - // CHECK: hipLaunchKernelGGL((axpy), dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y); + // CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy), dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y); axpy<<<1, kDataLen>>>(a, device_x, device_y); - // CHECK: hipLaunchKernelGGL((axpy), dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y); + // CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy), dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y); axpy<<<1, kDataLen>>>(a, device_x, device_y); - // CHECK: hipLaunchKernelGGL((axpy), dim3(1), dim3(kDataLen), 0, 0, a, TOKEN_PASTE(device, _x), device_y); + // CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy), dim3(1), dim3(kDataLen), 0, 0, a, TOKEN_PASTE(device, _x), device_y); axpy<<<1, kDataLen>>>(a, TOKEN_PASTE(device, _x), device_y); - // CHECK: hipLaunchKernelGGL((axpy), dim3(1), dim3(kDataLen), 0, 0, ARG_LIST_AS_MACRO); + // CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy), dim3(1), dim3(kDataLen), 0, 0, ARG_LIST_AS_MACRO); axpy<<<1, kDataLen>>>(ARG_LIST_AS_MACRO); - // CHECK: hipLaunchKernelGGL((KERNEL_NAME_MACRO), dim3(1), dim3(kDataLen), 0, 0, ARG_LIST_AS_MACRO); + // CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(KERNEL_NAME_MACRO), dim3(1), dim3(kDataLen), 0, 0, ARG_LIST_AS_MACRO); KERNEL_NAME_MACRO<<<1, kDataLen>>>(ARG_LIST_AS_MACRO); - // CHECK: hipLaunchKernelGGL((axpy), dim3(1), dim3(kDataLen), 0, 0, ARG_LIST_AS_MACRO); + // CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(axpy), 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); diff --git a/tests/hipify-clang/unit_tests/samples/square.cu b/tests/hipify-clang/unit_tests/samples/square.cu index 690d99f848..0c2ba94083 100644 --- a/tests/hipify-clang/unit_tests/samples/square.cu +++ b/tests/hipify-clang/unit_tests/samples/square.cu @@ -94,7 +94,7 @@ int main(int argc, char *argv[]) const unsigned threadsPerBlock = 256; printf ("info: launch 'vector_square' kernel\n"); - // CHECK: hipLaunchKernelGGL((vector_square), dim3(blocks), dim3(threadsPerBlock), 0, 0, C_d, A_d, N); + // CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(vector_square), dim3(blocks), dim3(threadsPerBlock), 0, 0, C_d, A_d, N); vector_square <<>> (C_d, A_d, N); printf ("info: copy Device2Host\n");