Commit Graph

54851 Commits

Author SHA1 Message Date
foreman 96d41be9e3 P4 to Git Change 1475542 by johtaylo@johtaylo-jtincrementor2-increment on 2017/10/27 03:00:05
SWDEV-2 - Change OpenCL version number from 2528 to 2529.

Affected files ...

... //depot/stg/opencl/drivers/opencl/runtime/utils/versions.hpp#2276 edit


[ROCm/clr commit: 9f9daf779e]
2017-10-27 03:07:50 -04:00
foreman 9f9daf779e P4 to Git Change 1475542 by johtaylo@johtaylo-jtincrementor2-increment on 2017/10/27 03:00:05
SWDEV-2 - Change OpenCL version number from 2528 to 2529.

Affected files ...

... //depot/stg/opencl/drivers/opencl/runtime/utils/versions.hpp#2276 edit
2017-10-27 03:07:50 -04:00
Siu Chi Chan bec0ff43e0 Merge remote-tracking branch 'origin/master' into HEAD
[ROCm/clr commit: cf98d0e37d]
2017-10-27 01:18:28 -04:00
Siu Chi Chan cf98d0e37d Merge remote-tracking branch 'origin/master' into HEAD 2017-10-27 01:18:28 -04:00
Siu Chi Chan a9789ddcda Merge remote-tracking branch 'origin/master' into HEAD 2017-10-27 01:18:28 -04:00
Ben Sander 731c1afea6 Merge pull request #198 from AlexVlx/feature_support_globals_for_module_api
Feature support globals for module api
2017-10-27 01:53:34 +02:00
Ben Sander b0aa15ee5f Merge pull request #198 from AlexVlx/feature_support_globals_for_module_api
Feature support globals for module api

[ROCm/clr commit: 8a64feef61]
2017-10-27 01:53:34 +02:00
Ben Sander 8a64feef61 Merge pull request #198 from AlexVlx/feature_support_globals_for_module_api
Feature support globals for module api
2017-10-27 01:53:34 +02:00
Ben Sander f288f24e95 Merge pull request #198 from AlexVlx/feature_support_globals_for_module_api
Feature support globals for module api
2017-10-27 01:53:34 +02:00
foreman 7f7c37e5a3 P4 to Git Change 1475427 by lmoriche@lmoriche_opencl_dev2 on 2017/10/26 18:17:19
SWDEV-126884 - OCL 2.1 Platform APIs
	- Add OpenCL 2.2 clSetProgramSpecializationConstant and clSetProgramReleaseCallback stubs.

Affected files ...

... //depot/stg/opencl/drivers/opencl/api/opencl/amdocl/cl_icd.cpp#31 edit
... //depot/stg/opencl/drivers/opencl/api/opencl/amdocl/cl_program.cpp#44 edit


[ROCm/clr commit: 41082b16ca]
2017-10-26 18:28:20 -04:00
foreman 41082b16ca P4 to Git Change 1475427 by lmoriche@lmoriche_opencl_dev2 on 2017/10/26 18:17:19
SWDEV-126884 - OCL 2.1 Platform APIs
	- Add OpenCL 2.2 clSetProgramSpecializationConstant and clSetProgramReleaseCallback stubs.

Affected files ...

... //depot/stg/opencl/drivers/opencl/api/opencl/amdocl/cl_icd.cpp#31 edit
... //depot/stg/opencl/drivers/opencl/api/opencl/amdocl/cl_program.cpp#44 edit
2017-10-26 18:28:20 -04:00
foreman aac9311992 P4 to Git Change 1475393 by lmoriche@lmoriche_opencl_dev2 on 2017/10/26 17:01:11
SWDEV-126884 - OCL 2.1 Platform APIs
	- Implement clGetKernelSubGroupInfo, clCloneKernel, clGetHostTimer and clGetHostAndDeviceTimer

Affected files ...

... //depot/stg/opencl/drivers/opencl/api/opencl/amdocl/cl_device.cpp#68 edit
... //depot/stg/opencl/drivers/opencl/api/opencl/amdocl/cl_execute.cpp#25 edit
... //depot/stg/opencl/drivers/opencl/api/opencl/amdocl/cl_program.cpp#43 edit
... //depot/stg/opencl/drivers/opencl/runtime/platform/kernel.cpp#24 edit
... //depot/stg/opencl/drivers/opencl/runtime/platform/kernel.hpp#19 edit


[ROCm/clr commit: c9b973e307]
2017-10-26 17:08:01 -04:00
foreman c9b973e307 P4 to Git Change 1475393 by lmoriche@lmoriche_opencl_dev2 on 2017/10/26 17:01:11
SWDEV-126884 - OCL 2.1 Platform APIs
	- Implement clGetKernelSubGroupInfo, clCloneKernel, clGetHostTimer and clGetHostAndDeviceTimer

Affected files ...

... //depot/stg/opencl/drivers/opencl/api/opencl/amdocl/cl_device.cpp#68 edit
... //depot/stg/opencl/drivers/opencl/api/opencl/amdocl/cl_execute.cpp#25 edit
... //depot/stg/opencl/drivers/opencl/api/opencl/amdocl/cl_program.cpp#43 edit
... //depot/stg/opencl/drivers/opencl/runtime/platform/kernel.cpp#24 edit
... //depot/stg/opencl/drivers/opencl/runtime/platform/kernel.hpp#19 edit
2017-10-26 17:08:01 -04:00
foreman 89ff1a059d P4 to Git Change 1475393 by lmoriche@lmoriche_opencl_dev2 on 2017/10/26 17:01:11
SWDEV-126884 - OCL 2.1 Platform APIs
	- Implement clGetKernelSubGroupInfo, clCloneKernel, clGetHostTimer and clGetHostAndDeviceTimer

Affected files ...

... //depot/stg/opencl/drivers/opencl/api/opencl/amdocl/cl_device.cpp#68 edit
... //depot/stg/opencl/drivers/opencl/api/opencl/amdocl/cl_execute.cpp#25 edit
... //depot/stg/opencl/drivers/opencl/api/opencl/amdocl/cl_program.cpp#43 edit
... //depot/stg/opencl/drivers/opencl/runtime/platform/kernel.cpp#24 edit
... //depot/stg/opencl/drivers/opencl/runtime/platform/kernel.hpp#19 edit


[ROCm/clr commit: 687e1fdd9c]
2017-10-26 17:08:01 -04:00
foreman 687e1fdd9c P4 to Git Change 1475393 by lmoriche@lmoriche_opencl_dev2 on 2017/10/26 17:01:11
SWDEV-126884 - OCL 2.1 Platform APIs
	- Implement clGetKernelSubGroupInfo, clCloneKernel, clGetHostTimer and clGetHostAndDeviceTimer

Affected files ...

... //depot/stg/opencl/drivers/opencl/api/opencl/amdocl/cl_device.cpp#68 edit
... //depot/stg/opencl/drivers/opencl/api/opencl/amdocl/cl_execute.cpp#25 edit
... //depot/stg/opencl/drivers/opencl/api/opencl/amdocl/cl_program.cpp#43 edit
... //depot/stg/opencl/drivers/opencl/runtime/platform/kernel.cpp#24 edit
... //depot/stg/opencl/drivers/opencl/runtime/platform/kernel.hpp#19 edit
2017-10-26 17:08:01 -04:00
Ben Sander 1058f074f3 Merge pull request #218 from ChrisKitching/nodiscard
Add [[nodiscard]] attribute to hipError_t in C++17 mode

[ROCm/clr commit: 9713dbb6f6]
2017-10-26 22:48:54 +02:00
Ben Sander 9713dbb6f6 Merge pull request #218 from ChrisKitching/nodiscard
Add [[nodiscard]] attribute to hipError_t in C++17 mode
2017-10-26 22:48:54 +02:00
Ben Sander e97f675397 Merge pull request #218 from ChrisKitching/nodiscard
Add [[nodiscard]] attribute to hipError_t in C++17 mode
2017-10-26 22:48:54 +02:00
Ben Sander 7ce81b1573 Merge pull request #223 from bensander/2x_bidir
Use 2X for bidir memory bandwidth calc

[ROCm/clr commit: f76bcf4045]
2017-10-26 21:49:06 +02:00
Ben Sander f76bcf4045 Merge pull request #223 from bensander/2x_bidir
Use 2X for bidir memory bandwidth calc
2017-10-26 21:49:06 +02:00
Ben Sander 772fe865fc Merge pull request #223 from bensander/2x_bidir
Use 2X for bidir memory bandwidth calc
2017-10-26 21:49:06 +02:00
Ben Sander fb169b4de6 Fix bug with peer-to-peer combined with context API
- Store context inside the tracker rather than using int deviceID that
  was always mapped to primary context
- IsPeerWatcher now based on device IDs rather than specific peers.


[ROCm/clr commit: a417241507]
2017-10-26 19:44:22 +00:00
Ben Sander a417241507 Fix bug with peer-to-peer combined with context API
- Store context inside the tracker rather than using int deviceID that
  was always mapped to primary context
- IsPeerWatcher now based on device IDs rather than specific peers.
2017-10-26 19:44:22 +00:00
Ben Sander 7d30f32332 Fix bug with peer-to-peer combined with context API
- Store context inside the tracker rather than using int deviceID that
  was always mapped to primary context
- IsPeerWatcher now based on device IDs rather than specific peers.
2017-10-26 19:44:22 +00:00
Aditya Atluri a4af82519a Enhance debug for copy pointers
- show more pointer tracking fields
- show pointer info before and after "tailoring'


[ROCm/clr commit: 698721be34]
2017-10-26 19:44:22 +00:00
Aditya Atluri 698721be34 Enhance debug for copy pointers
- show more pointer tracking fields
- show pointer info before and after "tailoring'
2017-10-26 19:44:22 +00:00
Aditya Atluri 5d646d0fe3 Enhance debug for copy pointers
- show more pointer tracking fields
- show pointer info before and after "tailoring'
2017-10-26 19:44:22 +00:00
Philip Yang 3501b2f40d Fix double free on fork after hsaKmtCloseKFD
Child process hsaKmtOpenKFD() call must re-initialize global variables
copied from parent process. This includes close all file handles, free
dynamically malloc buf. Double free issue is because destroy_device_
debugging_memory() free the memory in parent process hsaKmtCloseKFD()
but don't reset it to null pointer. As a result, child process free it
again. kfd_fd is closed in parent process but don't reset to 0, so
child process close it again.

Fix: reset kfd_fd to 0 after close, reset is_device_debugged pointer to 0
after free



Change-Id: I421b3decbcaa4111298b8e599aa16940d851a58c
Signed-off-by: Philip Yang <Philip.Yang@amd.com>
2017-10-26 15:36:15 -04:00
Philip Yang ac80bac82a Fix double free on fork after hsaKmtCloseKFD
Child process hsaKmtOpenKFD() call must re-initialize global variables
copied from parent process. This includes close all file handles, free
dynamically malloc buf. Double free issue is because destroy_device_
debugging_memory() free the memory in parent process hsaKmtCloseKFD()
but don't reset it to null pointer. As a result, child process free it
again. kfd_fd is closed in parent process but don't reset to 0, so
child process close it again.

Fix: reset kfd_fd to 0 after close, reset is_device_debugged pointer to 0
after free



Change-Id: I421b3decbcaa4111298b8e599aa16940d851a58c
Signed-off-by: Philip Yang <Philip.Yang@amd.com>


[ROCm/ROCR-Runtime commit: 3501b2f40d]
2017-10-26 15:36:15 -04:00
Evgeny Mankov d0216119e2 Merge branch 'master' of https://github.com/ROCm-Developer-Tools/HIP
[ROCm/clr commit: d98d287863]
2017-10-26 20:33:18 +03:00
Evgeny Mankov d98d287863 Merge branch 'master' of https://github.com/ROCm-Developer-Tools/HIP 2017-10-26 20:33:18 +03:00
Evgeny Mankov bb2b416337 Merge branch 'master' of https://github.com/ROCm-Developer-Tools/HIP 2017-10-26 20:33:18 +03:00
foreman 1653cb9849 P4 to Git Change 1475222 by gandryey@gera-lnx-rcf on 2017/10/26 12:45:19
SWDEV-79445 - Don't enable HostMemoryDirectAccess flag for persistent memory

Affected files ...

... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocmemory.cpp#27 edit


[ROCm/clr commit: 66e0d09af0]
2017-10-26 12:54:46 -04:00
foreman 66e0d09af0 P4 to Git Change 1475222 by gandryey@gera-lnx-rcf on 2017/10/26 12:45:19
SWDEV-79445 - Don't enable HostMemoryDirectAccess flag for persistent memory

Affected files ...

... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocmemory.cpp#27 edit
2017-10-26 12:54:46 -04:00
Chris Kitching 7360326705 Greatly enhance handling of macros in kernel launches
All but the most contrived use of macros is now properly handled -
have a look at the new testcases this commit adds. You can have
macros in kernel calls, macros spanning chunks of your arguments,
the call, call parameters, or callee can all be macros or
partially macros.


[ROCm/clr commit: 6491c2c3eb]
2017-10-26 17:28:46 +01:00
Chris Kitching 6491c2c3eb Greatly enhance handling of macros in kernel launches
All but the most contrived use of macros is now properly handled -
have a look at the new testcases this commit adds. You can have
macros in kernel calls, macros spanning chunks of your arguments,
the call, call parameters, or callee can all be macros or
partially macros.
2017-10-26 17:28:46 +01:00
Chris Kitching 094b2b9b05 Greatly enhance handling of macros in kernel launches
All but the most contrived use of macros is now properly handled -
have a look at the new testcases this commit adds. You can have
macros in kernel calls, macros spanning chunks of your arguments,
the call, call parameters, or callee can all be macros or
partially macros.
2017-10-26 17:28:46 +01:00
Chris Kitching 58428d739c Simplify how kernel launch expressions get translated
It seems like there was a lot of machinery here that is no longer
needed now we have hipLaunchKernelGGL (which doesn't require us
to insert an extra argument into kernel functions). We no longer
need to waste cycles scanning the AST for callees.

We can literally just do "Take the callee expression, and dump
it into the first argument of hipLaunchKernelGGL()".


[ROCm/clr commit: a35d30e0b7]
2017-10-26 17:28:30 +01:00
Chris Kitching a35d30e0b7 Simplify how kernel launch expressions get translated
It seems like there was a lot of machinery here that is no longer
needed now we have hipLaunchKernelGGL (which doesn't require us
to insert an extra argument into kernel functions). We no longer
need to waste cycles scanning the AST for callees.

We can literally just do "Take the callee expression, and dump
it into the first argument of hipLaunchKernelGGL()".
2017-10-26 17:28:30 +01:00
Chris Kitching eff86d975b Simplify how kernel launch expressions get translated
It seems like there was a lot of machinery here that is no longer
needed now we have hipLaunchKernelGGL (which doesn't require us
to insert an extra argument into kernel functions). We no longer
need to waste cycles scanning the AST for callees.

We can literally just do "Take the callee expression, and dump
it into the first argument of hipLaunchKernelGGL()".
2017-10-26 17:28:30 +01:00
Chris Kitching 74af29d66a Deduplicate preprocessor code
There's three functions here that all do the same thing...

There was also logic that looks for numeric literals and works
backwards to find the macro name from which they are expanded.
I previously introduced code that rewrites macro references at
expand-time in the `MacroExpands` callback, so that code is no
longer doing anything useful.


[ROCm/clr commit: 1ef68090ae]
2017-10-26 17:28:30 +01:00
Chris Kitching 1ef68090ae Deduplicate preprocessor code
There's three functions here that all do the same thing...

There was also logic that looks for numeric literals and works
backwards to find the macro name from which they are expanded.
I previously introduced code that rewrites macro references at
expand-time in the `MacroExpands` callback, so that code is no
longer doing anything useful.
2017-10-26 17:28:30 +01:00
Chris Kitching fd911e1839 Deduplicate preprocessor code
There's three functions here that all do the same thing...

There was also logic that looks for numeric literals and works
backwards to find the macro name from which they are expanded.
I previously introduced code that rewrites macro references at
expand-time in the `MacroExpands` callback, so that code is no
longer doing anything useful.
2017-10-26 17:28:30 +01:00
Chris Kitching b25b17b6b3 Rewrite _all_ CUDA macro identifiers in the preprocessor
Calls to macros that were themselves CUDA API calls were often
being missed - this applies the identifier transform to macro
names at the callsites, too.


[ROCm/clr commit: 30e7e7d919]
2017-10-26 17:27:56 +01:00
Chris Kitching 30e7e7d919 Rewrite _all_ CUDA macro identifiers in the preprocessor
Calls to macros that were themselves CUDA API calls were often
being missed - this applies the identifier transform to macro
names at the callsites, too.
2017-10-26 17:27:56 +01:00
Chris Kitching d1e26b2e7e Rewrite _all_ CUDA macro identifiers in the preprocessor
Calls to macros that were themselves CUDA API calls were often
being missed - this applies the identifier transform to macro
names at the callsites, too.
2017-10-26 17:27:56 +01:00
Chris Kitching cf50b4f97a Don't special-case source locations for calls in macros
The source location for a call that's inside a macro body will,
by default, point into the macro definition itself. The original
logic was causing macro invocations to be overwritten, as I
explain here:
https://github.com/ROCm-Developer-Tools/HIP/issues/207#issuecomment-337521851

The existing PPCallbacks code is correctly rewriting macro
definitions, so the practical effect of this change is that AST
rewrites on code that's expanded from macros are no-ops.

It might be a performance optimisation to put a short-circiut at
the top of the AST callbacks to abort when faced with code that
was expanded from macros.

It might yet prove wise to do absolutely everything at lex-time...


[ROCm/clr commit: f7e65c5334]
2017-10-26 17:26:37 +01:00
Chris Kitching f7e65c5334 Don't special-case source locations for calls in macros
The source location for a call that's inside a macro body will,
by default, point into the macro definition itself. The original
logic was causing macro invocations to be overwritten, as I
explain here:
https://github.com/ROCm-Developer-Tools/HIP/issues/207#issuecomment-337521851

The existing PPCallbacks code is correctly rewriting macro
definitions, so the practical effect of this change is that AST
rewrites on code that's expanded from macros are no-ops.

It might be a performance optimisation to put a short-circiut at
the top of the AST callbacks to abort when faced with code that
was expanded from macros.

It might yet prove wise to do absolutely everything at lex-time...
2017-10-26 17:26:37 +01:00
Chris Kitching 4a794ed8c0 Don't special-case source locations for calls in macros
The source location for a call that's inside a macro body will,
by default, point into the macro definition itself. The original
logic was causing macro invocations to be overwritten, as I
explain here:
https://github.com/ROCm-Developer-Tools/HIP/issues/207#issuecomment-337521851

The existing PPCallbacks code is correctly rewriting macro
definitions, so the practical effect of this change is that AST
rewrites on code that's expanded from macros are no-ops.

It might be a performance optimisation to put a short-circiut at
the top of the AST callbacks to abort when faced with code that
was expanded from macros.

It might yet prove wise to do absolutely everything at lex-time...
2017-10-26 17:26:37 +01:00
Chris Kitching af81909cab Prefer early-return to deep nesting
A chain of 7 closing braces is never a great sign :D

In the process it became apparant that the unsupported flag
was being silently ignored, causing users to be left with cuda
API calls in their programs with no warning given. This has been
rectified for consistency.


[ROCm/clr commit: c1f4612176]
2017-10-26 17:26:37 +01:00