From cedc7bb2f91b0a9636c659398f78eeb7f4b27504 Mon Sep 17 00:00:00 2001 From: Sarbojit Sarkar Date: Mon, 12 Oct 2020 02:44:58 -0400 Subject: [PATCH] Added sample with __gfx*__ macro used Change-Id: I80a474947ea1b755082dad5329d1aff3ee78840c [ROCm/hip-tests commit: f40e7d71cc977ad90ce95b6d99562d6710880dc9] --- .../2_Cookbook/14_gpu_arch/CMakeLists.txt | 20 ++++ .../samples/2_Cookbook/14_gpu_arch/Makefile | 20 ++++ .../2_Cookbook/14_gpu_arch/gpuarch.cpp | 98 +++++++++++++++++++ 3 files changed, 138 insertions(+) create mode 100644 projects/hip-tests/samples/2_Cookbook/14_gpu_arch/CMakeLists.txt create mode 100644 projects/hip-tests/samples/2_Cookbook/14_gpu_arch/Makefile create mode 100644 projects/hip-tests/samples/2_Cookbook/14_gpu_arch/gpuarch.cpp diff --git a/projects/hip-tests/samples/2_Cookbook/14_gpu_arch/CMakeLists.txt b/projects/hip-tests/samples/2_Cookbook/14_gpu_arch/CMakeLists.txt new file mode 100644 index 0000000000..066dda439a --- /dev/null +++ b/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) diff --git a/projects/hip-tests/samples/2_Cookbook/14_gpu_arch/Makefile b/projects/hip-tests/samples/2_Cookbook/14_gpu_arch/Makefile new file mode 100644 index 0000000000..a0bdcfe1ad --- /dev/null +++ b/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) diff --git a/projects/hip-tests/samples/2_Cookbook/14_gpu_arch/gpuarch.cpp b/projects/hip-tests/samples/2_Cookbook/14_gpu_arch/gpuarch.cpp new file mode 100644 index 0000000000..8d629b3015 --- /dev/null +++ b/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 + +#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(malloc(NBytes)); + hOutput = static_cast(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; +} \ No newline at end of file