SWDEV-2 - Change OpenCL version number from 2140 to 2141.
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/utils/versions.hpp#1887 edit
[ROCm/clr commit: 251374c4e4]
SWDEV-2 - Change OpenCL version number from 2139 to 2140.
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/utils/versions.hpp#1886 edit
[ROCm/clr commit: 1bf495769d]
SWDEV-2 - Change OpenCL version number from 2138 to 2139.
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/utils/versions.hpp#1885 edit
[ROCm/clr commit: ff7a43814b]
SWDEV-2 - Change OpenCL version number from 2137 to 2138.
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/utils/versions.hpp#1884 edit
[ROCm/clr commit: af51c1687b]
SWDEV-2 - Change OpenCL version number from 2136 to 2137.
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/utils/versions.hpp#1883 edit
[ROCm/clr commit: a0d9fe3c69]
commit 9548493fa754b3bf5c31cbdc2211db1e73e8c07c
Author: Jack Chung <whchung@gmail.com>
Date: Mon May 23 11:57:23 2016 +0800
Rename hipExternShared test to hipDynamicShared
Change-Id: I180d9d539420fb69cfc121eceaa7db9da03483b2
commit 827081f8244a38f010789d556db0c4ff7b6422d8
Author: Jack Chung <whchung@gmail.com>
Date: Mon May 23 11:56:27 2016 +0800
Rename HIP_DECLARE_EXTERN_SHARED to HIP_DYNAMIC_SHARED
Change-Id: I22362d179812ac547e0f11ba4e2bb999050e08ae
commit 4c277228ed41af187739610fa17eab1fb144c947
Author: Jack Chung <whchung@gmail.com>
Date: Thu May 19 17:49:52 2016 +0800
Adopt new interface to get dynamic LDS in hc.hpp
Change-Id: I47b433b714633a4c97df87c40a0b1d3386429a00
commit 5a36117d777064113a528dc47b42e8c8413baa97
Author: Jack Chung <whchung@gmail.com>
Date: Thu May 19 11:29:24 2016 +0800
Add test patterns for regular expression to match "extern __shared__"
These test patterns should better be saved as an individual test case, but I'm
not familiar with HIP test structures so I leave them as comments in hipify as
of now.
Change-Id: I7fee89c89b9e73de2133357a226ec0c769733531
commit 1b26284168c7f5339f63338fd0149bed5d994656
Author: Jack Chung <whchung@gmail.com>
Date: Thu May 19 11:25:23 2016 +0800
Add one HIP unit test to use HIP_DECLARE_EXTERN_SHARED
Change-Id: I4d9907815920693a74ea9d575fe26e7c67636109
commit 77b816ee5972b13d829d5bbcf06fbfd07acea2af
Author: Jack Chung <whchung@gmail.com>
Date: Wed May 18 19:18:59 2016 +0800
Adopt HIP_ prefix for DECLARE_EXTERN_SHARED macro
Change-Id: I555ded16b449b67d2e20904013d86fe1ded6a2be
commit ef0997939c3578a9ae11621bf21c0416f04d2622
Author: Jack Chung <whchung@gmail.com>
Date: Wed May 18 17:42:04 2016 +0800
Modify hipify to support converting extern __shared__ to DECLARE_EXTERN_SHARED macro
Added regular expression to search & replace extern __shared__ declarations to
DECLARE_EXTERN_SHARED macro.
Limitation:
- Won't work if "extern __shared__" is declared at global scope
Sample Usages:
extern __shared__ double foo[];
extern __shared__ unsigned int foo[];
extern volatile __shared__ double foo[];
extern volatile __shared__ unsigned int sdata[];
extern __shared__ volatile unsigned int sdata[];
extern __shared__ T s[];
extern __shared__ T::type s[];
extern __shared__ blah<T>::type s[];
extern __shared__ typename mapper<Float>::type s_data[];
extern __attribute__((used)) __shared__ typename mapper<Float>::type s_data[];
Change-Id: I2be0b7039adeddb789f5a2b067d403a43fdc3e26
commit 93ff268724493aedfacdcd5a5aa9a100f4ebaed0
Author: Jack Chung <whchung@gmail.com>
Date: Wed May 18 15:13:09 2016 +0800
Introduce DECLARE_EXTERN_SHARED macro to encapsulate "extern __shared__" decls
Change-Id: I93b2d37c763195b0ca9fd0afee78605a1e3272db
commit cff9c95412de343cc6405158b5acc4f1029267ff
Author: Jack Chung <whchung@gmail.com>
Date: Wed May 18 12:53:54 2016 +0800
Add __get_dynamic_groupbaseptr() to point to dynamic LDS
Change-Id: I97b548d8a691488057617c551a8f331cad7afc77
Change-Id: I84e7875b76fa1f59e860e19c93bd4209cdd1fd2c
[ROCm/clr commit: 65448e74ed]
SWDEV-2 - Change OpenCL version number from 2135 to 2136.
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/utils/versions.hpp#1882 edit
[ROCm/clr commit: 6494e2f8e7]
SWDEV-2 - Change OpenCL version number from 2134 to 2135.
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/utils/versions.hpp#1881 edit
[ROCm/clr commit: ebdad97cb2]
SWDEV-92067 - Put back full CL#1275679.
TC has added '//depot/stg/opencl/drivers/opencl/appprofiles/...' to the OpenCL Brahma Builds' client-spec.
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/device/appprofile.cpp#14 edit
[ROCm/clr commit: a90de181e6]
SWDEV-2 - Change OpenCL version number from 2133 to 2134.
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/utils/versions.hpp#1880 edit
[ROCm/clr commit: 9dffb49ac9]
SWDEV-76911 - Set output pipeline of multi-shader compilation in the same way driver sets. Fill output shader pointers only if they will be generated based on the input shaders, otherwise NULL.
ReviewBoardURL = http://dxreview.amd.com/r/20059/
Affected files ...
... //depot/stg/sc/Src/Dev/TestEngine.cpp#512 edit
[ROCm/clr commit: 610ee40c74]
SWDEV-76911 - Set output pipeline of multi-shader compilation in the same way driver sets. Fill output shader pointers only if they will be generated based on the input shaders, otherwise NULL.
ReviewBoardURL = http://dxreview.amd.com/r/20059/
Affected files ...
... //depot/stg/sc/Src/Dev/TestEngine.cpp#512 edit
[ROCm/clr commit: 436ca8ad4d]
SWDEV-92067 - Support app profiles on Brahma.
The Brahma kernel driver does not support parsing the app profile blob files, and does not have ADL. OGL has a library that reads and parses the blob file directly. This CL imports the OGL library into OCL runtime.
Affected files ...
... //depot/stg/opencl/drivers/opencl/appprofiles/brahma/apl.cpp#1 add
... //depot/stg/opencl/drivers/opencl/appprofiles/brahma/apl.hpp#1 add
... //depot/stg/opencl/drivers/opencl/appprofiles/brahma/aplexport.cpp#1 add
... //depot/stg/opencl/drivers/opencl/appprofiles/brahma/structanddefines.hpp#1 add
... //depot/stg/opencl/drivers/opencl/runtime/build/Makefile.runtime#63 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/appprofile.cpp#12 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/appprofile.hpp#9 edit
[ROCm/clr commit: 428dfba879]
SWDEV-2 - Change OpenCL version number from 2132 to 2133.
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/utils/versions.hpp#1879 edit
[ROCm/clr commit: abe1a9f942]
SWDEV-2 - Change OpenCL version number from 2131 to 2132.
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/utils/versions.hpp#1878 edit
[ROCm/clr commit: b90237baa1]
SWDEV-86035 - Add PAL backend to OpenCL
- Fix imagebuffer and SRGB failures in Conformance 2.0. The new PAL is required for all imagebuffer formats.
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/paldevice.cpp#8 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palresource.cpp#5 edit
[ROCm/clr commit: bb3678a080]
SWDEV-2 - Change OpenCL version number from 2130 to 2131.
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/utils/versions.hpp#1877 edit
[ROCm/clr commit: c87a217105]
SWDEV-2 - Change OpenCL version number from 2129 to 2130.
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/utils/versions.hpp#1876 edit
[ROCm/clr commit: 63118b2740]
SWDEV-2 - Change OpenCL version number from 2128 to 2129.
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/utils/versions.hpp#1875 edit
[ROCm/clr commit: 9456725e52]
SWDEV-2 - Change OpenCL version number from 2127 to 2128.
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/utils/versions.hpp#1874 edit
[ROCm/clr commit: acee9c2e2c]
SWDEV-95095 - Bristol Ridge parts report their OpenCL device ID name as "Carrizo"
- Detect the ASIC revision for the name swap with Bristo Ridgel instead of Carrizo
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/device/gpu/gpucompiler.cpp#154 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/gpu/gpudevice.cpp#548 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/gpu/gslbe/src/rt/GSLDevice.cpp#168 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/gpu/gslbe/src/rt/backend.h#13 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palcompiler.cpp#2 edit
[ROCm/clr commit: 9cbc946a1d]
SWDEV-79445 - Fix a TC regression after CL#1273440
- Add the both conditions to the restrictions: sorce and destination alignment
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/device/gpu/gslbe/src/rt/GSLDevice.cpp#167 edit
[ROCm/clr commit: 23ec683890]
SWDEV-2 - Change OpenCL version number from 2126 to 2127.
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/utils/versions.hpp#1873 edit
[ROCm/clr commit: d9cea8af16]
SWDEV-94881 - [OpenCL]increasing the size limit of image buffer to 2^27 which is for the moment at 65536
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/device/gpu/gpudefs.hpp#136 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/paldefs.hpp#3 edit
[ROCm/clr commit: 624f331317]
SWDEV-93545 - HSA HLC: target option interface between complib and BE and denorm control refactoring
Global variables and associated options to control target GPU for optimizations and fp32 denorm support are removed.
Instead standard llvm -mcpu=<cpu> is used to pass chip family name and fp32 denorm is turned into a subtarget feature.
Subtarget feature can be set for llc as standard -mattr=+fp32-denormals and corresponding code to pass feature string
to the BE is added to the compiler lib, mimicing what we used to have for AMDIL.
Device name HSAIL metadata will now reflect an actual GPU family passed to the HSAIL BE instead of "generic".
Denorm support can be switched on as a feature bit in the target mapping. It is on starting from VI. However, just
switching this bit for a family will not produce denorm supporting code. The option -cl-denorms-are-zero can be used
to override this and runtime passes it for configs where CL_FP_DENORM is not reported.
Currently CL_FP_DENORM is not reported for any device, however it can be changed with AMD_GPU_FORCE_SINGLE_FP_DENORM
environment variable. If set it will be honored only starting from VI as set in the target mapping.
Implemented isFMAFasterThanFMulAndFAdd to handle use of v_fma_f32 on GFX9 instead of a direct chip family check.
Testing: smoke, precheckin
Reviewed by Evgeny Mankov
Affected files ...
... //depot/stg/opencl/drivers/opencl/compiler/lib/backends/common/codegen.cpp#70 edit
... //depot/stg/opencl/drivers/opencl/compiler/lib/backends/common/linker.cpp#146 edit
... //depot/stg/opencl/drivers/opencl/compiler/lib/backends/common/opt_level.cpp#30 edit
... //depot/stg/opencl/drivers/opencl/compiler/lib/utils/v0_8/libUtils.cpp#18 edit
... //depot/stg/opencl/drivers/opencl/compiler/lib/utils/v0_8/libUtils.h#27 edit
... //depot/stg/opencl/drivers/opencl/compiler/lib/utils/v0_8/target_mappings.h#42 edit
... //depot/stg/opencl/drivers/opencl/compiler/lib/utils/v0_8/target_mappings_hsail.h#32 edit
... //depot/stg/opencl/drivers/opencl/compiler/lib/utils/v0_8/target_mappings_hsail64.h#27 edit
... //depot/stg/opencl/drivers/opencl/compiler/llvm/lib/Target/HSAIL/BRIGAsmPrinter.cpp#155 edit
... //depot/stg/opencl/drivers/opencl/compiler/llvm/lib/Target/HSAIL/BRIGAsmPrinter.h#69 edit
... //depot/stg/opencl/drivers/opencl/compiler/llvm/lib/Target/HSAIL/HSAIL.td#11 edit
... //depot/stg/opencl/drivers/opencl/compiler/llvm/lib/Target/HSAIL/HSAILISelDAGToDAG.cpp#71 edit
... //depot/stg/opencl/drivers/opencl/compiler/llvm/lib/Target/HSAIL/HSAILISelLowering.cpp#116 edit
... //depot/stg/opencl/drivers/opencl/compiler/llvm/lib/Target/HSAIL/HSAILISelLowering.h#29 edit
... //depot/stg/opencl/drivers/opencl/compiler/llvm/lib/Target/HSAIL/HSAILInstrInfo.cpp#42 edit
... //depot/stg/opencl/drivers/opencl/compiler/llvm/lib/Target/HSAIL/HSAILInstrInfo.h#18 edit
... //depot/stg/opencl/drivers/opencl/compiler/llvm/lib/Target/HSAIL/HSAILInstructions.td#22 edit
... //depot/stg/opencl/drivers/opencl/compiler/llvm/lib/Target/HSAIL/HSAILKernelManager.cpp#55 edit
... //depot/stg/opencl/drivers/opencl/compiler/llvm/lib/Target/HSAIL/HSAILSubtarget.cpp#14 edit
... //depot/stg/opencl/drivers/opencl/compiler/llvm/lib/Target/HSAIL/HSAILSubtarget.h#15 edit
... //depot/stg/opencl/drivers/opencl/compiler/llvm/lib/Target/HSAIL/HSAILTargetMachine.cpp#57 edit
... //depot/stg/opencl/drivers/opencl/tests/hsa/bin/test_driver.pl#29 edit
... //depot/stg/opencl/drivers/opencl/tests/hsa/tlst/llc_opt.tlst#97 edit
... //depot/stg/opencl/drivers/opencl/tests/hsa/tlst/ocl_features.tlst#58 edit
[ROCm/clr commit: 712c0c4009]
SWDEV-94161 - [Adobe Photoshop] TDR is observed while applying Blur effect on the selected portion of an Image
- Restore back the original condition for SDMA restriction (CL#1195487). The root cause wasn't identified and probably somewhere in GSL, since PAL backend works fine.
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/device/gpu/gslbe/src/rt/GSLDevice.cpp#166 edit
[ROCm/clr commit: 24b23e3f72]
SWDEV-85602 - HSA HLC: serialize linking
Since we do not have build referenced symbol map now during linking we do not have to keep all library modules
in memory during the whole link process. This change serializes linking to save some memory.
Testing: smoke, precheckin
Reviewed by Evgeny Mankov
Affected files ...
... //depot/stg/opencl/drivers/opencl/compiler/lib/backends/common/linker.cpp#145 edit
... //depot/stg/opencl/drivers/opencl/compiler/lib/linker/include/AMDResolveLinker.h#5 edit
... //depot/stg/opencl/drivers/opencl/compiler/lib/linker/lib/AMDResolveLinker.cpp#5 edit
... //depot/stg/opencl/drivers/opencl/compiler/lib/linker/tools/opencl-link/opencl-link.cpp#7 edit
[ROCm/clr commit: 1161118bc3]
SWDEV-2 - Change OpenCL version number from 2125 to 2126.
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/utils/versions.hpp#1872 edit
[ROCm/clr commit: 06ef7ffe6e]
SWDEV-2 - Change OpenCL version number from 2124 to 2125.
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/utils/versions.hpp#1871 edit
[ROCm/clr commit: 112e68f217]
Aditya Atluri
7SK
foreman
Maneesh Gupta
Jack Chung