SWDEV-79445 - OCL generic changes and code clean-up
- Use SDMA staging transfers for data upload if pinning fails. Fixes HIP failure in a test that uses the code segment data for uppload.
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palblit.cpp#26 edit
[ROCm/clr commit: 45d50bb738]
SWDEV-159881 - [OCL][ROCm] Add SVM coarse-grain buffer support with device memory
1. Use the system memory pool for coarse grain allocations on APUs
ReviewBoardURL = http://ocltc.amd.com/reviews/r/15671/diff/
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocmemory.cpp#39 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocvirtual.cpp#64 edit
[ROCm/clr commit: 1bf70582f0]
SWDEV-161959 - [ROCm QA][RAVEN] QCDGPU-S test is having ERROR -61: (clCreateKernel failed) on RAVEN
SWDEV-161983 - [ROCm QA][RAVEN] Cachebench test is failing with CL_INVALID_BUFFER_SIZE issue
SWDEV-161978 - [ROCm QA][RAVEN] PCIeBW is failing on -with error : 61, OpenCL error creating buffer !
SWDEV-161962 - [ROCm QA][RAVEN] rodinia->nw test has ERROR: clCreateBuffer input_item_set (size:67125249) => -61
- We make the global memory size at least 1G byte. This is to avoid issue/regression if sysconf API misbehaves.
ReviewBoardURL = http://ocltc.amd.com/reviews/r/15660/
Tests:
1. ocltst -m oclruntime.so -A oclruntime.exclude - PASS except SVM test (non regression)
2. TeamCity presubmission test (OpenCL) - PASS
3. Run test qcdgpu-s.sh : PASS
4. Run test cachebench-ocl : PASS
5. Run test PCIeBandwidth -c 0 -g 0 : PASS
6. Run test Rodinia/opencl/nw/run : PASS
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocdevice.cpp#98 edit
[ROCm/clr commit: d886bbdb2c]
SWDEV-161959 - [ROCm QA][RAVEN] QCDGPU-S test is having ERROR -61: (clCreateKernel failed) on RAVEN
SWDEV-161983 - [ROCm QA][RAVEN] Cachebench test is failing with CL_INVALID_BUFFER_SIZE issue
SWDEV-161978 - [ROCm QA][RAVEN] PCIeBW is failing on -with error : 61, OpenCL error creating buffer !
SWDEV-161962 - [ROCm QA][RAVEN] rodinia->nw test has ERROR: clCreateBuffer input_item_set (size:67125249) => -61
- Adjust Global memory size as half of the system physical memory size in APU for Rocm
- Similar to current DGPU calculation, environment GPU_SINGLE_ALLOC_PERCENT can be used to adjust max memory allocation size.
ReviewBoardURL = http://ocltc.amd.com/reviews/r/15659/
Tests:
1. ocltst -m oclruntime.so -A oclruntime.exclude - PASS except SVM test (non regression)
2. TeamCity presubmission test (OpenCL) - PASS
http://ocltc.amd.com:8111/viewModification.html?modId=106628&personal=true&init=1&tab=vcsModificationBuilds
3. Run test qcdgpu-s.sh : PASS
4. Run test cachebench-ocl : PASS
5. Run test PCIeBandwidth -c 0 -g 0 : PASS
6. Run test Rodinia/opencl/nw/run : PASS
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocdevice.cpp#97 edit
[ROCm/clr commit: f84d1d2d03]
SWDEV-160930 - Back out CL1594730 as it is failing OCL/LC and HIP tests
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palkernel.cpp#59 edit
[ROCm/clr commit: c6280baa3c]
SWDEV-160930 - SPECworkstation 3 benchmark GPU Compute tests fail
Root cause: Caffe compute benchmark fails within SPECWorkstation app because one of the Caffe's OCL kernel tries to launch a kernel with the local_work_size of 1024 causing the clEnqueueNDRangeKernel API to return CL_INVALID_WORK_GROUP_SIZE (i.e., the maximum allowable number is 256)
Proposed workaround: In order to run a kernel with a local_work_size of 1024, we check the number of used VGPRs in the Kernel and if the Kernel is not using all the available VGPRs we let the Kernel to use 1024 as the local_work_size.
ReviewURLBoard = http://ocltc.amd.com/reviews/r/15638/
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palkernel.cpp#58 edit
[ROCm/clr commit: 1dcac07a7e]
SWDEV-145570 - [HIP] Output Kernel name and mem arguments passed with LOG_LEVEL=3 for PAL and ROCm backends
ReviewBoardURL = http://ocltc.amd.com/reviews/r/15617/diff/
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palvirtual.cpp#120 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocvirtual.cpp#62 edit
[ROCm/clr commit: bc3e3217b6]
SWDEV-160552 - [CQE OCL][DTB][BD][Perf][QR][G] Soft hang is observed with ProRender after few subtests due to base driver.
- Reset the pointer to nullptr if PAL platform was destroyed
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/paldevice.cpp#102 edit
[ROCm/clr commit: 1a61f6b9e5]
SWDEV-160552 - [CQE OCL][DTB][BD][Perf][QR][G] Soft hang is observed with ProRender after few subtests due to base driver.
- Delete PAL platform object if runtime didn't find any supported device.
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/paldevice.cpp#101 edit
[ROCm/clr commit: 769cc7cddb]
SWDEV-145570 - [HIP] Increase the default size for SVM to 16GB for HIP
ReviewBoardURL = http://ocltc.amd.com/reviews/r/15578/diff/
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/paldevice.cpp#100 edit
[ROCm/clr commit: aa99feb153]
SWDEV-155434 - Add SQTT instrumentation tokens for OpenCL dispatches for RGP support
- Switch to the workgroup size report for the dispatch info.
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palvirtual.cpp#118 edit
[ROCm/clr commit: 19968ab78a]
SWDEV-132899 - [OCL][GFX10] removing gfx1001 (target for Navi10 Lite if xnack is enabled)
starting from gfx10, we will use the same target (e.g., gfx1000) no matter if xnack is enabled/disbaled and we pass "-xnack" option to Finalizer if xnack is enabled on the target
ReviewURLBoard = http://ocltc.amd.com/reviews/r/15492/
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/paldefs.hpp#37 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/paldevice.cpp#99 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palprogram.cpp#64 edit
[ROCm/clr commit: 9ddfeef105]
SWDEV-79445 - OCL generic changes and code clean-up
1. In SvmCopyMemoryCommand handle the case when both src and dst pointer don't belong to the SVM space.
ReviewBoardURL = http://ocltc.amd.com/reviews/r/15481/diff/
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/device/gpu/gpuvirtual.cpp#424 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palvirtual.cpp#117 edit
[ROCm/clr commit: ebe91eeace]
SWDEV-79445 - OCL generic changes and code clean-up
- Reset memory dependency if runtime invalidated L1 for the profiling logic workaround. Profiling can be enabled for wave limiter, which could cause L1 invalidation twice.
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palvirtual.cpp#116 edit
[ROCm/clr commit: c84525aa73]
SWDEV-155860 - [OCL][LC] 62 Failures running on OCL/LC Windows drivers
- Conformance test basic->vload_private was failing due to alignment assert check on akc->workitem_private_segment_byte_size.
- Determined needed to add alignment, so added alignment into the header file, and removed assert check (as alignment was done afterwards).
ReviewBoardURL = http://ocltc.amd.com/reviews/r/15451/diff/
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palkernel.cpp#57 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palkernel.hpp#18 edit
[ROCm/clr commit: 12a063779d]
SWDEV-159255 - [CQE OCL][ocltst][WIN] [DTB-Blocker] OCLMemoryInfo[0] a sub-test of ocltst oclruntime module is failed while running whole module and getting pass while running alone this test due to faulty CL#1576247
- Reduce extra size acceptance for the cache look-up, so it will satisfy the test condition.
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palresource.cpp#69 edit
[ROCm/clr commit: 75423e4972]
SWDEV-155438 - Produce RGP Queue Timings chunk for OpenCL RGP files
- Register SDMA queue in order to get SDMA timing. The RGP trace capture with SDMA may cause a HW hang occasionally
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palgpuopen.cpp#6 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palgpuopen.hpp#5 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/palvirtual.cpp#115 edit
[ROCm/clr commit: 8c44442f90]
SWDEV-79445 - OCL generic changes and code clean-up
1. std::bind2nd is deprecated in c++11 and removed in c++17. Use a range based for loop without any functor instead.
ReviewBoardURL = http://ocltc.amd.com/reviews/r/15430/diff/
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocvirtual.cpp#59 edit
[ROCm/clr commit: 7899017d49]
SWDEV-152361 - Fixed writeBuffer, writeBufferRect and copyBufferRect paths for ExternalPhysical memory in GPU backend, so that SDI copy can sync with writeMarker correctly in Linux.
http://ocltc.amd.com/reviews/r/15420/
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/device/gpu/gpublit.cpp#131 edit
[ROCm/clr commit: 509d540374]
SWDEV-154474 - [Task]: TAN Dual Realtime Queue Enhancement For Radeon Rays
- Use the new enum for RT1
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/device/gpu/gpuvirtual.cpp#423 edit
[ROCm/clr commit: 77bd31536a]
SWDEV-158585 - [ROCm] OCLPinnedMemory test is giving VM fault
1. If host ptr is present, use the CopyBuffer path instead of the ReadBuffer path
ReviewBoardURL = http://ocltc.amd.com/reviews/r/15403/diff/
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocvirtual.cpp#58 edit
[ROCm/clr commit: 017d0f3e03]
SWDEV-158730 - [CQE OCL][ocltst][WIN] OCLMemoryInfo[1] a sub-test of ocltst oclruntime module is failed
- Update free memory calculation for APU systems. There is still an issue in GSL with multiple alloc/free passes
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/device/gpu/gpudevice.cpp#595 edit
... //depot/stg/opencl/drivers/opencl/runtime/device/pal/paldevice.cpp#96 edit
[ROCm/clr commit: 0b08fa4a4e]
SWDEV-126897 - Remove un-necessary synchronize() for Device Enqueue in LC/ROCm path.
We don't need it since we have a hsa_signal_wait_acquire() to wait for the scheduler to finish. Moreover, synchronize() sends a barrier packet that could get stuck for some reason.
Affected files ...
... //depot/stg/opencl/drivers/opencl/runtime/device/rocm/rocblit.cpp#25 edit
[ROCm/clr commit: e2b18eda3e]
foreman