06:16 Duke``: Hi. I have an hybrid intel+amd gfx laptop. Both cards work in opengl with DRI_PRIME, however vulkaninfo is only listing the intel gfx. The AMD card is [Radeon HD 8670A/8670M/8690M / R5 M330 / M430 / Radeon 520 Mobile]: is it supported by mesa vulkan? if yes, is there something to do to enable the discrete card for vulkan? (i.e. like DRI_PRIME=1 for opengl)
09:40 linkmauve: Duke``, this sounds like a quite old card, check out the minimum requirements for radv.
09:41 linkmauve: IIRC it requires the amdgpu kernel driver, check out that you have that.
09:42 Duke``: [drm] amdgpu kernel modesetting enabled
09:42 Duke``: but yes, it's an old dell latitude e6440 (haswell), that's why I'm not sure for vulkan support
09:43 Duke``: but I don't know where to check for wether this amd card is supported or not
09:45 HdkR: radv minspec should be GCN
09:45 HdkR: And that's GCN 1.0
09:49 HdkR: could do `RADV_DEBUG=startup` and see if it prints something about the card
09:49 HdkR: while running vulkaninfo with that environment variable
09:55 HdkR: Your software could just be outdated though
10:59 FireBurn|Home: Duke: As someone with a similar setup your vulkaninfo should list both cards in the same output
11:00 FireBurn|Home: https://dpaste.com/G3JEVXB6K
11:01 FireBurn|Home: To force a specific driver you can point directly to it's json file like this:
11:01 FireBurn|Home: VK_ICD_FILENAMES=/usr/share/vulkan/icd.d/radeon_icd.x86_64.json vulkaninfo
11:01 FireBurn|Home: Your distro might have it in a differrent place
11:04 FireBurn: Does your distro install radv as default?
11:05 FireBurn: For me it's installed in /usr/lib64/libvulkan_radeon.so
13:55 EdB: how does it usually work for shader in mesa. Does images sizes get stored somewhere in a register on the card ?
14:19 maelcum: EdB: presumably it needs to be known to the hardware, for things like repeat mode, interpolation and mipmapping, which are at least partially done by fixed-function hardware. if it's programmed into registers or into a descriptor block in vram that is loaded into a hardware block on demand... i would guess the former.
14:19 maelcum: eh, i mean the latter :>
14:20 maelcum: similar to how paging works on a cpu
14:57 SneakySnake: Can anyone with radeon hardware confirm that the cube at the "platforming cube" section of this demo is rendering incorrectly? https://stianj.com/what-are-you-syncing-about/
14:58 SneakySnake: This is what it looks like with llvmpipe vs the hardware renderer https://i.imgur.com/8vJZNzc.png
14:58 SneakySnake: This is on an RX 580
16:35 FireBurn: I think it looked ok
16:35 FireBurn: M395X here on chromium
16:36 FireBurn: So technically its swiftshader converting to OpenGL
16:40 hanetzer: SneakySnake: currently over ssh, so really can't, but once I get a chance I'll do it :)
18:02 airlied: EdB: depends on hw
18:03 airlied: radeonsi has sizes in descriptor
18:17 EdB: airlied: is that upload with ac_add_arg ?
18:22 airlied: n9
18:22 airlied: no
18:23 airlied: its usually loaded from user sgprs indirectly
18:23 airlied: but ive no idea how rocm does it
18:24 airlied: in theory you could put it in the input args
18:24 airlied: or you put an index in there
18:25 airlied: and load it from a table pointed to by the dispatch
18:32 EdB: roc is full of #define to create those function :/
18:49 airlied: EdB: https://github.com/RadeonOpenCompute/ROCm-Device-Libs/blob/amd-stg-open/opencl/src/image/imwrap.cl would need to be ported to libclc
18:50 EdB: yes
18:50 airlied: for LLVM path
18:50 EdB: I'm trying to figuring out what it realy does
18:50 airlied: I assume it produces the newer amdgcn intrinsic
18:51 EdB: yes
18:51 EdB: it seems to read a buffer at position 10 to get width
18:52 EdB: but I would like to know wich one
18:53 airlied: also I'm not sure how code object v2 vs v3 works here
18:53 airlied: which makes it tricker to figure out the dispatch path
18:53 EdB: mesa is still v2
18:53 airlied: "Image and sample handles created by the ROCm runtime are 64-bit addresses of a
18:53 airlied: hardware 32-byte V# and 48 byte S# object respectively. In order to support the
18:53 airlied: HSA ``query_sampler`` operations two extra dwords are used to store the HSA BRIG
18:54 airlied: enumeration values for the queries that are not trivially deducible from the S#
18:54 airlied: representation."
18:54 airlied: but that is all v3 I think
18:57 EdB: I don't want to just CC the stuff. I would like to understand it. But really, I'm very not familiar with all this stuff
18:58 airlied: so it's likely the input args have a 64-bit global addr in them
18:58 airlied: which points to 32-byte + 48 byte descriptors
18:58 airlied:has looked a couple of times at fixing up images on the llvm path and gave up
18:59 EdB: but, but, you are my hero !
19:02 airlied: hehe, I was only mildly interested in CL back then and realised I probably was taking on too much :-P
19:04 EdB: when we launch a kernel, we put handle into the input args. That loks like what you descibe, to me
19:05 EdB: but, not sure about the content of those addr.
19:05 airlied: yeah so I think the current clover code sticks an index in there for r600
19:07 airlied: since for r600 an index make sense
19:07 EdB: for the printf stuff, I added a buffer id that I could read back from device with '__builtin_amdgcn_implicitarg_ptr()'
19:34 yshui`: Is there something I can read on the details of freesync/vrr? I don't quite understand how it works, and why it's a problem for compositors to have vrr enabled.
20:00 EdB: ok, so now if I try to figure out what Roc produce I have : implicit declaration of function '__ockl_image_width_2D' is invalid in OpenCL
20:05 EdB: ah, I forgot to take the okcl part
20:18 EdB: airlied: llvm gives me : bitcast %opencl.image2d_ro_t.0 addrspace(4)* %0 to i32 addrspace(4)*
20:18 EdB: getelementptr inbounds i32, i32 addrspace(4)* %2, i64 10
20:18 EdB: load i32, i32 addrspace(4)* %3, align 4, !tbaa !8
20:19 EdB: there is %2 =before the bitcast
20:20 EdB: so that match what you describe earlier
20:52 Duke``: <FireBurn|Home> https://dpaste.com/G3JEVXB6K <-- thx for info (I now know what it should output). When I force only radeon ICD, there's a vulkan initialization error. I'm using Ubuntu 20.04.
21:12 FireBurn: Duke: What does the error say? Is the .so file installed correctly?
21:17 Duke``: /build/vulkan-tools-KEbD_A/vulkan-tools-1.2.131.1+dfsg1/vulkaninfo/vulkaninfo.h:477: failed with ERROR_INITIALIZATION_FAILED
21:20 Duke``: the .so file seems installed correctly (/usr/lib/x86_64-linux-gnu/libvulkan_radeon.so)
21:21 Duke``: I have to go, I'll investigate that another day, thx anyway
21:24 FireBurn: Is there anything in your dmesg when that happens?