From 22094572e19c7d5fa6bc54f777c75c8deec57e82 Mon Sep 17 00:00:00 2001 From: Tao Sang Date: Mon, 4 Jan 2021 19:58:20 -0500 Subject: [PATCH] Remove hip-hcc codes: Part four Remove __HCC__, __HCC_ONLY__, __HCC_CPP__, __HCC_C__, __HCC_OR_HIP_CLANG__, __HIP_ROCclr__ and their guarded codes. Remove Hcc codes from directed_tests and samples. Remove __HIP_PLATFORM_HCC__ and __HIP_PLATFORM_NVCC__ from some files where they are not necessary. Add deprecation notice. Change-Id: I1ae467eafd749d6c25bca204c1724b026be21fce --- samples/0_Intro/hcc_dialects/.gitignore | 5 - samples/0_Intro/hcc_dialects/Makefile | 70 -------------- samples/0_Intro/hcc_dialects/vadd_hc_am.cpp | 91 ------------------- .../0_Intro/hcc_dialects/vadd_hc_array.cpp | 81 ----------------- samples/0_Intro/hcc_dialects/vadd_hc_array.hc | 62 ------------- .../hcc_dialects/vadd_hc_arrayview.cpp | 77 ---------------- samples/0_Intro/hcc_dialects/vadd_hip.cpp | 77 ---------------- .../hipDispatchLatency/hipDispatchLatency.cpp | 10 -- samples/1_Utils/hipInfo/hipInfo.cpp | 3 - .../12_cmake_hip_add_executable/Readme.md | 2 +- 10 files changed, 1 insertion(+), 477 deletions(-) delete mode 100644 samples/0_Intro/hcc_dialects/.gitignore delete mode 100644 samples/0_Intro/hcc_dialects/Makefile delete mode 100644 samples/0_Intro/hcc_dialects/vadd_hc_am.cpp delete mode 100644 samples/0_Intro/hcc_dialects/vadd_hc_array.cpp delete mode 100644 samples/0_Intro/hcc_dialects/vadd_hc_array.hc delete mode 100644 samples/0_Intro/hcc_dialects/vadd_hc_arrayview.cpp delete mode 100644 samples/0_Intro/hcc_dialects/vadd_hip.cpp diff --git a/samples/0_Intro/hcc_dialects/.gitignore b/samples/0_Intro/hcc_dialects/.gitignore deleted file mode 100644 index bce1cdf193..0000000000 --- a/samples/0_Intro/hcc_dialects/.gitignore +++ /dev/null @@ -1,5 +0,0 @@ -vadd_amp_arrayview -vadd_hc_am -vadd_hc_array -vadd_hc_arrayview -vadd_hip diff --git a/samples/0_Intro/hcc_dialects/Makefile b/samples/0_Intro/hcc_dialects/Makefile deleted file mode 100644 index 4a514b6691..0000000000 --- a/samples/0_Intro/hcc_dialects/Makefile +++ /dev/null @@ -1,70 +0,0 @@ -HCC_HOME?=/opt/rocm/hcc -HCC = $(HCC_HOME)/bin/hcc - -OPT=-O2 -HCC_CFLAGS= `$(HCC_HOME)/bin/hcc-config --cxxflags` ${OPT} -HCC_LDFLAGS= `$(HCC_HOME)/bin/hcc-config --ldflags` ${OPT} - -CPPAMP_CFLAGS= `$(HCC_HOME)/bin/clamp-config --cxxflags` -CPPAMP_LDFLAGS= `$(HCC_HOME)/bin/clamp-config --ldflags` - -HIP_PATH?= $(wildcard /opt/rocm/hip) -ifeq (,$(HIP_PATH)) - HIP_PATH=../../.. -endif -HIPCC=$(HIP_PATH)/bin/hipcc -HIP_PLATFORM=$(shell $(HIP_PATH)/bin/hipconfig --platform) - -ifneq (${HIP_PLATFORM}, hcc) -$(error hcc_dialects requires hcc compiler and only runs on hcc platform) -endif - - -TARGETS=vadd_hc_arrayview vadd_hc_array vadd_hc_am vadd_amp_arrayview vadd_hip - -all: $(TARGETS) - -clean: - rm -f $(TARGETS) *.o - -run: $(TARGETS) - @for t in $(TARGETS); do\ - echo "Running $$t"; \ - ./$$t; \ - done - - -# HCC version: -vadd_hc_arrayview.o: vadd_hc_arrayview.cpp - $(HCC) $(HCC_CFLAGS) -c $< -o $@ -vadd_hc_arrayview: vadd_hc_arrayview.o - $(HCC) $(HCC_LDFLAGS) $< -o $@ - - -# HCC version, using explicit arrays: -vadd_hc_array.o: vadd_hc_array.cpp - $(HCC) $(HCC_CFLAGS) -c $< -o $@ -vadd_hc_array: vadd_hc_array.o - $(HCC) $(HCC_LDFLAGS) $< -o $@ - - -# HCC version, using AM (accelerator memory) pointer -vadd_hc_am.o: vadd_hc_am.cpp - $(HCC) $(HCC_CFLAGS) -c $< -o $@ -vadd_hc_am: vadd_hc_am.o - $(HCC) $(HCC_LDFLAGS) -lhc_am $< -o $@ - - - -# HIP version: -vadd_hip.o: vadd_hip.cpp - $(HIPCC) -c $< -o $@ -vadd_hip: vadd_hip.o - $(HIPCC) $< -o $@ - - -# AMP version: -vadd_amp_arrayview.o: vadd_amp_arrayview.cpp - $(HCC) $(CPPAMP_CFLAGS) -c $< -o $@ -vadd_amp_arrayview: vadd_amp_arrayview.o - $(HCC) $(CPPAMP_LDFLAGS) $< -o $@ diff --git a/samples/0_Intro/hcc_dialects/vadd_hc_am.cpp b/samples/0_Intro/hcc_dialects/vadd_hc_am.cpp deleted file mode 100644 index 3fc0c8c27a..0000000000 --- a/samples/0_Intro/hcc_dialects/vadd_hc_am.cpp +++ /dev/null @@ -1,91 +0,0 @@ -/* -Copyright (c) 2015-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 WARRANTY OF ANY KIND, EXPRESS OR -IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, -FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE -AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER -LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, -OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN -THE SOFTWARE. -*/ - -// Simple test showing how to use HC syntax with AM (accelerator memory). -// AM provides a set of c-style memory management routines for allocating, -// freeing, and copying memory. am_alloc returns a device pointer -// which can only be used on the device. The programmer has full control -// over when data is copied. - -#if defined(HC_NEXT) - #include - #include -#else - #include - #include -#endif - -int main(int argc, char* argv[]) { - int sizeElements = 1000000; - size_t sizeBytes = sizeElements * sizeof(float); - bool pass = true; - - // Allocate host memory - float* A_h = (float*)malloc(sizeBytes); - float* B_h = (float*)malloc(sizeBytes); - float* C_h = (float*)malloc(sizeBytes); - - // Allocate device pointers: - // Unlike array_view, these must be explicitly managed by user: - hc::accelerator acc; // grab default accelerator where we want to allocate memory: - hc::accelerator_view av = acc.get_default_view(); - - float *A_d, *B_d, *C_d; - A_d = hc::am_alloc(sizeBytes, acc, 0); - B_d = hc::am_alloc(sizeBytes, acc, 0); - C_d = hc::am_alloc(sizeBytes, acc, 0); - - // Initialize host data - for (int i = 0; i < sizeElements; i++) { - A_h[i] = 1.618f * i; - B_h[i] = 3.142f * i; - C_h[i] = 0; - } - - av.copy(A_h, A_d, sizeBytes); // C++ copy H2D - av.copy(B_h, B_d, sizeBytes); // C++ copy H2D - - // Launch kernel onto AV. - // Because the kernel PFE and the copies are submitted to same AV, they will execute in order - // and we don't need additional synchronization to ensure the copies complete before the PFE - // begins. - hc::completion_future cf = - hc::parallel_for_each(av, hc::extent<1>(sizeElements), [=](hc::index<1> idx)[[hc]] { - int i = idx[0]; - C_d[i] = A_d[i] + B_d[i]; - }); - - - // This copy is in same AV as the kernel and thus will wait for the kernel to finish before - // executing. - av.copy(C_d, C_h, sizeBytes); // C++ copy D2H - - - for (int i = 0; i < sizeElements; i++) { - float ref = 1.618f * i + 3.142f * i; - if (C_h[i] != ref) { - printf("error:%d computed=%6.2f, reference=%6.2f\n", i, C_h[i], ref); - pass = false; - } - }; - if (pass) printf("PASSED!\n"); -} diff --git a/samples/0_Intro/hcc_dialects/vadd_hc_array.cpp b/samples/0_Intro/hcc_dialects/vadd_hc_array.cpp deleted file mode 100644 index 635b4bff21..0000000000 --- a/samples/0_Intro/hcc_dialects/vadd_hc_array.cpp +++ /dev/null @@ -1,81 +0,0 @@ -/* -Copyright (c) 2015-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 WARRANTY OF ANY KIND, EXPRESS OR -IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, -FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE -AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER -LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, -OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN -THE SOFTWARE. -*/ - -// Simple test showing how to use HC syntax with array. -// Array provides a type-safe C++ mechanism to allocate accelerator memory. -// Like array_view, hc::array provides multi-dimensional indexing capability, -// and is typed. However, unlike array_view, hc::array does not provide -// automatic data management capabilities - instead the programmer -// takes the reins and controls when copies are executed. - -#if defined(HC_NEXT) - #include -#else - #include -#endif - -int main(int argc, char* argv[]) { - int sizeElements = 1000000; - size_t sizeBytes = sizeElements * sizeof(float); - bool pass = true; - - // Allocate host memory - float* A_h = (float*)malloc(sizeBytes); - float* B_h = (float*)malloc(sizeBytes); - float* C_h = (float*)malloc(sizeBytes); - - // Allocate device arrays<> - // Unlike array_view, these must be explicitly managed by user: - hc::array A_d(sizeElements); - hc::array B_d(sizeElements); - hc::array C_d(sizeElements); - - // Initialize host data - for (int i = 0; i < sizeElements; i++) { - A_h[i] = 1.618f * i; - B_h[i] = 3.142f * i; - } - - hc::copy(A_h, A_d); // C++ copy H2D - hc::copy(B_h, B_d); // C++ copy H2D - - // Launch kernel onto default accelerator: - // array<> types are not implicitly copied, so we performed copies above. - hc::parallel_for_each(hc::extent<1>(sizeElements), [&](hc::index<1> idx)[[hc]] { - int i = idx[0]; - C_d[i] = A_d[i] + B_d[i]; - }); - - // HCC runtime knows that C_d depends on previous PFE and will force the copy to wait for the - // PFE to complte. - hc::copy(C_d, C_h); // C++ copy D2H - - - for (int i = 0; i < sizeElements; i++) { - float ref = 1.618f * i + 3.142f * i; - if (C_h[i] != ref) { - printf("error:%d computed=%6.2f, reference=%6.2f\n", i, C_h[i], ref); - pass = false; - } - }; - if (pass) printf("PASSED!\n"); -} diff --git a/samples/0_Intro/hcc_dialects/vadd_hc_array.hc b/samples/0_Intro/hcc_dialects/vadd_hc_array.hc deleted file mode 100644 index 491ba9568e..0000000000 --- a/samples/0_Intro/hcc_dialects/vadd_hc_array.hc +++ /dev/null @@ -1,62 +0,0 @@ -/* -Copyright (c) 2015-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 WARRANTY OF ANY KIND, EXPRESS OR -IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, -FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE -AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER -LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, -OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN -THE SOFTWARE. -*/ - -#if defined(HC_NEXT) - #include -#else - #include -#endif - -int main(int argc, char *argv[]) -{ - int size = 1000000; - bool pass = true; - - // Allocate auto-managed host/device views of data: - hc::array_view A(size); - hc::array_view B(size); - hc::array_view C(size); - - // Initialize host data - for (int i=0; i (size), - [=] (hc::index<1> idx) [[hc]] { - int i = idx[0]; - C[i] = A[i] + B[i]; - }); - - for (int i=0; i -#else - #include -#endif - -int main(int argc, char* argv[]) { - int sizeElements = 1000000; - bool pass = true; - - // Allocate auto-managed host/device views of data: - hc::array_view A(sizeElements); - hc::array_view B(sizeElements); - hc::array_view C(sizeElements); - - // Initialize host data - for (int i = 0; i < sizeElements; i++) { - A[i] = 1.618f * i; - B[i] = 3.142f * i; - } - C.discard_data(); // tell runtime not to copy CPU host data. - - - // Launch kernel onto default accelerator: - // The HCC runtime will ensure that A and B are available on the accelerator before launching - // the kernel. - hc::parallel_for_each(hc::extent<1>(sizeElements), [=](hc::index<1> idx)[[hc]] { - int i = idx[0]; - C[i] = A[i] + B[i]; - }); - - for (int i = 0; i < sizeElements; i++) { - float ref = 1.618f * i + 3.142f * i; - // Because C is an array_view, the HCC runtime will copy C back to host at first access - // here: - if (C[i] != ref) { - printf("error:%d computed=%6.2f, reference=%6.2f\n", i, C[i], ref); - pass = false; - } - }; - if (pass) printf("PASSED!\n"); -} diff --git a/samples/0_Intro/hcc_dialects/vadd_hip.cpp b/samples/0_Intro/hcc_dialects/vadd_hip.cpp deleted file mode 100644 index 79605acff2..0000000000 --- a/samples/0_Intro/hcc_dialects/vadd_hip.cpp +++ /dev/null @@ -1,77 +0,0 @@ -/* -Copyright (c) 2015-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 WARRANTY OF ANY KIND, EXPRESS OR -IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, -FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE -AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER -LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, -OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN -THE SOFTWARE. -*/ - -#include "hip/hip_runtime.h" - -__global__ void vadd_hip(const float* a, const float* b, float* c, int N) { - int idx = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x); - - if (idx < N) { - c[idx] = a[idx] + b[idx]; - } -} - - -int main(int argc, char* argv[]) { - int sizeElements = 1000000; - size_t sizeBytes = sizeElements * sizeof(float); - bool pass = true; - - // Allocate host memory - float* A_h = (float*)malloc(sizeBytes); - float* B_h = (float*)malloc(sizeBytes); - float* C_h = (float*)malloc(sizeBytes); - - // Allocate device memory: - float *A_d, *B_d, *C_d; - hipMalloc(&A_d, sizeBytes); - hipMalloc(&B_d, sizeBytes); - hipMalloc(&C_d, sizeBytes); - - // Initialize host memory - for (int i = 0; i < sizeElements; i++) { - A_h[i] = 1.618f * i; - B_h[i] = 3.142f * i; - } - - // H2D Copy - hipMemcpy(A_d, A_h, sizeBytes, hipMemcpyHostToDevice); - hipMemcpy(B_d, B_h, sizeBytes, hipMemcpyHostToDevice); - - // Launch kernel onto default accelerator - int blockSize = 256; // pick arbitrary block size - int blocks = (sizeElements + blockSize - 1) / blockSize; // round up to launch enough blocks - hipLaunchKernelGGL(vadd_hip, dim3(blocks), dim3(blockSize), 0, 0, A_d, B_d, C_d, sizeElements); - - // D2H Copy - hipMemcpy(C_h, C_d, sizeBytes, hipMemcpyDeviceToHost); - - // Verify - for (int i = 0; i < sizeElements; i++) { - float ref = 1.618f * i + 3.142f * i; - if (C_h[i] != ref) { - printf("error:%d computed=%6.2f, reference=%6.2f\n", i, C_h[i], ref); - pass = false; - } - }; - if (pass) printf("PASSED!\n"); -} diff --git a/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp b/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp index fb4ab8455c..f34c6852bd 100644 --- a/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp +++ b/samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp @@ -108,16 +108,6 @@ int main() { /* Measures latency to start & finish executing a kernel with GPU-scope visibility */ /***********************************************************************************/ - //Timing directly the dispatch -#if defined(__HIP_PLATFORM_AMD__) && GENERIC_GRID_LAUNCH == 1 && defined(__HCC__) - for (auto i = 0; i < TOTAL_RUN_COUNT; ++i) { - hipExtLaunchKernelGGL((EmptyKernel), dim3(NUM_GROUPS), dim3(GROUP_SIZE), 0, stream0, start, stop, 0); - hipEventSynchronize(stop); - hipEventElapsedTime(&results[i], start, stop); - } - print_timing("Timing directly single dispatch latency", results); -#endif - //Timing around the dispatch for (auto i = 0; i < TOTAL_RUN_COUNT; ++i) { hipEventRecord(start, 0); diff --git a/samples/1_Utils/hipInfo/hipInfo.cpp b/samples/1_Utils/hipInfo/hipInfo.cpp index e15ce583b9..7bdf7e7b54 100644 --- a/samples/1_Utils/hipInfo/hipInfo.cpp +++ b/samples/1_Utils/hipInfo/hipInfo.cpp @@ -48,9 +48,6 @@ THE SOFTWARE. } void printCompilerInfo() { -#ifdef __HCC__ - printf("compiler: hcc version=%s, workweek (YYWWD) = %u\n", __hcc_version__, __hcc_workweek__); -#endif #ifdef __NVCC__ printf("compiler: nvcc\n"); #endif diff --git a/samples/2_Cookbook/12_cmake_hip_add_executable/Readme.md b/samples/2_Cookbook/12_cmake_hip_add_executable/Readme.md index 9aa6b96fab..b2a4547670 100644 --- a/samples/2_Cookbook/12_cmake_hip_add_executable/Readme.md +++ b/samples/2_Cookbook/12_cmake_hip_add_executable/Readme.md @@ -28,7 +28,7 @@ If your project already modifies ```CMAKE_MODULE_PATH```, you will need to appen ## Using the hip_add_executable macro FindHIP provides the ```hip_add_executable``` macro that is similar to the ```cuda_add_executable``` macro that is provided by FindCUDA. The syntax is also similar. The ```hip_add_executable``` macro uses the hipcc wrapper as the compiler. -The macro supports specifying HCC-specific, CLANG-specific, NVCC-specific compiler options using the ```HCC_OPTIONS```, ```CLANG_OPTIONS``` and ```NVCC_OPTIONS``` keywords. +The macro supports specifying CLANG-specific, NVCC-specific compiler options using the ```CLANG_OPTIONS``` and ```NVCC_OPTIONS``` keywords. Common options targeting both compilers can be specificed after the ```HIPCC_OPTIONS``` keyword. ## How to build and run: