From db4bedeafc3596d5284a69e8fafc2d6cb8ae4923 Mon Sep 17 00:00:00 2001 From: "Sun, Peng" Date: Mon, 20 Mar 2017 17:03:21 -0500 Subject: [PATCH] revert workaround for square sample and update doc on GGL Change-Id: I731c68ca4111e7dc2e45bef51c4cad2c23fc81f8 [ROCm/clr commit: 14f8c7cf74b560657e5401899d3c2322ac7dd7bf] --- projects/clr/hipamd/docs/markdown/hip_faq.md | 8 +------- .../clr/hipamd/include/hip/hcc_detail/grid_launch_v2.hpp | 1 - .../clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h | 2 +- projects/clr/hipamd/include/hip/hcc_detail/host_defines.h | 2 +- .../clr/hipamd/samples/0_Intro/square/square.hipref.cpp | 2 +- 5 files changed, 4 insertions(+), 11 deletions(-) diff --git a/projects/clr/hipamd/docs/markdown/hip_faq.md b/projects/clr/hipamd/docs/markdown/hip_faq.md index 01b05ed223..5ad8d9e8a9 100644 --- a/projects/clr/hipamd/docs/markdown/hip_faq.md +++ b/projects/clr/hipamd/docs/markdown/hip_faq.md @@ -243,10 +243,4 @@ $HIP/include/hip/hcc_detail/hip_runtime_api.h $HIP/include/hip/hcc_detail/host_defines.h Or pass "-DGENERIC_GRID_LAUNCH=1" to hipcc at application compilation time. -There are some limitation/assumptions of GGL implementation right now: -1. GGL was only tested with Ubuntu16.04, assuming HCC and HIP only link against libstdc++ but not libc++. -2. GGL currently requires templated kernel fucntions passed to hipLaunchKernel to be specialized. e.g.: - -``` -hipLaunchKernel(vector_square, dim3(blocks), dim3(threadsPerBlock), 0, nullptr, C_d, A_d, N); -``` +GGL was only tested with Ubuntu16.04, assuming HCC and HIP only link against libstdc++ but not libc++. diff --git a/projects/clr/hipamd/include/hip/hcc_detail/grid_launch_v2.hpp b/projects/clr/hipamd/include/hip/hcc_detail/grid_launch_v2.hpp index 02f214bd67..9ce0722496 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/grid_launch_v2.hpp +++ b/projects/clr/hipamd/include/hip/hcc_detail/grid_launch_v2.hpp @@ -323,7 +323,6 @@ namespace glo_tests make_lambda_wrapper(kernel_name, __VA_ARGS__), \ ##__VA_ARGS__); \ } -#warning "GGL hipLaunchKernel Reached" #define hipLaunchKernel( \ kernel_name, \ num_blocks, \ diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h index 01fea548b8..2ff4a70802 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime_api.h @@ -32,7 +32,7 @@ THE SOFTWARE. #include #ifndef GENERIC_GRID_LAUNCH -#define GENERIC_GRID_LAUNCH 0 +#define GENERIC_GRID_LAUNCH 0 #endif #include diff --git a/projects/clr/hipamd/include/hip/hcc_detail/host_defines.h b/projects/clr/hipamd/include/hip/hcc_detail/host_defines.h index 6804ba464b..28c9268f59 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/host_defines.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/host_defines.h @@ -32,7 +32,7 @@ THE SOFTWARE. // Add guard to Generic Grid Launch method #ifndef GENERIC_GRID_LAUNCH -#define GENERIC_GRID_LAUNCH 0 +#define GENERIC_GRID_LAUNCH 0 #endif #ifdef __HCC__ diff --git a/projects/clr/hipamd/samples/0_Intro/square/square.hipref.cpp b/projects/clr/hipamd/samples/0_Intro/square/square.hipref.cpp index 118f8acf13..963ab63260 100644 --- a/projects/clr/hipamd/samples/0_Intro/square/square.hipref.cpp +++ b/projects/clr/hipamd/samples/0_Intro/square/square.hipref.cpp @@ -83,7 +83,7 @@ int main(int argc, char *argv[]) const unsigned threadsPerBlock = 256; printf ("info: launch 'vector_square' kernel\n"); - hipLaunchKernel(vector_square, dim3(blocks), dim3(threadsPerBlock), 0, nullptr, C_d, A_d, N); + hipLaunchKernel(vector_square, dim3(blocks), dim3(threadsPerBlock), 0, nullptr, C_d, A_d, N); printf ("info: copy Device2Host\n"); CHECK ( hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost));