I've just recently discovered AMD's equivalent to CUDA's __byte_perm intrinsic; amdgcn_ds_swizzle(Or at least I think its the equivalent of a byte permutation function). My problem is this: CUDA's byte perm takes in two unsigned 32 bit integers, and then permutes that based on the value of the selector argument (supplied as a hex value). However, AMD's swizzle function only takes in one single unsigned 32 bit integer, and one int that's named as "pattern". How do I utilize AMD's Swizzle intrinsic function?
AMD HCC Swizzle Intrinsic
866 views Asked by ligosan At
1
There are 1 answers
Related Questions in C++
- How to immediately apply DISPLAYCONFIG_SCALING display scaling mode with SetDisplayConfig and DISPLAYCONFIG_PATH_TARGET_INFO
- Why can't I use templates members in its specialization?
- How to fix "Access violation executing location" when using GLFW and GLAD
- Dynamic array of structures in C++/ cannot fill a dynamic array of doubles in structure from dynamic array of structures
- How do I apply the interface concept with the base-class in design?
- File refuses to compile std::erase() even if using -std=g++23
- How can I do a successful map when the number of elements to be mapped is not consistent in Thrust C++
- Can std::bit_cast be applied to an empty object?
- Unexpected inter-thread happens-before relationships from relaxed memory ordering
- How i can move element of dynamic vector in argument of function push_back for dynamic vector
- Brick Breaker Ball Bounce
- Thread-safe lock-free min where both operands can change c++
- Watchdog Timer Reset on ESP32 using Webservers
- How to solve compiler error: no matching function for call to 'dmhFS::dmhFS()' in my case?
- Conda CMAKE CXX Compiler error while compiling Pytorch
Related Questions in PARALLEL-PROCESSING
- How to calculate Matrix exponential with Tailor series PARALLEL using MPI c++
- Efficiently processing many small elements of a collection concurrently in Java
- Parallelize filling of Eigen Matrix in C++
- Memory efficient parallel repeated rarefaction with subsequent matrix addition of large data set
- How to publish messages to RabbitMQ by using Multi threading?
- Running a C++ Program with CMake, MPI and OpenCV
- Alternative approach to io.ReadAll to store memory consumption and send a PUT Request with valid data
- Parallelize nested loop with running sum in Fortran
- Can I use parfor within a parfeval in Matlab R2019b and if yes how?
- Parallel testing with cucumber, selenium and junit 5
- Parallel.ForEach vs ActionBlock
- Passing variable to foreach-object -parallel which is with in start-job
- dbatools SQL Functions Not Running In Parallel While SQL Server queries do in Powershell
- How do I run multiple instances of my Powershell function in parallel?
- Joblib.parallel vs concurrent.futures
Related Questions in GPGPU
- OpenCL dynamic parallelism enqueue_kernel() functionality
- Sign a PGP public key using a private key and password, then save the signed key to a file
- Passing arguments to OpenCL kernel, before execution finished
- CUDA kernel for finding the min and max index of values in a 1D array greater than particular threshold
- Cuda __device__ member function with explicit template declaration
- AMD GPU Compute with c++
- Why is webgpu on mac "max binding size" much smaller than reported "max buffer size"?
- Running multiple times a python script from different threads using different gpus
- GPGPU with Radeon Pro VII in Windows
- Pytorch Memory Management Issue
- Perform vector calculation on GPU in C++, regardless of brand
- Reinterpret cast on *shared memory*
- Can I really launch a library kernel (CUkernel) rather than an in-context kernel (CUfunction)?
- How to use shared memory in PyCuda, LogicError: cuModuleLoadDataEx failed: an illegal memory access was encountered
- What (if anything) is this GPU compute or shader pattern called?
Related Questions in GPU
- A deterministic GPU implementation of fused batch-norm backprop, when training is disabled, is not currently available
- What is the parameter for CLI YOLOv8 predict to use Intel GPU?
- Windows 10 TensorFlow cannot detect Nvidia GPU
- Is there a way to profile a CUDA kernel from another CUDA kernel
- Does Unity render invisible material?
- Quantization 4 bit and 8 bit - error in 'quantization_config'
- Pyarrow: ImportError: /lib/x86_64-linux-gnu/libc.so.6: version `GLIBC_2.28' not found
- How to setup SLI on two GTX 560Ti's
- How can I delete a process in CUDA?
- No GPU EC2 instances associated with AWS Batch
- access fan and it's speed, in linux mint on acer predator helios 300
- Why can CPU memory be specified and allocated during instance creation but not GPU memory on the cloud?
- Why do CUDA asynchronous errors occur? (occur on the linux OS)
- Pytorch how to use num_worker>0 for Dataloader when using multiple gpus
- Running PyTorch MPS acceleration on Apple M1, get "Placeholder storage has not been allocated on MPS device!" error, but all seems to be on device
Popular Questions
- How do I undo the most recent local commits in Git?
- How can I remove a specific item from an array in JavaScript?
- How do I delete a Git branch locally and remotely?
- Find all files containing a specific text (string) on Linux?
- How do I revert a Git repository to a previous commit?
- How do I create an HTML button that acts like a link?
- How do I check out a remote Git branch?
- How do I force "git pull" to overwrite local files?
- How do I list all files of a directory?
- How to check whether a string contains a substring in JavaScript?
- How do I redirect to another webpage?
- How can I iterate over rows in a Pandas DataFrame?
- How do I convert a String to an int in Java?
- Does Python have a string 'contains' substring method?
- How do I check if a string contains a specific word?
Trending Questions
- UIImageView Frame Doesn't Reflect Constraints
- Is it possible to use adb commands to click on a view by finding its ID?
- How to create a new web character symbol recognizable by html/javascript?
- Why isn't my CSS3 animation smooth in Google Chrome (but very smooth on other browsers)?
- Heap Gives Page Fault
- Connect ffmpeg to Visual Studio 2008
- Both Object- and ValueAnimator jumps when Duration is set above API LvL 24
- How to avoid default initialization of objects in std::vector?
- second argument of the command line arguments in a format other than char** argv or char* argv[]
- How to improve efficiency of algorithm which generates next lexicographic permutation?
- Navigating to the another actvity app getting crash in android
- How to read the particular message format in android and store in sqlite database?
- Resetting inventory status after order is cancelled
- Efficiently compute powers of X in SSE/AVX
- Insert into an external database using ajax and php : POST 500 (Internal Server Error)
ds_swizzle and __byte_perm do are a little bit different. One permutes a whole register across lanes and the later permutes any four bytes from two 32-bit regs.
AMD's ds_swizzle_b32 GCN instruction is actually swapping values with other lanes. You specify the 32-bit register in the lane you want to read and the 32-bit register you want to place it in. There is also a hard-coded value that specifies how these are to be swapped. A great explanation of ds_swizzle_b32 is here as user3528438 pointed out.
The __byte_perm does not swap data with other lanes. It just gathers any 4 bytes from two 32-bit registers in its own lane and stores it to a register. There is no cross-lane traffic.
I'm guessing the next question would be how to do a "byte permute" on AMD GCN hardware. The instruction for that is v_perm_b32. (see page 12-152 here) It basically selects any four bytes from two specified 32-bit registers.