From 9ab3b01ba650871a1aaa0b61f4c67dec48f54698 Mon Sep 17 00:00:00 2001 From: Sandeep Kumar Date: Wed, 14 Dec 2016 15:49:40 +0530 Subject: [PATCH 1/8] Fixes in Makefile of couple of samples - modified Makefile for hipblas_saxpy to replaced hcblas.so with hipblas.so as part of HCSWAP-100 - Resolved missing separator issue in peer2peer cookbook Makefile Change-Id: I678fea267eee1481f02da09379339ed78d3f95f2 [ROCm/hip-tests commit: 5ae4e8bd6783b63a4d6974c97e1d37b3d41b31fb] --- .../samples/2_Cookbook/8_peer2peer/Makefile | 13 +++++++------ .../samples/7_Advanced/hipblas_saxpy/Makefile | 2 +- 2 files changed, 8 insertions(+), 7 deletions(-) diff --git a/projects/hip-tests/samples/2_Cookbook/8_peer2peer/Makefile b/projects/hip-tests/samples/2_Cookbook/8_peer2peer/Makefile index 5cb7473921..0bf9e6f93e 100644 --- a/projects/hip-tests/samples/2_Cookbook/8_peer2peer/Makefile +++ b/projects/hip-tests/samples/2_Cookbook/8_peer2peer/Makefile @@ -1,6 +1,6 @@ HIP_PATH?= $(wildcard /opt/rocm/hip) ifeq (,$(HIP_PATH)) - HIP_PATH=../../.. + HIP_PATH=../../.. endif HIPCC=$(HIP_PATH)/bin/hipcc @@ -22,14 +22,15 @@ CXX=$(HIPCC) $(EXECUTABLE): $(OBJECTS) - $(HIPCC) $(OBJECTS) -o $@ + $(HIPCC) $(OBJECTS) -o $@ + test: $(EXECUTABLE) - $(EXECUTABLE) + $(EXECUTABLE) clean: -rm -f $(EXECUTABLE) -rm -f $(OBJECTS) -rm -f $(HIP_PATH)/src/*.o + rm -f $(EXECUTABLE) + rm -f $(OBJECTS) + rm -f $(HIP_PATH)/src/*.o diff --git a/projects/hip-tests/samples/7_Advanced/hipblas_saxpy/Makefile b/projects/hip-tests/samples/7_Advanced/hipblas_saxpy/Makefile index ed88be2dd0..8586e75d25 100644 --- a/projects/hip-tests/samples/7_Advanced/hipblas_saxpy/Makefile +++ b/projects/hip-tests/samples/7_Advanced/hipblas_saxpy/Makefile @@ -12,7 +12,7 @@ endif ifeq (${HIP_PLATFORM}, hcc) HCBLAS_ROOT?= $(wildcard /opt/rocm/hcblas) HIPCC_FLAGS += -stdlib=libc++ -I$(HCBLAS_ROOT)/include - LIBS = -L$(HCBLAS_ROOT)/lib -lhcblas + LIBS = -L$(HCBLAS_ROOT)/lib -lhipblas -rpath $(HIP_PATH)/lib endif From b6dc34732916236e3966d01841b33861cd5aa768 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Fri, 7 Apr 2017 15:40:09 +0530 Subject: [PATCH 2/8] Merge branch 'amd-develop' into amd-master Change-Id: I53d5a8916d769c4f0fe60d2ee3b240551da80b4f (cherry picked from commit 01c523f6c9fc5a66d5f6e60d5efe78c29c9fe317) [ROCm/hip-tests commit: 9fcc03e2b64dde04a82935ce2fd559a262498108] --- projects/hip-tests/samples/0_Intro/bit_extract/Makefile | 4 ---- projects/hip-tests/samples/1_Utils/hipCommander/Makefile | 3 --- .../hip-tests/samples/1_Utils/hipCommander/hipCommander.cpp | 1 - 3 files changed, 8 deletions(-) diff --git a/projects/hip-tests/samples/0_Intro/bit_extract/Makefile b/projects/hip-tests/samples/0_Intro/bit_extract/Makefile index 78f6a2faa8..08bca6e642 100644 --- a/projects/hip-tests/samples/0_Intro/bit_extract/Makefile +++ b/projects/hip-tests/samples/0_Intro/bit_extract/Makefile @@ -11,10 +11,6 @@ HIPCC=$(HIP_PATH)/bin/hipcc ifeq (${HIP_PLATFORM}, nvcc) HIPCC_FLAGS = -gencode=arch=compute_20,code=sm_20 endif -ifeq (${HIP_PLATFORM}, hcc) - HIPCC_FLAGS = -stdlib=libc++ -endif - EXE=bit_extract diff --git a/projects/hip-tests/samples/1_Utils/hipCommander/Makefile b/projects/hip-tests/samples/1_Utils/hipCommander/Makefile index e770c636a4..a411763b7f 100644 --- a/projects/hip-tests/samples/1_Utils/hipCommander/Makefile +++ b/projects/hip-tests/samples/1_Utils/hipCommander/Makefile @@ -10,9 +10,6 @@ OPT=-O3 CXXFLAGS = $(OPT) --std=c++11 HIP_PLATFORM=$(shell $(HIP_PATH)/bin/hipconfig --platform) -ifeq (${HIP_PLATFORM}, hcc) - CXXFLAGS += " -stdlib=libc++" -endif CODE_OBJECTS=nullkernel.hsaco diff --git a/projects/hip-tests/samples/1_Utils/hipCommander/hipCommander.cpp b/projects/hip-tests/samples/1_Utils/hipCommander/hipCommander.cpp index 0add1ce3e3..4b93180b18 100644 --- a/projects/hip-tests/samples/1_Utils/hipCommander/hipCommander.cpp +++ b/projects/hip-tests/samples/1_Utils/hipCommander/hipCommander.cpp @@ -11,7 +11,6 @@ #include #include #include -#include #endif #include From 182d5261b3b72821f0254f4c916e46cf1f525a86 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Mon, 24 Apr 2017 08:50:43 +0530 Subject: [PATCH 3/8] Merge branch 'amd-develop' into amd-master Change-Id: I312fb9d1181733ef5160d1e993e2ae57ced0f6b3 (cherry picked from commit 88fb807af081f31314031d4549e98d9b621cfc41) [ROCm/hip-tests commit: b8fd2f159a4003d204f60f24186e3627afef6c3a] --- projects/hip-tests/samples/0_Intro/square/Makefile | 1 + projects/hip-tests/samples/0_Intro/square/square.hipref.cpp | 2 +- 2 files changed, 2 insertions(+), 1 deletion(-) diff --git a/projects/hip-tests/samples/0_Intro/square/Makefile b/projects/hip-tests/samples/0_Intro/square/Makefile index 1e8cdba080..aa48cc5864 100644 --- a/projects/hip-tests/samples/0_Intro/square/Makefile +++ b/projects/hip-tests/samples/0_Intro/square/Makefile @@ -15,5 +15,6 @@ square.hip.out: square.hipref.cpp + clean: rm -f *.o *.out diff --git a/projects/hip-tests/samples/0_Intro/square/square.hipref.cpp b/projects/hip-tests/samples/0_Intro/square/square.hipref.cpp index 963ab63260..e694bfb8a4 100644 --- a/projects/hip-tests/samples/0_Intro/square/square.hipref.cpp +++ b/projects/hip-tests/samples/0_Intro/square/square.hipref.cpp @@ -83,7 +83,7 @@ int main(int argc, char *argv[]) const unsigned threadsPerBlock = 256; printf ("info: launch 'vector_square' kernel\n"); - hipLaunchKernel(vector_square, dim3(blocks), dim3(threadsPerBlock), 0, nullptr, C_d, A_d, N); + hipLaunchKernel(vector_square, dim3(blocks), dim3(threadsPerBlock), 0, 0, C_d, A_d, N); printf ("info: copy Device2Host\n"); CHECK ( hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost)); From 7a20b9afe154e82ab304c770ef700534a5bd95cd Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Thu, 29 Jun 2017 12:01:40 -0500 Subject: [PATCH 4/8] automate gcnarch detection Change-Id: Ibbad22db136f7f5e2be84c82e9169298a144cc77 [ROCm/hip-tests commit: 98905a7272b9fbc6b0c9712cd47577962b8a17c1] --- projects/hip-tests/samples/1_Utils/hipInfo/hipInfo.cpp | 1 + 1 file changed, 1 insertion(+) diff --git a/projects/hip-tests/samples/1_Utils/hipInfo/hipInfo.cpp b/projects/hip-tests/samples/1_Utils/hipInfo/hipInfo.cpp index cf4660eae7..0401745efd 100644 --- a/projects/hip-tests/samples/1_Utils/hipInfo/hipInfo.cpp +++ b/projects/hip-tests/samples/1_Utils/hipInfo/hipInfo.cpp @@ -129,6 +129,7 @@ void printDeviceProp (int deviceId) cout << setw(w1) << "arch.hasSurfaceFuncs: " << props.arch.hasSurfaceFuncs << endl; cout << setw(w1) << "arch.has3dGrid: " << props.arch.has3dGrid << endl; cout << setw(w1) << "arch.hasDynamicParallelism: " << props.arch.hasDynamicParallelism << endl; + cout << setw(w1) << "gcnArch: " << props.gcnArch << endl; int deviceCnt; hipGetDeviceCount(&deviceCnt); From 6442d0652f4886e986df1fb93000984f4a9395ec Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Wed, 5 Jul 2017 11:44:44 +0530 Subject: [PATCH 5/8] GPUOpen-ProfessionalCompute-Tools -> ROCm-Developer-Tools Change-Id: I9f5b29dd1097385acecb0c672770d8adca2fdcf7 [ROCm/hip-tests commit: 8252ae785b3772476b0a32c748cd6671d6140e5e] --- .../2_Cookbook/0_MatrixTranspose/Readme.md | 18 +++++++++--------- .../samples/2_Cookbook/10_inline_asm/Readme.md | 18 +++++++++--------- .../samples/2_Cookbook/1_hipEvent/Readme.md | 18 +++++++++--------- .../samples/2_Cookbook/2_Profiler/Readme.md | 16 ++++++++-------- .../2_Cookbook/3_shared_memory/Readme.md | 18 +++++++++--------- .../samples/2_Cookbook/4_shfl/Readme.md | 18 +++++++++--------- .../samples/2_Cookbook/5_2dshfl/Readme.md | 18 +++++++++--------- .../2_Cookbook/6_dynamic_shared/Readme.md | 18 +++++++++--------- .../samples/2_Cookbook/7_streams/Readme.md | 18 +++++++++--------- .../samples/2_Cookbook/9_unroll/Readme.md | 18 +++++++++--------- 10 files changed, 89 insertions(+), 89 deletions(-) diff --git a/projects/hip-tests/samples/2_Cookbook/0_MatrixTranspose/Readme.md b/projects/hip-tests/samples/2_Cookbook/0_MatrixTranspose/Readme.md index 5e9483b595..ab5dbdc958 100644 --- a/projects/hip-tests/samples/2_Cookbook/0_MatrixTranspose/Readme.md +++ b/projects/hip-tests/samples/2_Cookbook/0_MatrixTranspose/Readme.md @@ -7,7 +7,7 @@ This tutorial shows how to get write simple HIP application. We will write the s HIP is a C++ runtime API and kernel language that allows developers to create portable applications that can run on AMD and other GPU’s. Our goal was to rise above the lowest-common-denominator paths and deliver a solution that allows you, the developer, to use essential hardware features and maximize your application’s performance on GPU hardware. ## Requirement: -For hardware requirement and software installation [Installation](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/INSTALL.md) +For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md) ## prerequiste knowledge: @@ -90,11 +90,11 @@ Use the make command and execute it using ./exe Use hipcc to build the application, which is using hcc on AMD and nvcc on nvidia. ## More Info: -- [HIP FAQ](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_faq.md) -- [HIP Kernel Language](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_kernel_language.md) -- [HIP Runtime API (Doxygen)](http://gpuopen-professionalcompute-tools.github.io/HIP) -- [HIP Porting Guide](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_porting_guide.md) -- [HIP Terminology](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenL) -- [hipify-clang](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/hipify-clang/README.md) -- [Developer/CONTRIBUTING Info](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/CONTRIBUTING.md) -- [Release Notes](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/RELEASE.md) +- [HIP FAQ](https://github.com/ROCm-Developer-Tools/HIP/docs/markdown/hip_faq.md) +- [HIP Kernel Language](https://github.com/ROCm-Developer-Tools/HIP/docs/markdown/hip_kernel_language.md) +- [HIP Runtime API (Doxygen)](http://rocm-developer-tools.github.io/HIP) +- [HIP Porting Guide](https://github.com/ROCm-Developer-Tools/HIP/docs/markdown/hip_porting_guide.md) +- [HIP Terminology](https://github.com/ROCm-Developer-Tools/HIP/docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenL) +- [hipify-clang](https://github.com/ROCm-Developer-Tools/HIP/hipify-clang/README.md) +- [Developer/CONTRIBUTING Info](https://github.com/ROCm-Developer-Tools/HIP/CONTRIBUTING.md) +- [Release Notes](https://github.com/ROCm-Developer-Tools/HIP/RELEASE.md) diff --git a/projects/hip-tests/samples/2_Cookbook/10_inline_asm/Readme.md b/projects/hip-tests/samples/2_Cookbook/10_inline_asm/Readme.md index 8c98547220..0e64fe9c6e 100644 --- a/projects/hip-tests/samples/2_Cookbook/10_inline_asm/Readme.md +++ b/projects/hip-tests/samples/2_Cookbook/10_inline_asm/Readme.md @@ -15,7 +15,7 @@ For more information: [User Guide for AMDGPU Back-end](llvm.org/docs/AMDGPUUsage.html) ## Requirement: -For hardware requirement and software installation [Installation](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/INSTALL.md) +For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md) ## prerequiste knowledge: @@ -37,11 +37,11 @@ Use hipcc to build the application, which is using hcc on AMD and nvcc on nvidia ## More Info: -- [HIP FAQ](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_faq.md) -- [HIP Kernel Language](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_kernel_language.md) -- [HIP Runtime API (Doxygen)](http://gpuopen-professionalcompute-tools.github.io/HIP) -- [HIP Porting Guide](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_porting_guide.md) -- [HIP Terminology](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenL) -- [clang-hipify](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/clang-hipify/README.md) -- [Developer/CONTRIBUTING Info](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/CONTRIBUTING.md) -- [Release Notes](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/RELEASE.md) +- [HIP FAQ](https://github.com/ROCm-Developer-Tools/HIP/docs/markdown/hip_faq.md) +- [HIP Kernel Language](https://github.com/ROCm-Developer-Tools/HIP/docs/markdown/hip_kernel_language.md) +- [HIP Runtime API (Doxygen)](http://rocm-developer-tools.github.io/HIP) +- [HIP Porting Guide](https://github.com/ROCm-Developer-Tools/HIP/docs/markdown/hip_porting_guide.md) +- [HIP Terminology](https://github.com/ROCm-Developer-Tools/HIP/docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenL) +- [clang-hipify](https://github.com/ROCm-Developer-Tools/HIP/clang-hipify/README.md) +- [Developer/CONTRIBUTING Info](https://github.com/ROCm-Developer-Tools/HIP/CONTRIBUTING.md) +- [Release Notes](https://github.com/ROCm-Developer-Tools/HIP/RELEASE.md) diff --git a/projects/hip-tests/samples/2_Cookbook/1_hipEvent/Readme.md b/projects/hip-tests/samples/2_Cookbook/1_hipEvent/Readme.md index e3ec8ad780..ea4f3a67e9 100644 --- a/projects/hip-tests/samples/2_Cookbook/1_hipEvent/Readme.md +++ b/projects/hip-tests/samples/2_Cookbook/1_hipEvent/Readme.md @@ -7,7 +7,7 @@ This tutorial is follow-up of the previous one where we learn how to write our f Memory transfer and kernel execution are the most important parameter in parallel computing (specially HPC and machine learning). Memory bottlenecks is the main problem why we are not able to get the highest performance, therefore obtaining the memory transfer timing and kernel execution timing plays key role in application optimization. ## Requirement: -For hardware requirement and software installation [Installation](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/INSTALL.md) +For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md) ## prerequiste knowledge: @@ -64,11 +64,11 @@ Use the make command and execute it using ./exe Use hipcc to build the application, which is using hcc on AMD and nvcc on nvidia. ## More Info: -- [HIP FAQ](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_faq.md) -- [HIP Kernel Language](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_kernel_language.md) -- [HIP Runtime API (Doxygen)](http://gpuopen-professionalcompute-tools.github.io/HIP) -- [HIP Porting Guide](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_porting_guide.md) -- [HIP Terminology](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenL) -- [hipify-clang](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/hipify-clang/README.md) -- [Developer/CONTRIBUTING Info](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/CONTRIBUTING.md) -- [Release Notes](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/RELEASE.md) +- [HIP FAQ](https://github.com/ROCm-Developer-Tools/HIP/docs/markdown/hip_faq.md) +- [HIP Kernel Language](https://github.com/ROCm-Developer-Tools/HIP/docs/markdown/hip_kernel_language.md) +- [HIP Runtime API (Doxygen)](http://rocm-developer-tools.github.io/HIP) +- [HIP Porting Guide](https://github.com/ROCm-Developer-Tools/HIP/docs/markdown/hip_porting_guide.md) +- [HIP Terminology](https://github.com/ROCm-Developer-Tools/HIP/docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenL) +- [hipify-clang](https://github.com/ROCm-Developer-Tools/HIP/hipify-clang/README.md) +- [Developer/CONTRIBUTING Info](https://github.com/ROCm-Developer-Tools/HIP/CONTRIBUTING.md) +- [Release Notes](https://github.com/ROCm-Developer-Tools/HIP/RELEASE.md) diff --git a/projects/hip-tests/samples/2_Cookbook/2_Profiler/Readme.md b/projects/hip-tests/samples/2_Cookbook/2_Profiler/Readme.md index 92a8be228e..4059e42193 100644 --- a/projects/hip-tests/samples/2_Cookbook/2_Profiler/Readme.md +++ b/projects/hip-tests/samples/2_Cookbook/2_Profiler/Readme.md @@ -37,11 +37,11 @@ You can also print the HIP function strings to stderr using HIP_TRACE_API enviro Note this trace mode uses colors. "less -r" can handle raw control characters and will display the debug output in proper colors. ## More Info: -- [HIP FAQ](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_faq.md) -- [HIP Kernel Language](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_kernel_language.md) -- [HIP Runtime API (Doxygen)](http://gpuopen-professionalcompute-tools.github.io/HIP) -- [HIP Porting Guide](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_porting_guide.md) -- [HIP Terminology](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenL) -- [hipify-clang](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/hipify-clang/README.md) -- [Developer/CONTRIBUTING Info](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/CONTRIBUTING.md) -- [Release Notes](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/RELEASE.md) +- [HIP FAQ](https://github.com/ROCm-Developer-Tools/HIP/docs/markdown/hip_faq.md) +- [HIP Kernel Language](https://github.com/ROCm-Developer-Tools/HIP/docs/markdown/hip_kernel_language.md) +- [HIP Runtime API (Doxygen)](http://rocm-developer-tools.github.io/HIP) +- [HIP Porting Guide](https://github.com/ROCm-Developer-Tools/HIP/docs/markdown/hip_porting_guide.md) +- [HIP Terminology](https://github.com/ROCm-Developer-Tools/HIP/docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenL) +- [hipify-clang](https://github.com/ROCm-Developer-Tools/HIP/hipify-clang/README.md) +- [Developer/CONTRIBUTING Info](https://github.com/ROCm-Developer-Tools/HIP/CONTRIBUTING.md) +- [Release Notes](https://github.com/ROCm-Developer-Tools/HIP/RELEASE.md) diff --git a/projects/hip-tests/samples/2_Cookbook/3_shared_memory/Readme.md b/projects/hip-tests/samples/2_Cookbook/3_shared_memory/Readme.md index 6b9393397c..8b9e102ec9 100644 --- a/projects/hip-tests/samples/2_Cookbook/3_shared_memory/Readme.md +++ b/projects/hip-tests/samples/2_Cookbook/3_shared_memory/Readme.md @@ -7,7 +7,7 @@ Earlier we learned how to write our first hip program, in which we compute Matri As we mentioned earlier that Memory bottlenecks is the main problem why we are not able to get the highest performance, therefore minimizing the latency for memory access plays prominent role in application optimization. In this tutorial, we'll learn how to use static shared memory and will explain the dynamic one latter. ## Requirement: -For hardware requirement and software installation [Installation](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/INSTALL.md) +For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md) ## prerequiste knowledge: @@ -32,11 +32,11 @@ Use the make command and execute it using ./exe Use hipcc to build the application, which is using hcc on AMD and nvcc on nvidia. ## More Info: -- [HIP FAQ](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_faq.md) -- [HIP Kernel Language](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_kernel_language.md) -- [HIP Runtime API (Doxygen)](http://gpuopen-professionalcompute-tools.github.io/HIP) -- [HIP Porting Guide](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_porting_guide.md) -- [HIP Terminology](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenL) -- [clang-hipify](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/clang-hipify/README.md) -- [Developer/CONTRIBUTING Info](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/CONTRIBUTING.md) -- [Release Notes](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/RELEASE.md) +- [HIP FAQ](https://github.com/ROCm-Developer-Tools/HIP/docs/markdown/hip_faq.md) +- [HIP Kernel Language](https://github.com/ROCm-Developer-Tools/HIP/docs/markdown/hip_kernel_language.md) +- [HIP Runtime API (Doxygen)](http://rocm-developer-tools.github.io/HIP) +- [HIP Porting Guide](https://github.com/ROCm-Developer-Tools/HIP/docs/markdown/hip_porting_guide.md) +- [HIP Terminology](https://github.com/ROCm-Developer-Tools/HIP/docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenL) +- [clang-hipify](https://github.com/ROCm-Developer-Tools/HIP/clang-hipify/README.md) +- [Developer/CONTRIBUTING Info](https://github.com/ROCm-Developer-Tools/HIP/CONTRIBUTING.md) +- [Release Notes](https://github.com/ROCm-Developer-Tools/HIP/RELEASE.md) diff --git a/projects/hip-tests/samples/2_Cookbook/4_shfl/Readme.md b/projects/hip-tests/samples/2_Cookbook/4_shfl/Readme.md index da62901851..923d2f3837 100644 --- a/projects/hip-tests/samples/2_Cookbook/4_shfl/Readme.md +++ b/projects/hip-tests/samples/2_Cookbook/4_shfl/Readme.md @@ -15,7 +15,7 @@ Let's talk about Warp first. The kernel code is executed in groups of fixed numb ` float __shfl_xor (float var, int laneMask, int width=warpSize); ` ## Requirement: -For hardware requirement and software installation [Installation](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/INSTALL.md) +For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md) ## prerequiste knowledge: @@ -41,11 +41,11 @@ Use hipcc to build the application, which is using hcc on AMD and nvcc on nvidia please make sure you have a 3.0 or higher compute capable device in order to use warp shfl operations and add `-gencode arch=compute=30, code=sm_30` nvcc flag in the Makefile while using this application. ## More Info: -- [HIP FAQ](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_faq.md) -- [HIP Kernel Language](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_kernel_language.md) -- [HIP Runtime API (Doxygen)](http://gpuopen-professionalcompute-tools.github.io/HIP) -- [HIP Porting Guide](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_porting_guide.md) -- [HIP Terminology](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenL) -- [clang-hipify](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/clang-hipify/README.md) -- [Developer/CONTRIBUTING Info](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/CONTRIBUTING.md) -- [Release Notes](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/RELEASE.md) +- [HIP FAQ](https://github.com/ROCm-Developer-Tools/HIP/docs/markdown/hip_faq.md) +- [HIP Kernel Language](https://github.com/ROCm-Developer-Tools/HIP/docs/markdown/hip_kernel_language.md) +- [HIP Runtime API (Doxygen)](http://rocm-developer-tools.github.io/HIP) +- [HIP Porting Guide](https://github.com/ROCm-Developer-Tools/HIP/docs/markdown/hip_porting_guide.md) +- [HIP Terminology](https://github.com/ROCm-Developer-Tools/HIP/docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenL) +- [clang-hipify](https://github.com/ROCm-Developer-Tools/HIP/clang-hipify/README.md) +- [Developer/CONTRIBUTING Info](https://github.com/ROCm-Developer-Tools/HIP/CONTRIBUTING.md) +- [Release Notes](https://github.com/ROCm-Developer-Tools/HIP/RELEASE.md) diff --git a/projects/hip-tests/samples/2_Cookbook/5_2dshfl/Readme.md b/projects/hip-tests/samples/2_Cookbook/5_2dshfl/Readme.md index fba114152a..8efff49d8e 100644 --- a/projects/hip-tests/samples/2_Cookbook/5_2dshfl/Readme.md +++ b/projects/hip-tests/samples/2_Cookbook/5_2dshfl/Readme.md @@ -15,7 +15,7 @@ Let's talk about Warp first. The kernel code is executed in groups of fixed numb ` float __shfl_xor (float var, int laneMask, int width=warpSize); ` ## Requirement: -For hardware requirement and software installation [Installation](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/INSTALL.md) +For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md) ## prerequiste knowledge: @@ -41,11 +41,11 @@ Use hipcc to build the application, which is using hcc on AMD and nvcc on nvidia please make sure you have a 3.0 or higher compute capable device in order to use warp shfl operations and add `-gencode arch=compute=30, code=sm_30` nvcc flag in the Makefile while using this application. ## More Info: -- [HIP FAQ](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_faq.md) -- [HIP Kernel Language](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_kernel_language.md) -- [HIP Runtime API (Doxygen)](http://gpuopen-professionalcompute-tools.github.io/HIP) -- [HIP Porting Guide](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_porting_guide.md) -- [HIP Terminology](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenL) -- [clang-hipify](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/clang-hipify/README.md) -- [Developer/CONTRIBUTING Info](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/CONTRIBUTING.md) -- [Release Notes](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/RELEASE.md) +- [HIP FAQ](https://github.com/ROCm-Developer-Tools/HIP/docs/markdown/hip_faq.md) +- [HIP Kernel Language](https://github.com/ROCm-Developer-Tools/HIP/docs/markdown/hip_kernel_language.md) +- [HIP Runtime API (Doxygen)](http://rocm-developer-tools.github.io/HIP) +- [HIP Porting Guide](https://github.com/ROCm-Developer-Tools/HIP/docs/markdown/hip_porting_guide.md) +- [HIP Terminology](https://github.com/ROCm-Developer-Tools/HIP/docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenL) +- [clang-hipify](https://github.com/ROCm-Developer-Tools/HIP/clang-hipify/README.md) +- [Developer/CONTRIBUTING Info](https://github.com/ROCm-Developer-Tools/HIP/CONTRIBUTING.md) +- [Release Notes](https://github.com/ROCm-Developer-Tools/HIP/RELEASE.md) diff --git a/projects/hip-tests/samples/2_Cookbook/6_dynamic_shared/Readme.md b/projects/hip-tests/samples/2_Cookbook/6_dynamic_shared/Readme.md index a10fd56a95..15ea299a9c 100644 --- a/projects/hip-tests/samples/2_Cookbook/6_dynamic_shared/Readme.md +++ b/projects/hip-tests/samples/2_Cookbook/6_dynamic_shared/Readme.md @@ -7,7 +7,7 @@ Earlier we learned how to use static shared memory. In this tutorial, we'll expl As we mentioned earlier that Memory bottlenecks is the main problem why we are not able to get the highest performance, therefore minimizing the latency for memory access plays prominent role in application optimization. In this tutorial, we'll learn how to use dynamic shared memory. ## Requirement: -For hardware requirement and software installation [Installation](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/INSTALL.md) +For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md) ## prerequiste knowledge: @@ -37,11 +37,11 @@ Use the make command and execute it using ./exe Use hipcc to build the application, which is using hcc on AMD and nvcc on nvidia. ## More Info: -- [HIP FAQ](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_faq.md) -- [HIP Kernel Language](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_kernel_language.md) -- [HIP Runtime API (Doxygen)](http://gpuopen-professionalcompute-tools.github.io/HIP) -- [HIP Porting Guide](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_porting_guide.md) -- [HIP Terminology](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenL) -- [clang-hipify](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/clang-hipify/README.md) -- [Developer/CONTRIBUTING Info](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/CONTRIBUTING.md) -- [Release Notes](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/RELEASE.md) +- [HIP FAQ](https://github.com/ROCm-Developer-Tools/HIP/docs/markdown/hip_faq.md) +- [HIP Kernel Language](https://github.com/ROCm-Developer-Tools/HIP/docs/markdown/hip_kernel_language.md) +- [HIP Runtime API (Doxygen)](http://rocm-developer-tools.github.io/HIP) +- [HIP Porting Guide](https://github.com/ROCm-Developer-Tools/HIP/docs/markdown/hip_porting_guide.md) +- [HIP Terminology](https://github.com/ROCm-Developer-Tools/HIP/docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenL) +- [clang-hipify](https://github.com/ROCm-Developer-Tools/HIP/clang-hipify/README.md) +- [Developer/CONTRIBUTING Info](https://github.com/ROCm-Developer-Tools/HIP/CONTRIBUTING.md) +- [Release Notes](https://github.com/ROCm-Developer-Tools/HIP/RELEASE.md) diff --git a/projects/hip-tests/samples/2_Cookbook/7_streams/Readme.md b/projects/hip-tests/samples/2_Cookbook/7_streams/Readme.md index a75149925e..ca295d3f49 100644 --- a/projects/hip-tests/samples/2_Cookbook/7_streams/Readme.md +++ b/projects/hip-tests/samples/2_Cookbook/7_streams/Readme.md @@ -7,7 +7,7 @@ In all Earlier tutorial we used single stream, In this tutorial, we'll explain h The various instances of kernel to be executed on device in exact launch order defined by Host are called streams. We can launch multiple streams on a single device. We will learn how to learn two streams which can we scaled with ease. ## Requirement: -For hardware requirement and software installation [Installation](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/INSTALL.md) +For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md) ## prerequiste knowledge: @@ -47,11 +47,11 @@ Use the make command and execute it using ./exe Use hipcc to build the application, which is using hcc on AMD and nvcc on nvidia. ## More Info: -- [HIP FAQ](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_faq.md) -- [HIP Kernel Language](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_kernel_language.md) -- [HIP Runtime API (Doxygen)](http://gpuopen-professionalcompute-tools.github.io/HIP) -- [HIP Porting Guide](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_porting_guide.md) -- [HIP Terminology](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenL) -- [clang-hipify](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/clang-hipify/README.md) -- [Developer/CONTRIBUTING Info](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/CONTRIBUTING.md) -- [Release Notes](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/RELEASE.md) +- [HIP FAQ](https://github.com/ROCm-Developer-Tools/HIP/docs/markdown/hip_faq.md) +- [HIP Kernel Language](https://github.com/ROCm-Developer-Tools/HIP/docs/markdown/hip_kernel_language.md) +- [HIP Runtime API (Doxygen)](http://rocm-developer-tools.github.io/HIP) +- [HIP Porting Guide](https://github.com/ROCm-Developer-Tools/HIP/docs/markdown/hip_porting_guide.md) +- [HIP Terminology](https://github.com/ROCm-Developer-Tools/HIP/docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenL) +- [clang-hipify](https://github.com/ROCm-Developer-Tools/HIP/clang-hipify/README.md) +- [Developer/CONTRIBUTING Info](https://github.com/ROCm-Developer-Tools/HIP/CONTRIBUTING.md) +- [Release Notes](https://github.com/ROCm-Developer-Tools/HIP/RELEASE.md) diff --git a/projects/hip-tests/samples/2_Cookbook/9_unroll/Readme.md b/projects/hip-tests/samples/2_Cookbook/9_unroll/Readme.md index 3c2635c0eb..194eb9c7ff 100644 --- a/projects/hip-tests/samples/2_Cookbook/9_unroll/Readme.md +++ b/projects/hip-tests/samples/2_Cookbook/9_unroll/Readme.md @@ -8,7 +8,7 @@ Loop unrolling optimization hints can be specified with #pragma unroll and #prag Specifying #pragma unroll without a parameter directs the loop unroller to attempt to fully unroll the loop if the trip count is known at compile time and attempt to partially unroll the loop if the trip count is not known at compile time. ## Requirement: -For hardware requirement and software installation [Installation](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/INSTALL.md) +For hardware requirement and software installation [Installation](https://github.com/ROCm-Developer-Tools/HIP/INSTALL.md) ## prerequiste knowledge: @@ -38,11 +38,11 @@ Use hipcc to build the application, which is using hcc on AMD and nvcc on nvidia please make sure you have a 3.0 or higher compute capable device in order to use warp shfl operations and add `-gencode arch=compute=30, code=sm_30` nvcc flag in the Makefile while using this application. ## More Info: -- [HIP FAQ](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_faq.md) -- [HIP Kernel Language](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_kernel_language.md) -- [HIP Runtime API (Doxygen)](http://gpuopen-professionalcompute-tools.github.io/HIP) -- [HIP Porting Guide](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_porting_guide.md) -- [HIP Terminology](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenL) -- [clang-hipify](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/clang-hipify/README.md) -- [Developer/CONTRIBUTING Info](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/CONTRIBUTING.md) -- [Release Notes](https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/RELEASE.md) +- [HIP FAQ](https://github.com/ROCm-Developer-Tools/HIP/docs/markdown/hip_faq.md) +- [HIP Kernel Language](https://github.com/ROCm-Developer-Tools/HIP/docs/markdown/hip_kernel_language.md) +- [HIP Runtime API (Doxygen)](http://rocm-developer-tools.github.io/HIP) +- [HIP Porting Guide](https://github.com/ROCm-Developer-Tools/HIP/docs/markdown/hip_porting_guide.md) +- [HIP Terminology](https://github.com/ROCm-Developer-Tools/HIP/docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenL) +- [clang-hipify](https://github.com/ROCm-Developer-Tools/HIP/clang-hipify/README.md) +- [Developer/CONTRIBUTING Info](https://github.com/ROCm-Developer-Tools/HIP/CONTRIBUTING.md) +- [Release Notes](https://github.com/ROCm-Developer-Tools/HIP/RELEASE.md) From 4255c3cb44894c04cec9b69e0380a749e396e3be Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Fri, 21 Jul 2017 15:50:12 -0500 Subject: [PATCH 6/8] fixed device selection during compilation to use rocm_agent_enumerator 1. Changed hipcc to use rocm_agent_enumerator 2. Changed square sample test to use device variable [ROCm/hip-tests commit: 8e3e104313314ffb18a53fca1c4faadbcb4dd8ab] --- projects/hip-tests/samples/0_Intro/square/square.hipref.cpp | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/projects/hip-tests/samples/0_Intro/square/square.hipref.cpp b/projects/hip-tests/samples/0_Intro/square/square.hipref.cpp index e694bfb8a4..167cb135f4 100644 --- a/projects/hip-tests/samples/0_Intro/square/square.hipref.cpp +++ b/projects/hip-tests/samples/0_Intro/square/square.hipref.cpp @@ -54,9 +54,10 @@ int main(int argc, char *argv[]) float *A_h, *C_h; size_t N = 1000000; size_t Nbytes = N * sizeof(float); - + static int device = 0; + CHECK(hipSetDevice(device)); hipDeviceProp_t props; - CHECK(hipGetDeviceProperties(&props, 0/*deviceID*/)); + CHECK(hipGetDeviceProperties(&props, device/*deviceID*/)); printf ("info: running on device %s\n", props.name); #ifdef __HIP_PLATFORM_HCC__ printf ("info: architecture on AMD GPU device is: %d\n",props.gcnArch); From a1d0dda3010fffc5047b2178707cbbfe3f00e262 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Wed, 26 Jul 2017 18:52:53 -0500 Subject: [PATCH 7/8] Refactor dispatch latency test and fix several bugs. [ROCm/hip-tests commit: 6ac55d2b34c1fcebbbbb5270192071ab5cf1ebeb] --- .../hipDispatchLatency/ResultDatabase.cpp | 29 ++- .../hipDispatchLatency/hipDispatchLatency.cpp | 219 +++++++++--------- 2 files changed, 136 insertions(+), 112 deletions(-) diff --git a/projects/hip-tests/samples/1_Utils/hipDispatchLatency/ResultDatabase.cpp b/projects/hip-tests/samples/1_Utils/hipDispatchLatency/ResultDatabase.cpp index d207154e39..f6f2fab709 100644 --- a/projects/hip-tests/samples/1_Utils/hipDispatchLatency/ResultDatabase.cpp +++ b/projects/hip-tests/samples/1_Utils/hipDispatchLatency/ResultDatabase.cpp @@ -7,16 +7,23 @@ using namespace std; +#define SORT_BY_NAME 0 +#define SORT_RETAIN_ATTS_ORDER 1 + + bool ResultDatabase::Result::operator<(const Result &rhs) const { if (test < rhs.test) return true; if (test > rhs.test) return false; +#if (SORT_RETAIN_ATTS_ORDER == 0) + // For ties, sort by the value of the attribute: if (atts < rhs.atts) return true; if (atts > rhs.atts) return false; +#endif return false; // less-operator returns false on equal } @@ -189,7 +196,10 @@ void ResultDatabase::AddResult(const string &test_orig, void ResultDatabase::DumpDetailed(ostream &out) { vector sorted(results); - sort(sorted.begin(), sorted.end()); + +#if SORT_BY_NAME + stable_sort(sorted.begin(), sorted.end()); +#endif const int testNameW = 24 ; const int attW = 12; @@ -283,12 +293,15 @@ void ResultDatabase::DumpDetailed(ostream &out) void ResultDatabase::DumpSummary(ostream &out) { vector sorted(results); - sort(sorted.begin(), sorted.end()); - const int testNameW = 24 ; +#if SORT_BY_NAME + stable_sort(sorted.begin(), sorted.end()); +#endif + + const int testNameW = 32 ; const int attW = 12; const int fieldW = 9; - out << std::fixed << right << std::setprecision(4); + out << std::fixed << right << std::setprecision(2); // TODO: in big parallel runs, the "trials" are the procs // and we really don't want to print them all out.... @@ -334,8 +347,8 @@ void ResultDatabase::DumpSummary(ostream &out) } if (0) { out << endl - << "Note: results marked with (*) had missing values such as" << endl - << "might occur with a mixture of architectural capabilities." << endl; + << "Note: results marked with (*) had missing values such as" << endl + << "might occur with a mixture of architectural capabilities." << endl; } } @@ -381,7 +394,9 @@ void ResultDatabase::DumpCsv(string fileName) bool emptyFile; vector sorted(results); - sort(sorted.begin(), sorted.end()); +#if SORT_BY_NAME + stable_sort(sorted.begin(), sorted.end()); +#endif //Check to see if the file is empty - if so, add the headers emptyFile = this->IsFileEmpty(fileName); diff --git a/projects/hip-tests/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp b/projects/hip-tests/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp index b343386b5c..2a4f6ff649 100644 --- a/projects/hip-tests/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp +++ b/projects/hip-tests/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp @@ -25,15 +25,27 @@ THE SOFTWARE. #include #include"ResultDatabase.h" -#define check(msg, status) \ -if(status != hipSuccess){ \ - printf("%s failed.\n",#msg); \ - exit(1); \ +#define PRINT_PROGRESS 0 + +#define check(cmd) \ +{\ + hipError_t status = cmd;\ + if(status != hipSuccess){ \ + printf("error: '%s'(%d) from %s at %s:%d\n", \ + hipGetErrorString(status), status, #cmd,\ + __FILE__, __LINE__); \ + abort(); \ + }\ } #define LEN 1024*1024 -#define SIZE LEN * sizeof(float) -#define ITER 10120 + +#define NUM_GROUPS 1 +#define GROUP_SIZE 64 +#define TEST_ITERS 20 +#define DISPATCHES_PER_TEST 100 + +const unsigned p_tests = 0xfffffff; // HCC optimizes away fully NULL kernel calls, so run one that is nearly null: @@ -44,115 +56,112 @@ __global__ void NearlyNull(hipLaunchParm lp, float* Ad){ } +ResultDatabase resultDB; + + +void stopTest(hipEvent_t start, hipEvent_t stop, const char *msg, int iters) +{ + float mS = 0; + check(hipEventRecord(stop)); + check(hipDeviceSynchronize()); + check(hipEventElapsedTime(&mS, start, stop)); + resultDB.AddResult(std::string(msg), "", "uS", mS*1000/iters); + if (PRINT_PROGRESS & 0x1 ) { + std::cout<< msg <<"\t\t"< Date: Wed, 13 Sep 2017 12:57:37 +0530 Subject: [PATCH 8/8] Add more info for inline asm in hip kernel guide and cookbook readme [ROCm/hip-tests commit: 451f36a42aebfacf3e71d113574f3415574ba99c] --- .../samples/2_Cookbook/10_inline_asm/Readme.md | 15 ++++++++++++++- 1 file changed, 14 insertions(+), 1 deletion(-) diff --git a/projects/hip-tests/samples/2_Cookbook/10_inline_asm/Readme.md b/projects/hip-tests/samples/2_Cookbook/10_inline_asm/Readme.md index 0e64fe9c6e..7d0301bc74 100644 --- a/projects/hip-tests/samples/2_Cookbook/10_inline_asm/Readme.md +++ b/projects/hip-tests/samples/2_Cookbook/10_inline_asm/Readme.md @@ -27,10 +27,23 @@ We will be using the Simple Matrix Transpose application from the our very first ## asm() Assembler statement -We insert the GCN isa into the kernel using asm() Assembler statement. In the same sourcecode, we used for MatrixTranspose. We'll add the following: +In the same sourcecode, we used for MatrixTranspose. We'll add the following: ` asm volatile ("v_mov_b32_e32 %0, %1" : "=v" (out[x*width + y]) : "v" (in[y*width + x])); ` +GCN ISA In-line assembly, is supported. For example: + +``` +asm volatile ("v_mac_f32_e32 %0, %2, %3" : "=v" (out[i]) : "0"(out[i]), "v" (a), "v" (in[i])); +``` + +We insert the GCN isa into the kernel using `asm()` Assembler statement. +`volatile` keyword is used so that the optimizers must not change the number of volatile operations or change their order of execution relative to other volatile operations. +`v_mac_f32_e32` is the GCN instruction, for more information please refer - [AMD GCN3 ISA architecture manual](http://gpuopen.com/compute-product/amd-gcn3-isa-architecture-manual/) +Index for the respective operand in the ordered fashion is provided by `%` followed by position in the list of operands +`"v"` is the constraint code (for target-specific AMDGPU) for 32-bit VGPR register, for more info please refer - [Supported Constraint Code List for AMDGPU](https://llvm.org/docs/LangRef.html#supported-constraint-code-list) +Output Constraints are specified by an `"="` prefix as shown above ("=v"). This indicate that assemby will write to this operand, and the operand will then be made available as a return value of the asm expression. Input constraints do not have a prefix - just the constraint code. The constraint string of `"0"` says to use the assigned register for output as an input as well (it being the 0'th constraint). + ## How to build and run: Use the make command and execute it using ./exe Use hipcc to build the application, which is using hcc on AMD and nvcc on nvidia.