UByte4 data type

Accelerated Computing CUDA CUDA Programming and Performance
1

What is the equivalent of UByte4 which was valid in previous cuda versions?

There is a code which looks like

#define BSIZE 256
#define STAGEBLOCK1(index)	CUT_BANK_CHECKER( stageBlock1, index )
__global__ void aesEncrypt128( unsigned * result, unsigned * inData, int inputSize)
{
    unsigned bx		= blockIdx.x;
    unsigned tx		= threadIdx.x;
    __shared__ UByte4 stageBlock1[BSIZE];     // I missed this line in the first version of my post. Thanks to njuffa for mentioning that.
    STAGEBLOCK1(tx).uival	= inData[BSIZE * bx + tx ];
    ...
}

Based on what is stated in [1, 2], I fixed the third line as

stageBlock1[ tx ] = inData[BSIZE * bx + tx ];

However, I get this error:

error: no operator “=” matches these operands
operand types are: UByte4 = unsigned int

I can not find a documentation on the meaning of UByte4. Is it enough to change UByte4 to unsinged int (the meanings look the same)

Any help?

[1] What is CUT_BANK_CHECKER - CUDA Programming and Performance - NVIDIA Developer Forums
[2] memory bank conflict check in new cuda versions - CUDA Programming and Performance - NVIDIA Developer Forums

2

To the best of my knowledge, UByte4 has never been a valid CUDA type. This is likely coming from a non-CUDA header file, maybe the same that provided CUT_BANK_CHECKER, maybe a different one (a quick Google search shows UByte4 occurring in various graphics contexts). You might make faster progress if you track down the missing header file(s) for this project.

For what it is worth, I see no instance of UByte4 in the code you show above. It seems reasonable to assume that this is a uint32_t, so why not use that?

3

Sorry for the missed line. I fixed that.
Yes, I replaced that and the error has been disappeared.

Also, I don’t see any doc for uival. There is also a ubval which I don’t understand that.

I replaced

STAGEBLOCK1(tx).uival	= inData[BSIZE * bx + tx ];

with

stageBlock1[ tx ] = inData[BSIZE * bx + tx ];

and I think it is OK.

However, I see

unsigned op1 = STAGEBLOCK1( tx ).ubval[0];
unsigned op2 = STAGEBLOCK1( tx ).ubval[1];

I don’t know what to do with the index number of ubval. Any thought?

4

Sorry, I don’t know what ‘ubval’ and ‘uival’ are. Find the missing header file that defines all this. This might refer to a

union {
   unsigned int uival;
   unsigned char ubval[4];
}

but that’s merely a guess.

5

Are you working on the ispass2009 GitHub code by any chance? In this file:

https://github.com/gpgpu-sim/ispass2009-benchmarks/blob/master/AES/sbox_E.h

one finds:

union  UByte4 {
    unsigned int uival;
    unsigned char ubval[4];
};

An old copy of cutil.h (utility library prior to CUDA 5, with CUT_BANK_CHECKER) can be found in this GitHub project:

https://github.com/kashif/cuda-workshop/tree/master/cutil

I am providing this link for reference only, and will warn against using cutil.h in your own code.

6

Yes, I am working on that code. That is interesting because I greped for UByte4 in the files and didn’t notice that. Maybe because of many instances, I missed the definition in the search result. Thanks for that.

No, I am not planning to use the cutil.h and I am trying to port the code to the newer versions. So, the nightmare is that for some functions and macros in the old cuda, there is no equivalent in the newer versions. Therefore, I have to understand the concept and implement with new things. The bad part is to modify a code which has been written by someone else.

7

Once again: cutil.h was never part of CUDA, and NVIDIA specifically and repeatedly warned CUDA programmers not to use it in their own projects. Whoever wrote this AES implementation obviously didn’t pay attention. CUDA proper has deprecated certain old features, but they rarely go away.

cutil.h was a collection of utility functions the creators of the CUDA sample applications created for their own use, to keep these applications free of clutter, in order to focus attention on whatever concept each app demonstrates.

That (modifying other people’s code) is what software engineers do for a living. Most of the time at least.