This adds the missing compilation flags to hip::device so it can compile with hip-clang compiler.
Change-Id: Ie2b30ea606bfca385a0e84ae03ee0a8d828ad16a
[ROCm/clr commit: 631f82ab82]
Even though the runtime and driver texture object API is one to one, the structs used by these APIs are not. See hipResourceDesc vs HIP_RESOURCE_DESC differences.
These differences are not trivial and most likely won't be able to handled by hipify, so we need new API entry points.
Change-Id: Id4bcb1ad0ae15378dbdb5a2ed07e5ea30f320082
[ROCm/clr commit: a09fadecf2]
This patch is a workaround to support user pitch for hipMemcpy{2D/3D}.
Historically OpenCL didn't support pitch with clEnqueueFillBuffer(), so neither did we in VDI. Adding it now will be slightly nontrivial, since the fill kernel and runtime in many places will need to be modified.
As a temporary workaround for cases when pitch > width, we can just enqueue a fill for each row separately. This implementation is slow, but it satisfies the correctness criteria.
Change-Id: Idfeca349288b51d6ff84a7cf001fb63c6a66818a
[ROCm/clr commit: 6e0722a5d0]
Fixes SWDEV-207362,
The output file name should not contribute to picking up the right flags for the compiler. This fix solves issues when the output has conflicting extensions which confuses hipcc to treat them as the source files and add the required flags for them.
PS: Output file refers to the file followed by -o
Change-Id: I1095966c11143ad73e81fabc35b4e9de5d3afada
Example: hipcc test.o -o test.hip will add the flags for .hip compilation ignoring the fact that it is an output file
[ROCm/clr commit: 36f24a40e5]
- Use symbol value as the qeury key. Compared to the symbol name, the
symbol value is more robust as developers may use unqualified or
qualified identifiers. It also removes the mangling and/or demangling
requirement for the runtime API.
Change-Id: I9d4259f3842612c7cc98551269fc2092d8b5c19e
[ROCm/clr commit: a14695d4eb]
This is cherry-picked from PR#1947 that was committed to the
github repo. It allows printf to work with hip-clang and HCC
runtime.
Change-Id: I754753250ea1e694cf3441722e2d4c9d25fa75bc
[ROCm/clr commit: c014c733c6]
Adjust the origin of the copy if the user passes a pointer that wasn't allocated by the runtime.
Change-Id: I0aeb20195ed730857a461a53f537626ec2573fd1
[ROCm/clr commit: 86c969fcea]
The only requirment from hipMallocPitch() is that the returned pitch is aligned to the HW image pitch alignment. There is no restriction on the size of the allocation, since the memory might not be used for images.
Change-Id: I97438e5fe4012ca4721b14b85f514dbac803c17c
[ROCm/clr commit: 111897dbea]
Fix issues of missing kernel function symbols and missing argument list via
using __hipRegister* functions.
Then the following tests can pass,
directed_tests/runtimeApi/module/hipFuncGetAttributes
directed_tests/runtimeApi/module/hipExtLaunchMultiKernelMultiDevice
directed_tests/gcc/LaunchKernel
Change-Id: I52135b61e8283eb4f9f10f77895151e4e55418d9
[ROCm/clr commit: 4a113bb669]
Instead of using the sampler field force_degamma to perform sRGB->linear conversion during pixel sampling, we use an appropriate image format instead. The overhead of this is having to create an image view when creating a texture object from an array.
Change-Id: I1ca368c312c1fd4b6f784a3a1b35b5eeb28070ff
[ROCm/clr commit: 23211f05d6]
This avoids the use of extra casts when obtaining a texture object handle.
Change-Id: I42df22bdad0ab9ac6c33cb8b282dee65fe7cfd6e
[ROCm/clr commit: 6cfbe19160]
The struct consists of a union - only the active object should be read.
Change-Id: I1c40965b61518acd91a2dcbae92a015ac9be346a
[ROCm/clr commit: 1cacb5c9b7]
This also adds declarations of all the missing texture APIs.
hipTexRefSet*() functions need to take a textureReference as a ptr for type erasure to work. Runtime has been modified to accomodate this.
This change only applies to VDI.
Change-Id: Icf43cc5bd44dfc2c39084b7fe56d5a793bf7319f
[ROCm/clr commit: 428b56e411]
We don't program the numChannels and format members (these are HCC specific), so printing these will only display garbage.
Change-Id: I83dc8be9a3cae2659c64f4594d07c05330d2dd14
[ROCm/clr commit: d662abaf55]
All we have to do is align the ptr to HW requirments an if it's not zero, then return the offset to the user.
We currently don't have anywhere to store this offset, so hipGetTextureAlignmentOffset() will still always return 0.
Change-Id: If31998127d99a2a3222a026d88249519d6102505
[ROCm/clr commit: 9823ccf5b0]
The hipExtModuleLaunchKernel and hipModuleLoadDataMultiThreaded tests keeps randomly failing on Jenkins.
Change-Id: I87e5d54fb7429c14ff1dcecb20e03a7816670fae
[ROCm/clr commit: 684cdd4a18]
Change-Id: If419d2fad490d0ed50eb1315af809fc1deda1ce3
SWDEV-227875: Add a lock in streams to lock when the callback is call so we make sure things aren't moving forward in the stream
[ROCm/clr commit: 9f81498bb5]
Change-Id: I1f107fc8a5c586cd571f0280ed8716c5f89d25b7
SWDEV-227875: Need to add a dummy marker in case the stream is empty.
[ROCm/clr commit: 31df9b358d]
What Cuda refers to "linear texture memory" is the OpenCL equivalent of CL_MEM_OBJECT_IMAGE1D_BUFFER. For these types of allocations we should create a typed buffer instead of an image.
Currently there is no check in the texture fetch functions as to what kind of SRD is written into the texture object, so any kind of incorrect programming will cause the TA to hang. Fortunately for us, every one writes correct code :)
Change-Id: I80dab85a992f2c0754ebf303d40ac6b5e045c7c1
[ROCm/clr commit: 78f7954765]
Paul Fultz II
Saleel Kudchadker
Vladislav Sytchenko
Reshabh Sharma
Michael Hong Bin Liao
Christophe Paquot
kjayapra-amd
Anusha Godavarthy Surya
Tao Sang
Sameer Sahasrabuddhe
Payam Ghafari
jujiang
Sarbojit Sarkar