From 329e2182d6e1f05ae4fe292c164ee4599bb764b8 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 --- docs/markdown/hip_faq.md | 8 +------- include/hip/hcc_detail/grid_launch_v2.hpp | 1 - include/hip/hcc_detail/hip_runtime_api.h | 2 +- include/hip/hcc_detail/host_defines.h | 2 +- samples/0_Intro/square/square.hipref.cpp | 2 +- 5 files changed, 4 insertions(+), 11 deletions(-) diff --git a/docs/markdown/hip_faq.md b/docs/markdown/hip_faq.md index 01b05ed223..5ad8d9e8a9 100644 --- a/docs/markdown/hip_faq.md +++ b/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/include/hip/hcc_detail/grid_launch_v2.hpp b/include/hip/hcc_detail/grid_launch_v2.hpp index 02f214bd67..9ce0722496 100644 --- a/include/hip/hcc_detail/grid_launch_v2.hpp +++ b/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/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index 01fea548b8..2ff4a70802 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/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/include/hip/hcc_detail/host_defines.h b/include/hip/hcc_detail/host_defines.h index 6804ba464b..28c9268f59 100644 --- a/include/hip/hcc_detail/host_defines.h +++ b/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/samples/0_Intro/square/square.hipref.cpp b/samples/0_Intro/square/square.hipref.cpp index 118f8acf13..963ab63260 100644 --- a/samples/0_Intro/square/square.hipref.cpp +++ b/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));