[docs] Attribute details for launch kernel greater than 256 (#1319)

* Add attribute detail for kernel launch group > 256

* Add example


[ROCm/hip commit: 6d6b89d247]

projects/hip/docs/markdown/hip_porting_guide.md

@@ -34,6 +34,7 @@ and provides practical suggestions on how to port CUDA code and work through com
   * [Choosing HIP File Extensions](#choosing-hip-file-extensions)
 - [Workarounds](#workarounds)
   * [warpSize](#warpsize)
+  * [Kernel launch with group size > 256](#kernel-launch-with-group-size--256)
 - [memcpyToSymbol](#memcpytosymbol)
 - [threadfence_system](#threadfence_system)
   * [Textures and Cache Control](#textures-and-cache-control)
@@ -412,6 +413,14 @@ run hipcc when appropriate.
 ### warpSize
 Code should not assume a warp size of 32 or 64.  See [Warp Cross-Lane Functions](hip_kernel_language.md#warp-cross-lane-functions) for information on how to write portable wave-aware code.
 
+### Kernel launch with group size > 256
+Kernel code should use ``` __attribute__((amdgpu_flat_work_group_size(<min>,<max>)))```.
+
+For example:
+```
+__global__ void dot(double *a,double *b,const int n) __attribute__((amdgpu_flat_work_group_size(1, 512)))
+```
+
 ## memcpyToSymbol
 
 HIP support for hipMemcpyToSymbol is complete.  This feature allows a kernel