Error: asm operand type size(8) does not match type/size implied by constraint 'r'

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi, I got this error when I’m providing an u32 input to the ptx asm code

"r"(3254828548)

the device is AWS p3.2xlarge with 1xV100, ubuntu22.04, cuda12.2.

how could I correct the code or environment? thanks!

2

it would be useful to see the line or lines of code generating the error. generally speaking, to write correct inline PTX, you should follow this guide. The constraint r implies a single 32-bit register and therefore a single 32-bit quantity. That particular number would fit in an unsigned 32-bit quantity but not a signed one. It might be as simple as putting a U after the constant (e.g. 3254828548U) , but it would be better to see the actual code.

3

I am unable to reproduce this issue, see my test case below.

#include <stdio.h>
#include <stdlib.h>

__device__ unsigned int loadconst (void)

{
    unsigned int res;
#if 0
    asm ("mov.b32 %0,%1;\n\t" : "=r"(res) : "r"(0xC200BE04));
#else
    asm ("mov.b32 %0,%1;\n\t" : "=r"(res) : "r"(3254828548));
#endif
    return res;
}

__global__ void kernel (void)
{
    unsigned int t = loadconst();
    printf ("t=%u\n", t);
}

int main (void)
{
    kernel<<<1,1>>>();
    cudaDeviceSynchronize();
    return EXIT_SUCCESS;
}
4

thanks Robert! It works with this extra U.
Could the constant 3254828548 be signed by default?

5

unfortunately when I compiled your code with the cmd nvcc -o main main.cu.
The compiler still complains:

main.cu(11): error: asm operand type size(8) does not match type/size implied by constraint 'r'
      asm ("mov.b32 %0,%1;\n\t" : "=r"(res) : "r"(3254828548));
                                              ^

1 error detected in the compilation of "main.cu".

But when I added U to it as 3254828548U, it passed.
how could it be so different on my side?

6

Without study and understanding of the code, there is no way to tell if this is a “correct” solution. I leave that to you.

In C++ anyway, all decimal integer constants are signed by default. Furthermore, with respect to type, I believe they will be roughly speaking of the first type that they fit into: int, then long, then long long.

The signed value 3254828548 will not fit into an int type (I stated this already), and the next types (long or long long) that it will fit into on this linux platform are both 8-byte types, which is (I believe) the origin of the error message.

1 Like
7

Different compiler versions, I would guess. I don’t have the latest nvcc installed. The compiler in CUDA 12.2 is compliant with the ISO-C++ standard.

Per the C++ standard (I looked at ISO-C++11, section 2.14.2 Integer literals; this should not have changed in later versions), an unsuffixed decimal integer literal constant will be assigned the first type into which it fits from the following list of types: int, long int, long long int. This explains your original observation. Since the number does not fit into an int, it is assigned a 64-bit integer type, which then mismatches the r binding attribute.

However, when the numerically same literal constant is presented as an unsuffixed octal or hexadecimal number, the list of types considered is int, unsigned int, long int, unsigned long int, long long int, unsigned long long int. Why the C++ standard specifies a different list of types depending on the input format I do not know. Presumably there is a rationale for this, but it may not be documented.

So you should be able to use unsuffixed hexadecimal 0xC200BE04, which is assigned type unsigned int, or suffixed decimal 3254828548u, which is also assigned type unsigned int.

1 Like
8

thanks guys for your time! problem solved:D

9

This topic was automatically closed 14 days after the last reply. New replies are no longer allowed.