Merge 'master' into 'amd-master'
Change-Id: I2c8409732f0c2397d6c65f1c39c0fb847bd1bcd8
[ROCm/hip commit: 0450d99367]
Этот коммит содержится в:
Jenkins
@@ -32,30 +32,30 @@
|
||||
## <a name="dependencies"></a> Dependencies
|
||||
|
||||
`hipify-clang` requires:
|
||||
1. **LLVM+CLANG** of at least version 3.8.0; the latest stable and recommended release: **6.0.1 on Windows**, and **8.0.0 on Linux**.
|
||||
1. [**LLVM+CLANG**](http://releases.llvm.org) of at least version [3.8.0](http://releases.llvm.org/download.html#3.8.0); the latest stable and recommended release: [**6.0.1**](http://releases.llvm.org/download.html#6.0.1) on **Windows**, and [**8.0.0**](http://releases.llvm.org/download.html#8.0.0) on **Linux**.
|
||||
|
||||
2. **CUDA** at least version 7.0, the latest supported version is **9.0 on Windows**, and **10.0 on Linux**.
|
||||
2. **CUDA** at least version [7.0](https://developer.nvidia.com/cuda-toolkit-70), the latest supported version is [**9.0**](https://developer.nvidia.com/cuda-90-download-archive) on **Windows**, and [**10.0**](https://developer.nvidia.com/cuda-10.0-download-archive) on **Linux**.
|
||||
|
||||
If the target CUDA is 9.1, 9.2 or 10.0, to work on Windows you may apply patches*: [for LLVM 7.0.0](patches/patch_for_clang_7.0.0_bug_38811.zip), [for LLVM 7.0.1](patches/patch_for_clang_7.0.1_bug_38811.zip), [for LLVM 8.0.0](patches/patch_for_clang_8.0.0_bug_38811.zip).
|
||||
If the target CUDA is [9.1](https://developer.nvidia.com/cuda-91-download-archive), [9.2](https://developer.nvidia.com/cuda-92-download-archive) or [10.0](https://developer.nvidia.com/cuda-10.0-download-archive), to work on Windows you may apply patches*: [for LLVM 7.0.0](patches/patch_for_clang_7.0.0_bug_38811.zip), [for LLVM 7.0.1](patches/patch_for_clang_7.0.1_bug_38811.zip), [for LLVM 8.0.0](patches/patch_for_clang_8.0.0_bug_38811.zip).
|
||||
|
||||
| **LLVM release version** | **CUDA latest supported version** | **Windows** | **Linux** |
|
||||
|:------------------------:|:---------------------------------:|:------------:|:---------:|
|
||||
| 3.8.0 | 7.5 | + | + |
|
||||
| 3.8.1 | 7.5 | + | + |
|
||||
| 3.9.0 | 7.5 | + | + |
|
||||
| 3.9.1 | 7.5 | + | + |
|
||||
| 4.0.0 | 8.0 | + | + |
|
||||
| 4.0.1 | 8.0 | + | + |
|
||||
| 5.0.0 | 8.0 | + | + |
|
||||
| 5.0.1 | 8.0 | + | + |
|
||||
| 5.0.2 | 8.0 | + | + |
|
||||
| 6.0.0 | 9.0 | + | + |
|
||||
| **6.0.1** | **9.0** | + <br/> **LATEST STABLE RELEASE** | + |
|
||||
| 7.0.0 | 9.2 | - <br/> not working due to <br/> the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) <br/>+<br/>[patch](patches/patch_for_clang_7.0.0_bug_38811.zip)*</br> | - <br/> not working due to <br/> the clang's bug [36384](https://bugs.llvm.org/show_bug.cgi?id=36384) |
|
||||
| 7.0.1 | 9.2 | - <br/> not working due to <br/> the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) <br/>+<br/>[patch](patches/patch_for_clang_7.0.1_bug_38811.zip)*</br> | - <br/> not working due to <br/> the clang's bug [36384](https://bugs.llvm.org/show_bug.cgi?id=36384) |
|
||||
| 7.1.0 | 9.2 (?) | - <br/> LLVM 7.1.0 <br/> is not yet released | - <br/> LLVM 7.1.0 <br/> is not yet released |
|
||||
| **8.0.0** | **10.0** | - <br/> not working due to <br/> the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) <br/>+<br/>[patch](patches/patch_for_clang_8.0.0_bug_38811.zip)*</br> | + <br/> **LATEST STABLE RELEASE** |
|
||||
| 8.0.1 | 10.1 (?) | - <br/> LLVM 8.0.1 <br/> is not yet released | - <br/> LLVM 8.0.1 <br/> is not yet released |
|
||||
| **LLVM release version** | **CUDA latest supported version** | **Windows** | **Linux** |
|
||||
|:----------------------------------------------------------:|:-------------------------------------------------------------------:|:------------:|:---------:|
|
||||
| [3.8.0](http://releases.llvm.org/download.html#3.8.0) | [7.5](https://developer.nvidia.com/cuda-75-downloads-archive) | + | + |
|
||||
| [3.8.1](http://releases.llvm.org/download.html#3.8.1) | [7.5](https://developer.nvidia.com/cuda-75-downloads-archive) | + | + |
|
||||
| [3.9.0](http://releases.llvm.org/download.html#3.9.0) | [7.5](https://developer.nvidia.com/cuda-75-downloads-archive) | + | + |
|
||||
| [3.9.1](http://releases.llvm.org/download.html#3.9.1) | [7.5](https://developer.nvidia.com/cuda-75-downloads-archive) | + | + |
|
||||
| [4.0.0](http://releases.llvm.org/download.html#4.0.0) | [8.0](https://developer.nvidia.com/cuda-80-ga2-download-archive) | + | + |
|
||||
| [4.0.1](http://releases.llvm.org/download.html#4.0.1) | [8.0](https://developer.nvidia.com/cuda-80-ga2-download-archive) | + | + |
|
||||
| [5.0.0](http://releases.llvm.org/download.html#5.0.0) | [8.0](https://developer.nvidia.com/cuda-80-ga2-download-archive) | + | + |
|
||||
| [5.0.1](http://releases.llvm.org/download.html#5.0.1) | [8.0](https://developer.nvidia.com/cuda-80-ga2-download-archive) | + | + |
|
||||
| [5.0.2](http://releases.llvm.org/download.html#5.0.2) | [8.0](https://developer.nvidia.com/cuda-80-ga2-download-archive) | + | + |
|
||||
| [6.0.0](http://releases.llvm.org/download.html#6.0.0) | [9.0](https://developer.nvidia.com/cuda-90-download-archive) | + | + |
|
||||
| [**6.0.1**](http://releases.llvm.org/download.html#6.0.1) | [**9.0**](https://developer.nvidia.com/cuda-90-download-archive) | + <br/> **LATEST STABLE RELEASE** | + |
|
||||
| [7.0.0](http://releases.llvm.org/download.html#7.0.0) | [9.2](https://developer.nvidia.com/cuda-92-download-archive) | - <br/> not working due to <br/> the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) <br/>+<br/>[patch](patches/patch_for_clang_7.0.0_bug_38811.zip)*</br> | - <br/> not working due to <br/> the clang's bug [36384](https://bugs.llvm.org/show_bug.cgi?id=36384) |
|
||||
| [7.0.1](http://releases.llvm.org/download.html#7.0.1) | [9.2](https://developer.nvidia.com/cuda-92-download-archive) | - <br/> not working due to <br/> the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) <br/>+<br/>[patch](patches/patch_for_clang_7.0.1_bug_38811.zip)*</br> | - <br/> not working due to <br/> the clang's bug [36384](https://bugs.llvm.org/show_bug.cgi?id=36384) |
|
||||
| 7.1.0 | 9.2 (?) | - <br/> LLVM 7.1.0 <br/> is not yet released | - <br/> LLVM 7.1.0 <br/> is not yet released |
|
||||
| [**8.0.0**](http://releases.llvm.org/download.html#8.0.0) | [**10.0**](https://developer.nvidia.com/cuda-10.0-download-archive) | - <br/> not working due to <br/> the clang's bug [38811](https://bugs.llvm.org/show_bug.cgi?id=38811) <br/>+<br/>[patch](patches/patch_for_clang_8.0.0_bug_38811.zip)*</br> | + <br/> **LATEST STABLE RELEASE** |
|
||||
| 8.0.1 | 10.1 (?) | - <br/> LLVM 8.0.1 <br/> is not yet released | - <br/> LLVM 8.0.1 <br/> is not yet released |
|
||||
|
||||
`*` Download the patch and unpack it into your LLVM distributive directory; a few header files will be overwritten; rebuilding of LLVM is not needed.
|
||||
|
||||
@@ -317,7 +317,7 @@ LLVM 7.0.0 - 8.0.0 (with patch*), CUDA 7.5 - 10.0, cudnn-7.0.5.15 - cudnn-7.5.1.
|
||||
|
||||
Build system for the above configurations:
|
||||
|
||||
Python 3.6 (min), cmake 3.12.3 (min), Visual Studio 15.5 2017 (min).
|
||||
Python 3.6 (min), cmake 3.12.3 (min), Visual Studio 15.5.2 2017 - 16.0.4 2019.
|
||||
|
||||
Here is an example of building `hipify-clang` with testing support on `Windows 10` by `Visual Studio 15 2017`:
|
||||
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
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
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
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
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
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
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
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
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
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
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
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
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2017 Advanced Micro Devices, Inc. All rights reserved.
|
||||
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
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015 - present Advanced Micro Devices, Inc. All rights reserved.
|
||||
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
|
||||
@@ -20,7 +20,6 @@ OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
THE SOFTWARE.
|
||||
*/
|
||||
|
||||
|
||||
#include "hip/hip_runtime.h"
|
||||
#include "hip/hip_runtime_api.h"
|
||||
#include <iostream>
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015 - present Advanced Micro Devices, Inc. All rights reserved.
|
||||
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
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015 - present Advanced Micro Devices, Inc. All rights reserved.
|
||||
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
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2017 - present Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2017-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
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2017 - present Advanced Micro Devices, Inc. All rights reserved.
|
||||
Copyright (c) 2017-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
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
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
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
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
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
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
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
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
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
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
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
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
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015 - present Advanced Micro Devices, Inc. All rights reserved.
|
||||
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
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015 - present Advanced Micro Devices, Inc. All rights reserved.
|
||||
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
|
||||
|
||||
+1
-1
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
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
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
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
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
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
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
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
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
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
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
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
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
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
|
||||
@@ -19,6 +19,7 @@ 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
|
||||
|
||||
@@ -1,19 +1,22 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
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 WARRANUMTY OF ANY KIND, EXPRESS OR
|
||||
IMPLIED, INUMCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
FITNUMESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANUMY CLAIM, DAMAGES OR OTHER
|
||||
LIABILITY, WHETHER INUM AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
OUT OF OR INUM CONUMECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
|
||||
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.
|
||||
*/
|
||||
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
/*
|
||||
Copyright (c) 2015-2016 Advanced Micro Devices, Inc. All rights reserved.
|
||||
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
|
||||
|
||||
+1
-1
@@ -37,7 +37,7 @@ using namespace std;
|
||||
|
||||
__global__ void matrixTranspose_static_shared(float* out, float* in,
|
||||
const int width) {
|
||||
// CHECK-NOT: HIP_DYNAMIC_SHARED(int, sharedMem);
|
||||
// CHECK-NOT: HIP_DYNAMIC_SHARED(float, sharedMem);
|
||||
// CHECK: __shared__ float sharedMem[WIDTH * WIDTH];
|
||||
__shared__ float sharedMem[WIDTH * WIDTH];
|
||||
|
||||
|
||||
+225
@@ -0,0 +1,225 @@
|
||||
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
|
||||
/*
|
||||
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 <iostream>
|
||||
// CHECK: #include <hip/hip_runtime.h>
|
||||
#include <cuda.h>
|
||||
#include <assert.h>
|
||||
#define WIDTH 32
|
||||
|
||||
#define NUM (WIDTH * WIDTH)
|
||||
|
||||
#define THREADS_PER_BLOCK_X 4
|
||||
#define THREADS_PER_BLOCK_Y 4
|
||||
#define THREADS_PER_BLOCK_Z 1
|
||||
|
||||
using namespace std;
|
||||
|
||||
#define KNRM "\x1B[0m"
|
||||
#define KRED "\x1B[31m"
|
||||
|
||||
#define failed(...) \
|
||||
printf("%serror: ", KRED); \
|
||||
printf(__VA_ARGS__); \
|
||||
printf("\n"); \
|
||||
printf("error: TEST FAILED\n%s", KNRM); \
|
||||
abort();
|
||||
|
||||
// CHECK: hipError_t localError = error;
|
||||
// CHECK: if ((localError != hipSuccess) && (localError != hipErrorPeerAccessAlreadyEnabled) &&
|
||||
// CHECK: (localError != hipErrorPeerAccessNotEnabled )) {
|
||||
// CHECK: printf("%serror: '%s'(%d) from %s at %s:%d%s\n", KRED, hipGetErrorString(localError),
|
||||
#define CUDACHECK(error) \
|
||||
{ \
|
||||
cudaError_t localError = error; \
|
||||
if ((localError != cudaSuccess) && (localError != cudaErrorPeerAccessAlreadyEnabled) && \
|
||||
(localError != cudaErrorPeerAccessNotEnabled )) { \
|
||||
printf("%serror: '%s'(%d) from %s at %s:%d%s\n", KRED, cudaGetErrorString(localError), \
|
||||
localError, #error, __FILE__, __LINE__, KNRM); \
|
||||
failed("API returned error code."); \
|
||||
} \
|
||||
}
|
||||
|
||||
void checkPeer2PeerSupport() {
|
||||
int gpuCount;
|
||||
int canAccessPeer;
|
||||
// CHECK: CUDACHECK(hipGetDeviceCount(&gpuCount));
|
||||
CUDACHECK(cudaGetDeviceCount(&gpuCount));
|
||||
for (int currentGpu = 0; currentGpu < gpuCount; currentGpu++) {
|
||||
// CHECK: CUDACHECK(hipSetDevice(currentGpu));
|
||||
CUDACHECK(cudaSetDevice(currentGpu));
|
||||
for (int peerGpu = 0; peerGpu < currentGpu; peerGpu++) {
|
||||
if (currentGpu != peerGpu) {
|
||||
// CHECK: CUDACHECK(hipDeviceCanAccessPeer(&canAccessPeer, currentGpu, peerGpu));
|
||||
CUDACHECK(cudaDeviceCanAccessPeer(&canAccessPeer, currentGpu, peerGpu));
|
||||
printf("currentGpu#%d canAccessPeer: peerGpu#%d=%d\n", currentGpu, peerGpu,
|
||||
canAccessPeer);
|
||||
}
|
||||
// CHECK: CUDACHECK(hipSetDevice(peerGpu));
|
||||
CUDACHECK(cudaSetDevice(peerGpu));
|
||||
// CHECK: CUDACHECK(hipDeviceReset());
|
||||
CUDACHECK(cudaDeviceReset());
|
||||
}
|
||||
// CHECK: CUDACHECK(hipSetDevice(currentGpu));
|
||||
CUDACHECK(cudaSetDevice(currentGpu));
|
||||
// CHECK: CUDACHECK(hipDeviceReset());
|
||||
CUDACHECK(cudaDeviceReset());
|
||||
}
|
||||
}
|
||||
|
||||
void enablePeer2Peer(int currentGpu, int peerGpu) {
|
||||
int canAccessPeer;
|
||||
// Must be on a multi-gpu system:
|
||||
assert(currentGpu != peerGpu);
|
||||
// CHECK: CUDACHECK(hipSetDevice(currentGpu));
|
||||
CUDACHECK(cudaSetDevice(currentGpu));
|
||||
// CHECK: hipDeviceCanAccessPeer(&canAccessPeer, currentGpu, peerGpu);
|
||||
cudaDeviceCanAccessPeer(&canAccessPeer, currentGpu, peerGpu);
|
||||
if (canAccessPeer == 1) {
|
||||
// CHECK: CUDACHECK(hipDeviceEnablePeerAccess(peerGpu, 0));
|
||||
CUDACHECK(cudaDeviceEnablePeerAccess(peerGpu, 0));
|
||||
} else
|
||||
printf("peer2peer transfer not possible between the selected gpu devices");
|
||||
}
|
||||
|
||||
void disablePeer2Peer(int currentGpu, int peerGpu) {
|
||||
int canAccessPeer;
|
||||
// Must be on a multi-gpu system:
|
||||
assert(currentGpu != peerGpu);
|
||||
// CHECK: CUDACHECK(hipSetDevice(currentGpu));
|
||||
CUDACHECK(cudaSetDevice(currentGpu));
|
||||
// CHECK: hipDeviceCanAccessPeer(&canAccessPeer, currentGpu, peerGpu);
|
||||
cudaDeviceCanAccessPeer(&canAccessPeer, currentGpu, peerGpu);
|
||||
if (canAccessPeer == 1) {
|
||||
// CHECK: CUDACHECK(hipDeviceDisablePeerAccess(peerGpu));
|
||||
CUDACHECK(cudaDeviceDisablePeerAccess(peerGpu));
|
||||
} else
|
||||
printf("peer2peer disable not required");
|
||||
}
|
||||
|
||||
__global__ void matrixTranspose_static_shared(float* out, float* in,
|
||||
const int width) {
|
||||
// CHECK-NOT: HIP_DYNAMIC_SHARED(float, sharedMem);
|
||||
// CHECK: __shared__ float sharedMem[WIDTH * WIDTH];
|
||||
__shared__ float sharedMem[WIDTH * WIDTH];
|
||||
int x = blockDim.x * blockIdx.x + threadIdx.x;
|
||||
int y = blockDim.y * blockIdx.y + threadIdx.y;
|
||||
sharedMem[y * width + x] = in[x * width + y];
|
||||
__syncthreads();
|
||||
out[y * width + x] = sharedMem[y * width + x];
|
||||
}
|
||||
|
||||
__global__ void matrixTranspose_dynamic_shared(float* out, float* in,
|
||||
const int width) {
|
||||
// declare dynamic shared memory
|
||||
// CHECK-NOT: extern __shared__
|
||||
// CHECK: HIP_DYNAMIC_SHARED(float, sharedMem)
|
||||
extern __shared__ float sharedMem[];
|
||||
int x = blockDim.x * blockIdx.x + threadIdx.x;
|
||||
int y = blockDim.y * blockIdx.y + threadIdx.y;
|
||||
sharedMem[y * width + x] = in[x * width + y];
|
||||
__syncthreads();
|
||||
out[y * width + x] = sharedMem[y * width + x];
|
||||
}
|
||||
|
||||
int main() {
|
||||
checkPeer2PeerSupport();
|
||||
int gpuCount;
|
||||
int currentGpu, peerGpu;
|
||||
// CHECK: CUDACHECK(hipGetDeviceCount(&gpuCount));
|
||||
CUDACHECK(cudaGetDeviceCount(&gpuCount));
|
||||
if (gpuCount < 2) {
|
||||
printf("Peer2Peer application requires atleast 2 gpu devices");
|
||||
return 0;
|
||||
}
|
||||
currentGpu = 0;
|
||||
peerGpu = (currentGpu + 1);
|
||||
printf("currentGpu=%d peerGpu=%d (Total no. of gpu = %d)\n", currentGpu, peerGpu, gpuCount);
|
||||
float *data[2], *TransposeMatrix[2], *gpuTransposeMatrix[2], *randArray;
|
||||
int width = WIDTH;
|
||||
randArray = (float*)malloc(NUM * sizeof(float));
|
||||
for (int i = 0; i < NUM; i++) {
|
||||
randArray[i] = (float)i * 1.0f;
|
||||
}
|
||||
enablePeer2Peer(currentGpu, peerGpu);
|
||||
// CHECK: CUDACHECK(hipSetDevice(currentGpu));
|
||||
CUDACHECK(cudaSetDevice(currentGpu));
|
||||
TransposeMatrix[0] = (float*)malloc(NUM * sizeof(float));
|
||||
// CHECK: hipMalloc((void**)&gpuTransposeMatrix[0], NUM * sizeof(float));
|
||||
cudaMalloc((void**)&gpuTransposeMatrix[0], NUM * sizeof(float));
|
||||
// CHECK: hipMalloc((void**)&data[0], NUM * sizeof(float));
|
||||
cudaMalloc((void**)&data[0], NUM * sizeof(float));
|
||||
// CHECK: hipMemcpy(data[0], randArray, NUM * sizeof(float), hipMemcpyHostToDevice);
|
||||
cudaMemcpy(data[0], randArray, NUM * sizeof(float), cudaMemcpyHostToDevice);
|
||||
dim3 dimGrid(WIDTH / THREADS_PER_BLOCK_X, WIDTH / THREADS_PER_BLOCK_Y);
|
||||
dim3 dimBlock(THREADS_PER_BLOCK_X, THREADS_PER_BLOCK_Y);
|
||||
// CHECK: hipLaunchKernelGGL(matrixTranspose_static_shared, dim3(dimGrid), dim3(dimBlock), 0, 0, gpuTransposeMatrix[0], data[0], width);
|
||||
matrixTranspose_static_shared <<<dimGrid, dimBlock>>>(gpuTransposeMatrix[0], data[0], width);
|
||||
// CHECK: CUDACHECK(hipSetDevice(peerGpu));
|
||||
CUDACHECK(cudaSetDevice(peerGpu));
|
||||
TransposeMatrix[1] = (float*)malloc(NUM * sizeof(float));
|
||||
// CHECK: hipMalloc((void**)&gpuTransposeMatrix[1], NUM * sizeof(float));
|
||||
cudaMalloc((void**)&gpuTransposeMatrix[1], NUM * sizeof(float));
|
||||
// CHECK: hipMalloc((void**)&data[1], NUM * sizeof(float));
|
||||
cudaMalloc((void**)&data[1], NUM * sizeof(float));
|
||||
// CHECK: hipMemcpy(data[1], gpuTransposeMatrix[0], NUM * sizeof(float), hipMemcpyDeviceToDevice);
|
||||
cudaMemcpy(data[1], gpuTransposeMatrix[0], NUM * sizeof(float), cudaMemcpyDeviceToDevice);
|
||||
// CHECK: hipLaunchKernelGGL(matrixTranspose_dynamic_shared, dim3(dimGrid), dim3(dimBlock), sizeof(float) * WIDTH * WIDTH, 0, gpuTransposeMatrix[1], data[1], width);
|
||||
matrixTranspose_dynamic_shared <<<dimGrid, dimBlock, sizeof(float) * WIDTH * WIDTH>>>(gpuTransposeMatrix[1], data[1], width);
|
||||
// CHECK: hipMemcpy(TransposeMatrix[1], gpuTransposeMatrix[1], NUM * sizeof(float), hipMemcpyDeviceToHost);
|
||||
cudaMemcpy(TransposeMatrix[1], gpuTransposeMatrix[1], NUM * sizeof(float), cudaMemcpyDeviceToHost);
|
||||
// CHECK: hipDeviceSynchronize();
|
||||
cudaDeviceSynchronize();
|
||||
disablePeer2Peer(currentGpu, peerGpu);
|
||||
// verify the results
|
||||
int errors = 0;
|
||||
double eps = 1.0E-6;
|
||||
for (int i = 0; i < NUM; i++) {
|
||||
if (std::abs(randArray[i] - TransposeMatrix[1][i]) > eps) {
|
||||
printf("%d cpu: %f gpu peered data %f\n", i, randArray[i], TransposeMatrix[1][i]);
|
||||
errors++;
|
||||
}
|
||||
}
|
||||
if (errors != 0) {
|
||||
printf("FAILED: %d errors\n", errors);
|
||||
} else {
|
||||
printf("Peer2Peer PASSED!\n");
|
||||
}
|
||||
free(randArray);
|
||||
for (int i = 0; i < 2; i++) {
|
||||
// CHECK: hipFree(data[i]);
|
||||
cudaFree(data[i]);
|
||||
// CHECK: hipFree(gpuTransposeMatrix[i]);
|
||||
cudaFree(gpuTransposeMatrix[i]);
|
||||
free(TransposeMatrix[i]);
|
||||
}
|
||||
// CHECK: CUDACHECK(hipSetDevice(peerGpu));
|
||||
CUDACHECK(cudaSetDevice(peerGpu));
|
||||
// CHECK: CUDACHECK(hipDeviceReset());
|
||||
CUDACHECK(cudaDeviceReset());
|
||||
// CHECK: CUDACHECK(hipSetDevice(currentGpu));
|
||||
CUDACHECK(cudaSetDevice(currentGpu));
|
||||
// CHECK: CUDACHECK(hipDeviceReset());
|
||||
CUDACHECK(cudaDeviceReset());
|
||||
return 0;
|
||||
}
|
||||
Ссылка в новой задаче
Block a user