From 579df3902689ef60a117fa8ea26e447f05a9d891 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Wed, 14 Aug 2019 04:29:04 -0700 Subject: [PATCH] [docs] Attribute details for launch kernel greater than 256 (#1319) * Add attribute detail for kernel launch group > 256 * Add example [ROCm/hip commit: 6d6b89d2472a666c83eeaabc2cf74f6e23fa9584] --- projects/hip/docs/markdown/hip_porting_guide.md | 9 +++++++++ 1 file changed, 9 insertions(+) diff --git a/projects/hip/docs/markdown/hip_porting_guide.md b/projects/hip/docs/markdown/hip_porting_guide.md index a2b830ebf1..8f35ec776d 100644 --- a/projects/hip/docs/markdown/hip_porting_guide.md +++ b/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(,)))```. + +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