From 3bd73d6280709a6c02eeef080e33ee53699cae69 Mon Sep 17 00:00:00 2001 From: "Sun, Peng" Date: Mon, 20 Mar 2017 15:44:28 -0500 Subject: [PATCH] Add document for switching to GGL in hip_faq.md Change-Id: I83d9fd3e76d21ab572949c3a446ac3898acb3ded --- hipamd/docs/markdown/hip_faq.md | 15 +++++++++++++++ 1 file changed, 15 insertions(+) diff --git a/hipamd/docs/markdown/hip_faq.md b/hipamd/docs/markdown/hip_faq.md index d7235c4ebc..1543241f6d 100644 --- a/hipamd/docs/markdown/hip_faq.md +++ b/hipamd/docs/markdown/hip_faq.md @@ -235,3 +235,18 @@ Unlike CUDA, in HCC, for functions defined in the header files, the keyword of " Thus, if failed to define "static" keyword, you might see a lot of "symbol multiply defined!" errors at compilation. The workaround is to explicitly add the keyword of "static" before any functions that were defined as "__forceinline__". + +### How do I enable HIP Generic Grid Launch option? +Generic Grid Launch(GGL) provide a second choice for kernel launch. +To enable it, either change the default value of GENERIC_GRID_LAUNCH to 1 in the following to header files and rebuild HIP: +$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); +```