Added sample with __gfx*__ macro used

Change-Id: I80a474947ea1b755082dad5329d1aff3ee78840c


[ROCm/hip-tests commit: f40e7d71cc]

projects/hip-tests/samples/2_Cookbook/14_gpu_arch/CMakeLists.txt

@@ -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)

projects/hip-tests/samples/2_Cookbook/14_gpu_arch/Makefile

@@ -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)

projects/hip-tests/samples/2_Cookbook/14_gpu_arch/gpuarch.cpp

@@ -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;
+}