Инструкция PTX prmt предоставляет функциональные возможности машинной инструкции PRMT. Режим по умолчанию для инструкции prmt используется, когда не указан ни один из специальных режимов .f4e, .b4e, .rc8, .ecl, .ecr, .rc16.
Режим по умолчанию имеет два суббайтовых подрежима, управляемых старшим значащим битом из 4 -битное поле селектора для каждого из восьми байтов источника. Обычно используемый подрежим должен иметь ноль msb поля селектора, что означает, что целевой байт дословно копируется из указанного исходного байта. Этот подрежим открывается через функцию устройства intrinsi c __byte_perm() и обычно используется для извлечения, вставки и перестановки байтов или выполнения сдвигов битов, кратных 8. Пример использования можно увидеть в этот ответ .
Другим подрежимом является special , в котором вместо копирования всего исходного байта он реплицирует старший значащий бит указанного исходного байта в целевом байте. Для этого msb поля селектора должен быть установлен в единицу. Программисты должны использовать встроенную сборку PTX для доступа к этой функции.
Я не проектировал аппаратное обеспечение графического процессора, поэтому не могу сказать, почему был реализован этот подрежим. Обычно полезно, когда msb каждого байта служит логическим значением, которое нужно преобразовать в маску для всего байта. Это, в свою очередь, обычно полезно для побайтной обработки в 32-битном регистре. Обратите внимание, что CUDA включает в себя множество встроенных функций устройства для такой обработки, и разборка подтвердит, что подрежим репликации msb режима по умолчанию prmt используется для многих из них.
Полностью проработанный пример, Эмуляция операции paddsb (побайтное сложение со знаковым насыщением) показана ниже. Обратите внимание на использование prmt с репликацией msb внутри masked_sign_to_byte_mask().
#include <stdio.h>
#include <stdlib.h>
#include <stdint.h>
#if (__CUDACC__)
#define __HOST__ __host__
#define __DEVICE__ __device__
#else // __CUDACC__
#define __HOST__
#define __DEVICE__
#endif // __CUDACC__
#define MSB_MASK (0x80808080U) // mask for msb of each byte
// r = (a ^ b) & ~c
__HOST__ __DEVICE__ uint32_t lop3_14 (uint32_t a, uint32_t b, uint32_t c)
{
uint32_t r;
#if (__CUDA_ARCH__ >= 300)
asm ("lop3.b32 %0,%1,%2,%3,0x14;\n\t" : "=r"(r) : "r"(a), "r"(b), "r"(c));
#else // __CUDA_ARCH__
r = (a ^ b) & ~c;
#endif // __CUDA_ARCH__
return r;
}
// r = (a ^ b) & c
__HOST__ __DEVICE__ uint32_t lop3_28 (uint32_t a, uint32_t b, uint32_t c)
{
uint32_t r;
#if (__CUDA_ARCH__ >= 300)
asm ("lop3.b32 %0,%1,%2,%3,0x28;\n\t" : "=r"(r) : "r"(a), "r"(b), "r"(c));
#else // __CUDA_ARCH__
r = (a ^ b) & c;
#endif // __CUDA_ARCH__
return r;
}
// r = a ^ (~b & c)
__HOST__ __DEVICE__ uint32_t lop3_d2 (uint32_t a, uint32_t b, uint32_t c)
{
uint32_t r;
#if (__CUDA_ARCH__ >= 300)
asm ("lop3.b32 %0,%1,%2,%3,0xd2;\n\t" : "=r"(r) : "r"(a), "r"(b), "r"(c));
#else // __CUDA_ARCH__
r = a ^ (~b & c);
#endif // __CUDA_ARCH__
return r;
}
// r = (a & c) | (b & ~c)
__HOST__ __DEVICE__ uint32_t lop3_f4 (uint32_t a, uint32_t b, uint32_t c)
{
uint32_t r;
#if (__CUDA_ARCH__ >= 300)
asm ("lop3.b32 %0,%1,%2,%3,0xf4;\n\t" : "=r"(r) : "r"(a), "r"(b), "r"(c));
#else // __CUDA_ARCH__
r = (a & c) | (b & ~c);
#endif // __CUDA_ARCH__
return r;
}
__HOST__ __DEVICE__ uint32_t masked_sign_to_byte_mask (uint32_t a)
{
#if (__CUDA_ARCH__ >= 200)
asm ("prmt.b32 %0,%0,0,0xba98;" : "+r"(a)); // convert MSBs to masks
#else
a = a & MSB_MASK;
a = a + a - (a >> 7); // extend MSBs to full byte to create mask
#endif
return a;
}
__HOST__ __DEVICE__ uint32_t masked_select (uint32_t a, uint32_t b, uint32_t m)
{
#if (__CUDA_ARCH__ >= 300)
return lop3_f4 (a, b, m);
#elif 0
return (((a)&(m))|((b)&(~(m))));
#else
return((((a)^(b))&(m))^(b));
#endif
}
/*
my_paddsb() performs byte-wise addition with signed saturation. In the
case of overflow, positive results are clamped at 127, while negative
results are clamped at -128.
*/
__HOST__ __DEVICE__ uint32_t my_paddsb (uint32_t a, uint32_t b)
{
uint32_t sum, res, ofl, sga, msk;
res = (a & ~MSB_MASK) + (b & ~MSB_MASK);
sum = a ^ b;
ofl = lop3_14 (res, a, sum); // ofl = (res ^ a) & ~sum
sga = masked_sign_to_byte_mask (a); // sign(a)-mask
msk = masked_sign_to_byte_mask (ofl);// overflow-mask
res = lop3_d2 (res, ~MSB_MASK, sum); // res = res ^ (MSB_MASK & sum)
sga = lop3_28 (sga, ~MSB_MASK, msk); // sga = (sga ^ ~MSB_MASK) & msk
res = masked_select (sga, res, msk); // res = (sga & msk) | (res & ~msk)
return res;
}
__global__ void kernel (uint32_t a, uint32_t b)
{
printf ("GPU: %08x\n", my_paddsb (a, b));
}
int main (void)
{
uint32_t a = 0x12ef70a0;
uint32_t b = 0x34cd6090;
kernel<<<1,1>>>(a, b);
cudaDeviceSynchronize();
printf ("CPU: %08x\n", my_paddsb (a, b));
return EXIT_SUCCESS;
}