AI/rocm-systems
2
0
Förgrening 0
Kod Ärenden Pull-förfrågningar Actions Packages Projekt Släpp Wiki Aktiviteter

Commit-Graf

  • 28f87f7d2e This switches HIP from its currently convoluted macro + pfe based dispatch mechanism to a more natural one partially based on the existing module API. The basic idea is that HCC will always correctly emit __global__ functions: as empty-bodied stubs, on host, and as kernels, on device. It then becomes trivial to obtain the mangled name on host, at dispatch, from the function's address, and then to use the mangled name to retrieve the kernel. This should address all problems stemming from serialisation, dubious mismatches due to the manufactured functor, macro-isms et al. It also immediately enables support for generalised globals as a consequence of that being available in the module API. Finally, it will make debug much easier, since the actual names of the __global__ functions will automatically be used in traces etc. One detail is that due to how dispatch works now (hipLaunchKernel and hipLaunchKernelGGL are themselves variadic function templates which deduce the function type of the callee), in certain cases it may be necesssary to insert explicit casts to ensure that the variadic argument list selects a viable overload - this can be observed in some unit tests. Eventually we may be able to remove this limitation, but for now it does not appear terribly onerous. The code is not extremely HIPpie, nor is it fully optimised, but rather is intended as a starting point for the HIP team to make its own. Alex Voicu 2017-11-01 15:09:59 +00:00
  • 70a41e7dac This switches HIP from its currently convoluted macro + pfe based dispatch mechanism to a more natural one partially based on the existing module API. The basic idea is that HCC will always correctly emit __global__ functions: as empty-bodied stubs, on host, and as kernels, on device. It then becomes trivial to obtain the mangled name on host, at dispatch, from the function's address, and then to use the mangled name to retrieve the kernel. This should address all problems stemming from serialisation, dubious mismatches due to the manufactured functor, macro-isms et al. It also immediately enables support for generalised globals as a consequence of that being available in the module API. Finally, it will make debug much easier, since the actual names of the __global__ functions will automatically be used in traces etc. One detail is that due to how dispatch works now (hipLaunchKernel and hipLaunchKernelGGL are themselves variadic function templates which deduce the function type of the callee), in certain cases it may be necesssary to insert explicit casts to ensure that the variadic argument list selects a viable overload - this can be observed in some unit tests. Eventually we may be able to remove this limitation, but for now it does not appear terribly onerous. The code is not extremely HIPpie, nor is it fully optimised, but rather is intended as a starting point for the HIP team to make its own. Alex Voicu 2017-11-01 15:09:59 +00:00
  • c2482d1255 This switches HIP from its currently convoluted macro + pfe based dispatch mechanism to a more natural one partially based on the existing module API. The basic idea is that HCC will always correctly emit __global__ functions: as empty-bodied stubs, on host, and as kernels, on device. It then becomes trivial to obtain the mangled name on host, at dispatch, from the function's address, and then to use the mangled name to retrieve the kernel. This should address all problems stemming from serialisation, dubious mismatches due to the manufactured functor, macro-isms et al. It also immediately enables support for generalised globals as a consequence of that being available in the module API. Finally, it will make debug much easier, since the actual names of the __global__ functions will automatically be used in traces etc. One detail is that due to how dispatch works now (hipLaunchKernel and hipLaunchKernelGGL are themselves variadic function templates which deduce the function type of the callee), in certain cases it may be necesssary to insert explicit casts to ensure that the variadic argument list selects a viable overload - this can be observed in some unit tests. Eventually we may be able to remove this limitation, but for now it does not appear terribly onerous. The code is not extremely HIPpie, nor is it fully optimised, but rather is intended as a starting point for the HIP team to make its own. Alex Voicu 2017-11-01 15:09:59 +00:00
  • 6f25720c66 Merge 'master' into 'amd-master' Jenkins 2017-11-01 04:42:40 -05:00
  • 4401c8c390 Merge 'master' into 'amd-master' Jenkins 2017-11-01 04:42:40 -05:00
  • 267f2040f4 Merge 'master' into 'amd-master' Jenkins 2017-11-01 04:42:40 -05:00
  • 7d55fc385b Merge 'master' into 'amd-master' Jenkins 2017-11-01 04:42:40 -05:00
  • f1764ece0c P4 to Git Change 1477155 by johtaylo@johtaylo-jtincrementor2-increment on 2017/11/01 03:00:04 foreman 2017-11-01 03:05:32 -04:00
  • b3181d4b6c P4 to Git Change 1477155 by johtaylo@johtaylo-jtincrementor2-increment on 2017/11/01 03:00:04 foreman 2017-11-01 03:05:32 -04:00
  • f7ea0599f8 P4 to Git Change 1476932 by kzhuravl@kzhuravl-fiji-ocllc on 2017/10/31 16:10:43 foreman 2017-10-31 16:19:51 -04:00
  • 5c24331d20 P4 to Git Change 1476932 by kzhuravl@kzhuravl-fiji-ocllc on 2017/10/31 16:10:43 foreman 2017-10-31 16:19:51 -04:00
  • 3698a97a23 P4 to Git Change 1476932 by kzhuravl@kzhuravl-fiji-ocllc on 2017/10/31 16:10:43 foreman 2017-10-31 16:19:51 -04:00
  • d52e0bf530 P4 to Git Change 1476932 by kzhuravl@kzhuravl-fiji-ocllc on 2017/10/31 16:10:43 foreman 2017-10-31 16:19:51 -04:00
  • a068301408 Add asic id for gfx906 on emulator Shaoyun Liu 2017-10-31 14:47:03 -04:00
  • 55fc06dac3 Add asic id for gfx906 on emulator Shaoyun Liu 2017-10-31 14:47:03 -04:00
  • 2094b92a3e Centos/RHEL - remove usage of constexpr since libc++ doesn't enable ctor for constexpr pair in C++11 Siu Chi Chan 2017-10-31 18:16:12 +00:00
  • c1a4ff7b20 Centos/RHEL - remove usage of constexpr since libc++ doesn't enable ctor for constexpr pair in C++11 Siu Chi Chan 2017-10-31 18:16:12 +00:00
  • d938703f9b Centos/RHEL - remove usage of constexpr since libc++ doesn't enable ctor for constexpr pair in C++11 Siu Chi Chan 2017-10-31 18:16:12 +00:00
  • 99d32a195f Centos/RHEL - remove usage of constexpr since libc++ doesn't enable ctor for constexpr pair in C++11 Siu Chi Chan 2017-10-31 18:16:12 +00:00
  • b04f17be8f use libc++ on centos 7 and rhel 7.4 Siu Chi Chan 2017-10-31 18:14:07 +00:00
  • cb454d5366 use libc++ on centos 7 and rhel 7.4 Siu Chi Chan 2017-10-31 18:14:07 +00:00
  • b6265074fa use libc++ on centos 7 and rhel 7.4 Siu Chi Chan 2017-10-31 18:14:07 +00:00
  • cd6996614f use libc++ on centos 7 and rhel 7.4 Siu Chi Chan 2017-10-31 18:14:07 +00:00
  • c1639486c6 Merge pull request #197 from bensander/update_coherency_tests Maneesh Gupta 2017-10-31 17:26:50 +05:30
  • a6458616f5 Merge pull request #197 from bensander/update_coherency_tests Maneesh Gupta 2017-10-31 17:26:50 +05:30
  • e6e90e9cfa Merge pull request #197 from bensander/update_coherency_tests Maneesh Gupta 2017-10-31 17:26:50 +05:30
  • f27c2c1715 Merge pull request #197 from bensander/update_coherency_tests Maneesh Gupta 2017-10-31 17:26:50 +05:30
  • 6b6d28479d Merge pull request #241 from ROCm-Developer-Tools/multi_host Maneesh Gupta 2017-10-31 16:35:12 +05:30
  • fc5a86d595 Merge pull request #241 from ROCm-Developer-Tools/multi_host Maneesh Gupta 2017-10-31 16:35:12 +05:30
  • 07b52ec043 Merge pull request #241 from ROCm-Developer-Tools/multi_host Maneesh Gupta 2017-10-31 16:35:12 +05:30
  • 4d85b6ab29 Merge pull request #241 from ROCm-Developer-Tools/multi_host Maneesh Gupta 2017-10-31 16:35:12 +05:30
  • 25389f4fcc Merge pull request #242 from bensander/time_hip_apis Maneesh Gupta 2017-10-31 16:33:46 +05:30
  • 5e1a721a05 Merge pull request #242 from bensander/time_hip_apis Maneesh Gupta 2017-10-31 16:33:46 +05:30
  • b93e9ae1b7 Merge pull request #242 from bensander/time_hip_apis Maneesh Gupta 2017-10-31 16:33:46 +05:30
  • 7026b73905 Merge pull request #242 from bensander/time_hip_apis Maneesh Gupta 2017-10-31 16:33:46 +05:30
  • e5639605da Merge 'master' into 'amd-master' Jenkins 2017-10-31 04:10:41 -05:00
  • 151a5f7f64 Merge 'master' into 'amd-master' Jenkins 2017-10-31 04:10:41 -05:00
  • 77093c4ae2 Merge 'master' into 'amd-master' Jenkins 2017-10-31 04:10:41 -05:00
  • d526239184 Merge 'master' into 'amd-master' Jenkins 2017-10-31 04:10:41 -05:00
  • cb47089c17 Make HostQueue::queue_count_ a portable atomic type. Sean Keely 2017-10-31 02:34:47 -05:00
  • d93f92f42d Make HostQueue::queue_count_ a portable atomic type. Sean Keely 2017-10-31 02:34:47 -05:00
  • df9feaa10c P4 to Git Change 1476622 by johtaylo@johtaylo-jtincrementor2-increment on 2017/10/31 03:00:04 foreman 2017-10-31 03:08:53 -04:00
  • a3857339e3 P4 to Git Change 1476622 by johtaylo@johtaylo-jtincrementor2-increment on 2017/10/31 03:00:04 foreman 2017-10-31 03:08:53 -04:00
  • 2c80f84e37 P4 to Git Change 1476496 by gandryey@gera-w8 on 2017/10/30 17:36:16 foreman 2017-10-30 17:44:53 -04:00
  • bea710d2e5 P4 to Git Change 1476496 by gandryey@gera-w8 on 2017/10/30 17:36:16 foreman 2017-10-30 17:44:53 -04:00
  • 4df6b371ab Add ns-level timer for HIP API routines Ben Sander 2017-10-30 20:18:41 +00:00
  • dc7d993a02 Add ns-level timer for HIP API routines Ben Sander 2017-10-30 20:18:41 +00:00
  • e88ef63bc8 Add ns-level timer for HIP API routines Ben Sander 2017-10-30 20:18:41 +00:00
  • 7e908bdec8 Add ns-level timer for HIP API routines Ben Sander 2017-10-30 20:18:41 +00:00
  • c949da237a Inital codes to remove x86_64 dependency in HIP source build Wen-Heng (Jack) Chung 2017-10-30 14:02:50 -05:00
  • b98dede49a Inital codes to remove x86_64 dependency in HIP source build Wen-Heng (Jack) Chung 2017-10-30 14:02:50 -05:00
  • afb75b0b2c Inital codes to remove x86_64 dependency in HIP source build Wen-Heng (Jack) Chung 2017-10-30 14:02:50 -05:00
  • 92fb244841 Inital codes to remove x86_64 dependency in HIP source build Wen-Heng (Jack) Chung 2017-10-30 14:02:50 -05:00
  • 44b9be069e Update README.md Gregory Stoner 2017-10-30 15:11:36 -05:00
  • 2948425072 Update README.md Gregory Stoner 2017-10-30 15:11:36 -05:00
  • 583c976672 Update README.md Gregory Stoner 2017-10-30 15:11:36 -05:00
  • 1c8838ecef Update README.md Gregory Stoner 2017-10-30 15:09:42 -05:00
  • 0fa8f6bdf4 Update README.md Gregory Stoner 2017-10-30 15:09:42 -05:00
  • df6210e389 Update README.md Gregory Stoner 2017-10-30 15:09:42 -05:00
  • 04f03d9aff Rename license.txt to License.txt Gregory Stoner 2017-10-30 15:05:33 -05:00
  • 31ace0d733 Rename license.txt to License.txt Gregory Stoner 2017-10-30 15:05:33 -05:00
  • 48e40cc418 Rename license.txt to License.txt Gregory Stoner 2017-10-30 15:05:33 -05:00
  • fe2680611e Create license.txt Gregory Stoner 2017-10-30 15:04:53 -05:00
  • e8d77cf34a Create license.txt Gregory Stoner 2017-10-30 15:04:53 -05:00
  • a669f562a7 Create license.txt Gregory Stoner 2017-10-30 15:04:53 -05:00
  • a4ee15064b P4 to Git Change 1476364 by gandryey@gera-w8 on 2017/10/30 14:40:10 foreman 2017-10-30 15:00:41 -04:00
  • c2105d8f16 P4 to Git Change 1476364 by gandryey@gera-w8 on 2017/10/30 14:40:10 foreman 2017-10-30 15:00:41 -04:00
  • f4e0940c7e Update README.md cfreehill 2017-10-30 12:15:12 -05:00
  • 4e12dfaa73 Update README.md cfreehill 2017-10-30 12:15:12 -05:00
  • 716a8a0661 Update README.md cfreehill 2017-10-30 12:15:12 -05:00
  • e83de0d06b first commit Chris Freehill 2017-10-30 12:12:40 -05:00
  • 1a7ef1a6dd first commit Chris Freehill 2017-10-30 12:12:40 -05:00
  • e948ee5e00 first commit Chris Freehill 2017-10-30 12:12:40 -05:00
  • cb107b30f1 Merge pull request #222 from bensander/fix_device_prop Ben Sander 2017-10-30 17:58:48 +01:00
  • 7f96edc89a Merge pull request #222 from bensander/fix_device_prop Ben Sander 2017-10-30 17:58:48 +01:00
  • fd5b9cb650 Merge pull request #222 from bensander/fix_device_prop Ben Sander 2017-10-30 17:58:48 +01:00
  • 7e8b39fc96 Merge pull request #222 from bensander/fix_device_prop Ben Sander 2017-10-30 17:58:48 +01:00
  • 51ee7807db Merge pull request #222 from bensander/fix_device_prop Ben Sander 2017-10-30 17:58:48 +01:00
  • 2e8ec71e40 Merge pull request #222 from bensander/fix_device_prop Ben Sander 2017-10-30 17:58:48 +01:00
  • ec29721b16 Check for null copyEngine before looking at peers. Ben Sander 2017-10-30 16:58:03 +00:00
  • 4c7b2be1c2 Check for null copyEngine before looking at peers. Ben Sander 2017-10-30 16:58:03 +00:00
  • ccb8b441a7 Check for null copyEngine before looking at peers. Ben Sander 2017-10-30 16:58:03 +00:00
  • d610f16c47 Check for null copyEngine before looking at peers. Ben Sander 2017-10-30 16:58:03 +00:00
  • a6d17854f8 Merge pull request #226 from scchan/add_printf3 Ben Sander 2017-10-30 17:08:18 +01:00
  • 4debd30c3d Merge pull request #226 from scchan/add_printf3 Ben Sander 2017-10-30 17:08:18 +01:00
  • 8242c45eda Merge pull request #226 from scchan/add_printf3 Ben Sander 2017-10-30 17:08:18 +01:00
  • f8843ae415 Merge pull request #226 from scchan/add_printf3 Ben Sander 2017-10-30 17:08:18 +01:00
  • 239d505d41 Merge 'master' into 'amd-master' Jenkins 2017-10-30 04:10:43 -05:00
  • a8859cd3e1 Merge 'master' into 'amd-master' Jenkins 2017-10-30 04:10:43 -05:00
  • 41da2a7706 Merge 'master' into 'amd-master' Jenkins 2017-10-30 04:10:43 -05:00
  • c62c05aadb Merge 'master' into 'amd-master' Jenkins 2017-10-30 04:10:43 -05:00
  • 9486eadfbe [newTests]Adding tests for device APIs Phaneendr-kumar Lanka 2017-10-30 14:34:24 +05:30
  • 71434383a3 [newTests]Adding tests for device APIs Phaneendr-kumar Lanka 2017-10-30 14:34:24 +05:30
  • a36c3c6d8b [newTests]Adding tests for device APIs Phaneendr-kumar Lanka 2017-10-30 14:34:24 +05:30
  • 511de63bcb [newTests]Adding tests for device APIs Phaneendr-kumar Lanka 2017-10-30 14:34:24 +05:30
  • 3e9a0561c0 Add APIs to support debugging vm fault Qingchuan Shi 2017-09-29 12:45:24 -04:00
  • ce6aee01ed Add APIs to support debugging vm fault Qingchuan Shi 2017-09-29 12:45:24 -04:00
  • 976ed873c0 P4 to Git Change 1475971 by johtaylo@johtaylo-jtincrementor2-increment on 2017/10/28 03:00:06 foreman 2017-10-28 03:09:38 -04:00
  • f0ac05164a P4 to Git Change 1475971 by johtaylo@johtaylo-jtincrementor2-increment on 2017/10/28 03:00:06 foreman 2017-10-28 03:09:38 -04:00
  • f0ae7312f5 P4 to Git Change 1475897 by gandryey@gera-w8 on 2017/10/27 16:28:25 foreman 2017-10-27 16:36:05 -04:00