AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Remove hip-hcc codes: Part four

Browse Source 
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
This commit is contained in:
Tao Sang
2021-01-04 19:58:20 -05:00
committed by Tao Sang
parent 693b89b3dd
commit 22094572e1
10 changed files with 1 additions and 477 deletions
Download Patch File Download Diff File Expand all files Collapse all files
samples/0_Intro/hcc_dialects/.gitignore
-5
View File
@@ -1,5 +0,0 @@
vadd_amp_arrayview
vadd_hc_am
vadd_hc_array
vadd_hc_arrayview
vadd_hip
samples/0_Intro/hcc_dialects/Makefile
-70
View File
@@ -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 $@
samples/0_Intro/hcc_dialects/vadd_hc_am.cpp
-91
View File
@@ -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 <hc/hc.hpp>
#include <hc/hc_am.hpp>
#else
#include <hc.hpp>
#include <hc_am.hpp>
#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");
}
samples/0_Intro/hcc_dialects/vadd_hc_array.cpp
-81
View File
@@ -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 <hc/hc.hpp>
#else
#include <hc.hpp>
#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<float> A_d(sizeElements);
hc::array<float> B_d(sizeElements);
hc::array<float> 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");
}
samples/0_Intro/hcc_dialects/vadd_hc_array.hc
-62
View File
@@ -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 <hc/hc.hpp>
#else
#include <hc.hpp>
#endif
int main(int argc, char *argv[])
{
int size = 1000000;
bool pass = true;
// Allocate auto-managed host/device views of data:
hc::array_view<float> A(size);
hc::array_view<float> B(size);
hc::array_view<float> C(size);
// Initialize host data
for (int i=0; i<size; 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:
hc::parallel_for_each(hc::extent<1> (size),
[=] (hc::index<1> idx) [[hc]] {
int i = idx[0];
C[i] = A[i] + B[i];
});
for (int i=0; i<size; i++) {
float ref= 1.618f * i + 3.142f * i;
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");
}
samples/0_Intro/hcc_dialects/vadd_hc_arrayview.cpp
-77
View File
@@ -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.
*/
// Simple test showing how to use HC syntax with array_view.
// The code uses AMP's array_view class, which provides automatic data synchronization
// of data between the host and the accelerator. As noted below, the HCC runtime
// will automatically copy data to and from the host, without the user needing
// to manually perform such copies. This is an excellent mode for developers
// new to GPU programming and matches the memory models provided by recent systems where
// CPU and GPU share the same memory pool. Advanced programmers may prefer
// more explicit control over the data movement - shown in the other vadd_hc_array and
// vadd_hc_am examples.
// This example shows the similarity between C++AMP and and HC for simple cases where
// implicit data transfer is used - really the only difference is the namespace.
// Other examples show some of the more advanced controls.
#if defined(HC_NEXT)
#include <hc/hc.hpp>
#else
#include <hc.hpp>
#endif
int main(int argc, char* argv[]) {
int sizeElements = 1000000;
bool pass = true;
// Allocate auto-managed host/device views of data:
hc::array_view<float> A(sizeElements);
hc::array_view<float> B(sizeElements);
hc::array_view<float> 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");
}
samples/0_Intro/hcc_dialects/vadd_hip.cpp
-77
View File
@@ -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");
}
samples/1_Utils/hipDispatchLatency/hipDispatchLatency.cpp
-10
View File
@@ -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);
samples/1_Utils/hipInfo/hipInfo.cpp
-3
View File
@@ -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
samples/2_Cookbook/12_cmake_hip_add_executable/Readme.md
+1 -1
View File
@@ -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: