Added sample with __gfx*__ macro used

Change-Id: I80a474947ea1b755082dad5329d1aff3ee78840c [ROCm/hip commit: a45f6f7ff3]
2020-10-12 02:44:58 -04:00
Commit 665440e499
@@ -52,7 +52,7 @@
 - [In-Line Assembly](#in-line-assembly)
 - [C++ Support](#c-support)
 - [Kernel Compilation](#kernel-compilation)
-
+- [GFX Arch specific kernel](#gfx-arch-specific-kernel)
 <!-- tocstop -->

 ## Introduction
@@ -803,3 +803,5 @@ The file format for binary is `.co` which means Code Object. The following comma

 Note that one important fact to remember when using binary code objects is that the number of arguments to the kernel are different on HCC and NVCC path. Refer to the sample in samples/0_Intro/module_api for differences in the arguments to be passed to the kernel.

+## gfx-arch-specific-kernel
+Clang defined '__gfx*__' macros can be used to execute gfx arch specific codes inside the kernel. Refer to the sample 14_gpu_arch in samples/2_Cookbook.
@@ -0,0 +1,20 @@
+project(gpuarch)
+
+cmake_minimum_required(VERSION 3.10)
+
+# Search for rocm in common locations
+list(APPEND CMAKE_PREFIX_PATH /opt/rocm/hip /opt/rocm)
+
+# Find hip
+find_package(hip)
+
+# Set compiler and linker
+set(CMAKE_CXX_COMPILER ${HIP_HIPCC_EXECUTABLE})
+set(CMAKE_CXX_LINKER   ${HIP_HIPCC_EXECUTABLE})
+set(CMAKE_BUILD_TYPE Release)
+
+# Create the excutable
+add_executable(gpuarch gpuarch.cpp)
+
+# Link with HIP
+target_link_libraries(gpuarch hip::host)
@@ -0,0 +1,20 @@
+HIP_PATH?= $(wildcard /opt/rocm/hip)
+ifeq (,$(HIP_PATH))
+	HIP_PATH=../../..
+endif
+HIPCC=$(HIP_PATH)/bin/hipcc
+
+EXE=./gpuarch
+
+.PHONY: test
+
+all: test
+
+$(EXE): gpuarch.cpp
+	$(HIPCC) $^ -o $@
+
+test: $(EXE)
+	$(EXE)
+
+clean:
+	rm -f *.o $(EXE)
@@ -0,0 +1,98 @@
+/*
+Copyright (c) 2020-Present Advanced Micro Devices, Inc. All rights reserved.
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
+IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT.  IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include "hip/hip_runtime.h"
+#include <iostream>
+
+#define THREADS_PER_BLOCK   64
+#define BLOCKS_PER_GRID     4
+#define SIZE                (BLOCKS_PER_GRID * THREADS_PER_BLOCK)
+#define NOT_SUPPORTED       -99  // dummy number indicates unsupported operation
+
+#define HIP_STATUS_CHECK(status)                                                                   \
+  if (status != hipSuccess) {                                                                      \
+    std::cout << "Got Status: " << status << " at Line: " << __LINE__ << std::endl;                \
+    exit(0);                                                                                       \
+  }
+
+// Using __gfx*__ macro one can have GPU architecture specific code flow
+// For example: If below kernel runs on gfx908 it will increment 'in' by 'value' and store into
+// 'out'
+//              but it will update with "NOT_SUPPORTED" for any other gfx archs.
+__global__ void incrementKernel(int32_t* in, int32_t* out, int32_t value, size_t buffSize) {
+  int index = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
+  if (index < buffSize) {
+#if defined(__gfx908__)
+    out[index] = in[index] + value;
+#else
+    out[index] = NOT_SUPPORTED;
+#endif
+  }
+}
+
+int main() {
+  int32_t incrementValue = 10;
+  // Host pointers
+  int32_t* hInput = nullptr;
+  int32_t* hOutput = nullptr;
+  // Device pointers
+  int32_t* dInput = nullptr;
+  int32_t* dOutput = nullptr;
+
+  size_t NBytes = SIZE * sizeof(int32_t);
+
+  hInput = static_cast<int32_t*>(malloc(NBytes));
+  hOutput = static_cast<int32_t*>(malloc(NBytes));
+
+  HIP_STATUS_CHECK(hipMalloc(&dInput, NBytes));
+  HIP_STATUS_CHECK(hipMalloc(&dOutput, NBytes));
+
+  // Initialize host input/output buffers
+  for (int i = 0; i < SIZE; ++i) {
+    hInput[i] = i;
+    hOutput[i] = 0;
+  }
+
+  // Initialize device input buffer
+  HIP_STATUS_CHECK(hipMemcpy(dInput, hInput, NBytes, hipMemcpyHostToDevice));
+
+  // Launch kernel
+  hipLaunchKernelGGL(incrementKernel, dim3(BLOCKS_PER_GRID), dim3(THREADS_PER_BLOCK), 0, 0, dInput,
+                     dOutput, incrementValue, SIZE);
+
+  // Copy result back to host buffer
+  HIP_STATUS_CHECK(hipMemcpy(hOutput, dOutput, NBytes, hipMemcpyDeviceToHost));
+
+  bool flag = true;
+  // verify data
+  for (int i = 0; i < SIZE; ++i) {
+    if (hOutput[i] != NOT_SUPPORTED && hOutput[i] != (hInput[i] + incrementValue)) {
+      std::cout << "Error : Data mismatch found";
+      exit(0);
+    } else if (hOutput[i] == NOT_SUPPORTED) {
+      flag &= false;
+    }
+  }
+  if (flag == false) {
+    std::cout << "Error: Kernel is supported for gfx908 architecture\n";
+  } else {
+    std::cout << "success\n";
+  }
+  return 0;
+}