AI/rocm-systems
2
0
Форкнуть 0
Код Задачи Запросы на слияние Действия Пакеты Проекты Релизы Вики Активность

Merge 'master' into 'amd-master'

Просмотр исходного кода 
Change-Id: I69f30c3cbac0dd973147b53f92f7bfaed7610566


[ROCm/hip commit: 1bc3cc369c]
Этот коммит содержится в:
Jenkins
2017-12-06 04:10:57 -06:00
родитель e547608c5e 0835c6f2db
Коммит 3aea692c8d
9 изменённых файлов: 249 добавлений и 176 удалений
Скачать Patch файл Скачать Diff файл Показать все файлы Свернуть все файлы
projects/hip/RELEASE.md
+9 -1
Просмотреть файл
@@ -11,7 +11,15 @@ We have attempted to document known bugs and limitations - in particular the [HI
===================================================================================================
Release: 1.5
Date:
- printf support in device code
- Support threadIdx, blockIdx, blockDim directly (no need for hipify conversions in kernels.) HIP
Kernel syntax is now identical to CUDA kernel syntax - no need for extra parms or conversions.
- Refactor launch syntax. HIP now extracts kernels from the executable and launches them using the
existing module interface. Kernels dispatch no longer flows through HCC. Result is faster
kernel launches and with less resource usage (no signals required).
- Remove requirement for manual "serializers" previously required when passing complex structures
into kernels.
- Remove need for manual destructors
- Provide printf in device code
- Support for globals when using module API
- hipify-clang now supports using newer versions of clang
- HIP texture support equivalent to CUDA texture driver APIs
projects/hip/bin/hipify-perl
-19
Просмотреть файл
@@ -309,25 +309,6 @@ while (@ARGV) {
$ft{'mem'} += s/\bcudaMallocPitch\b/hipMallocPitch/g;
#--------
# Coordinate Indexing and Dimensions:
$ft{'coord_func'} += s/\bthreadIdx\.x\b/hipThreadIdx_x/g;
$ft{'coord_func'} += s/\bthreadIdx\.y\b/hipThreadIdx_y/g;
$ft{'coord_func'} += s/\bthreadIdx\.z\b/hipThreadIdx_z/g;
$ft{'coord_func'} += s/\bblockIdx\.x\b/hipBlockIdx_x/g;
$ft{'coord_func'} += s/\bblockIdx\.y\b/hipBlockIdx_y/g;
$ft{'coord_func'} += s/\bblockIdx\.z\b/hipBlockIdx_z/g;
$ft{'coord_func'} += s/\bblockDim\.x\b/hipBlockDim_x/g;
$ft{'coord_func'} += s/\bblockDim\.y\b/hipBlockDim_y/g;
$ft{'coord_func'} += s/\bblockDim\.z\b/hipBlockDim_z/g;
$ft{'coord_func'} += s/\bgridDim\.x\b/hipGridDim_x/g;
$ft{'coord_func'} += s/\bgridDim\.y\b/hipGridDim_y/g;
$ft{'coord_func'} += s/\bgridDim\.z\b/hipGridDim_z/g;
#--------
# Events
$ft{'event'} += s/\bcudaEvent_t\b/hipEvent_t/g;
projects/hip/docs/markdown/hip_bugs.md
-134
Просмотреть файл
@@ -2,7 +2,6 @@
<!-- toc -->
- [Errors related to undefined reference to `__hcLaunchKernel__***__grid_launch_parm**`](#errors-related-to-undefined-reference-to-__hclaunchkernel____grid_launch_parm)
- [Can't find kernels inside dynamic linked library](#cant-find-kernels-inside-dynamic-linked-library)
- [What is the current limitation of HIP Generic Grid Launch method?](#what-is-the-current-limitation-of-hip-generic-grid-launch-method)
- [Errors related to `no matching constructor`](#errors-related-to-no-matching-constructor)
@@ -10,139 +9,6 @@
<!-- tocstop -->
### Errors related to undefined reference to `__hcLaunchKernel__***__grid_launch_parm**`
Some common code practices may lead to hipcc generating a error with the form :
```
undefined reference to `__hcLaunchKernel__ZN15vecAddNamespace6vecAddIidEEv16grid_launch_parmPT0_S3_S3_T_
```
Or:
```
error: weak declaration cannot have internal linkage
```
Suggested workarounds:
- Avoid use of static with kernel definition:
```c++
static __global__ MyKernel
```
- Avoid defining kernels in anonymous namespace :
```c++
namespace {
__global__ MyKernel
}
```
### Can't find kernels inside dynamic linked library
HCC requires use of the "-Bdynamic" flag when creating a dynamic library which contains kernels. The dynamic flag causes the symbols to be created with a signature which allows HCC to discover and load the kernels in the dynamic library. This flag is often not set by default and must be added to the link step of the library. If not done, HCC will be unable to find the kernels defined in the library, and will emit a message such as:
```
HSADevice::CreateKernel(): Unable to create kernel"
```
To correct, add the following flag to hcc or hipcc:
```
$ hipcc -Wl,-Bsymbolic ...
```
Ensure there is no space in the "Wl,-Bsymbolic" option.
### What is the current limitation of HIP Generic Grid Launch method?
1. __global__ functions cannot be marked as static or put in an unnamed namespace i.e. they cannot be given internal linkage (this would clash with __attribute__((weak)));
2. using the macro based dispatch mechanism i.e. hipLaunchKernel* only works for functions that take no more than 20 arguments (this limit can be increased up to 126, and is temporary until we can enable C++14 mode and use variadic generic lambdas); no such limitation applies do dispatching directly through grid_launch.
### Errors related to `no matching constructor`
The symptom is the compiler would complain about errors like `no matching constructor` for classes/structs passed as arguments into a GPU kernel. Often, this is caused by a design limitation in HCC where array-typed member variables inside a class/struct can’t be correctly passed into GPU kernels. To mitigate this issue, a custom serializer/deserializer pair is provided.
For example, `Foo` in the code snippets below contains an array-typed member variable `table`, which would fail the compiler if used as a kernel argument.
```
struct Foo {
float _data;
// table is an array, which makes foo
int table[3];
};
```
A workaround is to provide a custom serializer on host side which appends the contents of the array as kernel arguments, and a custome deserializaer on the device path to reconstruct the array inside the GPU kernels.
The deserializer can not be a function template, and should have scalar-typed parameters of the number equals to the length of the array-typed member variable. For example:
```
struct Foo {
float _data;
int _table[3];
#ifdef __HCC__
// user-provided CPU serializer
// Append the contents of the array member as kernel arguments
__attribute__((annotate(“serialize”)))
void __cxxamp_serialize(Kalmar::Serialize &s) const {
s.Append(sizeof(float), &_data);
for (int i = 0; i < 3; ++i)
s.Append(sizeof(int), &_table[i]);
}
// user-provided GPU deserializer
// table has 3 int elements, so deserializer must have 3 int parameters.
__attribute__((annotate(“user_deserialize”)))
Foo(float d, int x0, int x1, int x2) [[cpu]][[hc]] {
_data = d;
_table[0] = x0;
_table[1] = x1;
_table[2] = x2;
}
#endif
};
```
Rather than create serializer functions, another workaround is to pass the member fields from the structure as simple data types.
Note a class or struct can contain only one "user_deserialize" constructor.
For types which contain arrays which are based on template parameter, you can use partial template instantiation to implement one constructor per specialization.
However, an easier approach may be to create one user_deserializer which processes the maximum supported dimension.
This will take more memory in the structure and also require additional kernel arguments, but this may have little performance impact and the conversion is easier than partial template specialization. An example:
```
#define MAX_Dim 4
template<typename T, int Dim> struct MyArray {
T* dataPtr_;
//int size_[Dim]; // Original code with template-sized Dims
int size_[MAX_dim]; // Workaround code - allocate an array big enough for all dims so one serializer works.
...
#ifdef __HCC__
__attribute__((annotate("serialize")))
void __cxxamp_serialize(Kalmar::Serialize &s) const {
s.Append(sizeof(float), &_dataPtr);
for (int i=0; i<MAX_Dim; i++) {
s.Append(sizeof(size_[0]), &size_[i]);
}
}
__attribute__((annotate("user_deserialize")))
MyArray(T* data, int size0, int size1, int size2, int size3) [[cpu]][[hc]] {
data_ = data;
size_[0] = size0;
size_[1] = size1;
size_[2] = size2;
size_[3] = size3;
}
#endif
```
### HIP is more restrictive in enforcing restrictions
projects/hip/samples/0_Intro/square/Makefile
+16 -12
Просмотреть файл
@@ -1,20 +1,24 @@
HIP_PATH?= $(wildcard /opt/rocm/hip)
ifeq (,$(HIP_PATH))
HIP_PATH=../../..
endif
HIP_PLATFORM=$(shell $(HIP_PATH)/bin/hipconfig --platform)
HIPCC=$(HIP_PATH)/bin/hipcc
all: square.hip.out
square.cuda.out : square.cu
nvcc square.cu -o $@
#hipify square.cu > square.cpp
# Then review & finish port in square.cpp
#
square.hip.out: square.hipref.cpp
$(HIPCC) $(CXXFLAGS) square.hipref.cpp -o $@
ifeq (${HIP_PLATFORM}, nvcc)
SOURCES=square.cu
else
SOURCES=square.cpp
endif
all: square.out
# Step
square.cpp: square.cu
$(HIP_PATH)/bin/hipify-perl square.cu > square.cpp
square.out: $(SOURCES)
$(HIPCC) $(CXXFLAGS) $(SOURCES) -o $@
clean:
rm -f *.o *.out
rm -f *.o *.out square.cpp
projects/hip/samples/0_Intro/square/README.md
+6 -9
Просмотреть файл
@@ -1,16 +1,13 @@
# Square.md
Simple test which shows how to use hipify to port CUDA code to HIP.
See related [blog](http://gpuopen.com/hip-to-be-squared-an-introductory-hip-tutorial) that explains the example.
See related [blog](http://gpuopen.com/hip-to-be-squared-an-introductory-hip-tutorial) that explains the example.
Now it is even simpler and requires no manual modification to the hipified source code - just hipify and compile:
1. Add hip/bin path to the PATH :
<code>export PATH=$PATH:[MYHIP]/bin</code>
2. Do <code>$ hipify square.cu > square.cpp </code>
3. Manually edit square.cpp to add hipLaunchParms lp to kernel parms:
<code>vector_square(hipLaunchParm lp, T *C_d, const T *A_d, size_t N)</code>
(see square.hipref.cpp for the correct output after running hipify and the above manual step)
4. make
2. <code>$ make </code>
Make runs these steps. This can be performed on either CUDA or AMD platform:
<code>hipify-perl square.cu > square.cpp </code> # convert cuda code to hip code
<code>hipcc square.cpp</code> # compile into executable
projects/hip/samples/0_Intro/square/square.cu
+1 -1
Просмотреть файл
@@ -38,7 +38,7 @@ THE SOFTWARE.
*/
template <typename T>
__global__ void
vector_square(T *C_d, const T *A_d, size_t N)
vector_square(T *C_d, T *A_d, size_t N)
{
size_t offset = (blockIdx.x * blockDim.x + threadIdx.x);
size_t stride = blockDim.x * gridDim.x ;
projects/hip/samples/2_Cookbook/12_cmake_hip_add_executable/CMakeLists.txt
+28
Просмотреть файл
@@ -0,0 +1,28 @@
cmake_minimum_required(VERSION 2.8.3)
if(NOT DEFINED HIP_PATH)
if(NOT DEFINED ENV{HIP_PATH})
set(HIP_PATH "/opt/rocm/hip" CACHE PATH "Path to which HIP has been installed")
else()
set(HIP_PATH $ENV{HIP_PATH} CACHE PATH "Path to which HIP has been installed")
endif()
endif()
set(CMAKE_MODULE_PATH "${HIP_PATH}/cmake" ${CMAKE_MODULE_PATH})
project(12_cmake)
find_package(HIP QUIET)
if(HIP_FOUND)
message(STATUS "Found HIP: " ${HIP_VERSION})
else()
message(FATAL_ERROR "Could not find HIP. Ensure that HIP is either installed in /opt/rocm/hip or the variable HIP_PATH is set to point to the right location.")
endif()
set(MY_SOURCE_FILES MatrixTranspose.cpp)
set(MY_TARGET_NAME MatrixTranspose)
set(MY_HIPCC_OPTIONS)
set(MY_HCC_OPTIONS)
set(MY_NVCC_OPTIONS)
set_source_files_properties(${MY_SOURCE_FILES} PROPERTIES HIP_SOURCE_PROPERTY_FORMAT 1)
hip_add_executable(${MY_TARGET_NAME} ${MY_SOURCE_FILES} HIPCC_OPTIONS ${MY_HIPCC_OPTIONS} HCC_OPTIONS ${MY_HCC_OPTIONS} NVCC_OPTIONS ${MY_NVCC_OPTIONS})
projects/hip/samples/2_Cookbook/12_cmake_hip_add_executable/MatrixTranspose.cpp
+136
Просмотреть файл
@@ -0,0 +1,136 @@
/*
Copyright (c) 2015-2016 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<iostream>
// hip header file
#include "hip/hip_runtime.h"
#define WIDTH 1024
#define NUM (WIDTH*WIDTH)
#define THREADS_PER_BLOCK_X 4
#define THREADS_PER_BLOCK_Y 4
#define THREADS_PER_BLOCK_Z 1
// Device (Kernel) function, it must be void
// hipLaunchParm provides the execution configuration
__global__ void matrixTranspose(hipLaunchParm lp,
float *out,
float *in,
const int width)
{
int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
out[y * width + x] = in[x * width + y];
}
// CPU implementation of matrix transpose
void matrixTransposeCPUReference(
float * output,
float * input,
const unsigned int width)
{
for(unsigned int j=0; j < width; j++)
{
for(unsigned int i=0; i < width; i++)
{
output[i*width + j] = input[j*width + i];
}
}
}
int main() {
float* Matrix;
float* TransposeMatrix;
float* cpuTransposeMatrix;
float* gpuMatrix;
float* gpuTransposeMatrix;
hipDeviceProp_t devProp;
hipGetDeviceProperties(&devProp, 0);
std::cout << "Device name " << devProp.name << std::endl;
int i;
int errors;
Matrix = (float*)malloc(NUM * sizeof(float));
TransposeMatrix = (float*)malloc(NUM * sizeof(float));
cpuTransposeMatrix = (float*)malloc(NUM * sizeof(float));
// initialize the input data
for (i = 0; i < NUM; i++) {
Matrix[i] = (float)i*10.0f;
}
// allocate the memory on the device side
hipMalloc((void**)&gpuMatrix, NUM * sizeof(float));
hipMalloc((void**)&gpuTransposeMatrix, NUM * sizeof(float));
// Memory transfer from host to device
hipMemcpy(gpuMatrix, Matrix, NUM*sizeof(float), hipMemcpyHostToDevice);
// Lauching kernel from host
hipLaunchKernel(matrixTranspose,
dim3(WIDTH/THREADS_PER_BLOCK_X, WIDTH/THREADS_PER_BLOCK_Y),
dim3(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y),
0, 0,
gpuTransposeMatrix , gpuMatrix, WIDTH);
// Memory transfer from device to host
hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM*sizeof(float), hipMemcpyDeviceToHost);
// CPU MatrixTranspose computation
matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH);
// verify the results
errors = 0;
double eps = 1.0E-6;
for (i = 0; i < NUM; i++) {
if (std::abs(TransposeMatrix[i] - cpuTransposeMatrix[i]) > eps ) {
errors++;
}
}
if (errors!=0) {
printf("FAILED: %d errors\n",errors);
} else {
printf ("PASSED!\n");
}
//free the resources on device side
hipFree(gpuMatrix);
hipFree(gpuTransposeMatrix);
//free the resources on host side
free(Matrix);
free(TransposeMatrix);
free(cpuTransposeMatrix);
return errors;
}
projects/hip/samples/2_Cookbook/12_cmake_hip_add_executable/Readme.md
+53
Просмотреть файл
@@ -0,0 +1,53 @@
## hip_add_executable ###
This tutorial shows how to use the FindHIP cmake module and create an executable using ```hip_add_executable``` macro.
## Including FindHIP cmake module in the project
Since FindHIP cmake module is not yet a part of the default cmake distribution, ```CMAKE_MODULE_PATH``` needs to be updated to contain the path to FindHIP.cmake.
The simplest approach is to use
```
set(CMAKE_MODULE_PATH "/opt/rocm/hip/cmake" ${CMAKE_MODULE_PATH})
find_package(HIP)
```
A more generic solution that allows for a user specified location for the HIP installation would look something like
```
if(NOT DEFINED HIP_PATH)
if(NOT DEFINED ENV{HIP_PATH})
set(HIP_PATH "/opt/rocm/hip" CACHE PATH "Path to which HIP has been installed")
else()
set(HIP_PATH $ENV{HIP_PATH} CACHE PATH "Path to which HIP has been installed")
endif()
endif()
set(CMAKE_MODULE_PATH "${HIP_PATH}/cmake" ${CMAKE_MODULE_PATH})
find_package(HIP)
```
If your project already modifies ```CMAKE_MODULE_PATH```, you will need to append the path to FindHIP.cmake instead of replacing it.
## 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, NVCC-specific compiler options using the ```HCC_OPTIONS``` and ```NVCC_OPTIONS``` keywords.
Common options targeting both compilers can be specificed after the ```HIPCC_OPTIONS``` keyword.
## How to build and run:
Use the following commands to build and execute the sample
```
mkdir build
cd build
cmake ..
make
./MatrixTranspose
```
## More Info:
- [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)