intptr_t support in llvm

Hi all,

I tried compiling an opencl kernel with intptr_t datatype. However it gives error by default using LLVM/Clang 3.2 for nvptx . To allow usage, I tried following typedefs ( by looking at tests in llvm sources ).

typedef int intptr_t;
and
typedef __typeof( (int*) 0) intptr_t;

Both definitions compiles the code, however gives following warning.
warning: incompatible pointer to integer conversion assigning to ‘int’ from ‘int *’; dereference with *

Is it safe to use the above typedefs ? The generated code seems correct for x86 machines (with llvm instruction ptrtoint). Or is there better way to make sure that the semantics of intprt_t are preserved in clang/llvm for all archs ?

The code I tried to compile is

__kernel void intptr_t_kernel(int *a, int b)
{
intptr_t c = (intptr_t)a;
b = (int
) c;
*a += b;
}

and the code generated is

define ptx_kernel void @intptr_t_kernel(i32* %a, i32* nocapture %b) nounwind noinline {
entry:
%0 = load i32* %a, align 4, !tbaa !1
%add.ptr = getelementptr inbounds i32* %a, i32 %0
%conv = ptrtoint i32* %add.ptr to i32
store i32 %conv, i32* %a, align 4, !tbaa !1
ret void
}

Thanks a lot for help.

Regards,
Ankur

Hi all,

I tried compiling an opencl kernel with intptr_t datatype. However it gives error by default using LLVM/Clang 3.2 for nvptx . To allow usage, I tried following typedefs ( by looking at tests in llvm sources ).

typedef int intptr_t;
and
typedef __typeof( (int*) 0) intptr_t;

Both definitions compiles the code, however gives following warning.
warning: incompatible pointer to integer conversion assigning to ‘int’ from ‘int *’; dereference with *

Is it safe to use the above typedefs ? The generated code seems correct for x86 machines (with llvm instruction ptrtoint). Or is there better way to make sure that the semantics of intprt_t are preserved in clang/llvm for all archs ?

The code I tried to compile is

__kernel void intptr_t_kernel(int *a, int b)
{
intptr_t c = (intptr_t)a;
b = (int
) c;
*a += b;
}

and the code generated is

define ptx_kernel void @intptr_t_kernel(i32* %a, i32* nocapture %b) nounwind noinline {
entry:
%0 = load i32* %a, align 4, !tbaa !1
%add.ptr = getelementptr inbounds i32* %a, i32 %0
%conv = ptrtoint i32* %add.ptr to i32
store i32 %conv, i32* %a, align 4, !tbaa !1
ret void
}

Thanks a lot for help.

Regards,
Ankur

Leaving aside any OpenCL specific issues (you can use pointers in a more restricted way than in general C), are you sure that the final line is doing what you want (I think that’s what the warning is about). You’re adding the int* address held in b to the int value stored at the address held in a. C will let you do this primarily because it’s got quite weak rules about which conversions you must do explicitly, so this will go through on a machine where the int width is the same as an int* width.

Does the error go away if you do *a+=*b ?

I don’t think either of your typedefs will reliably work (ie, in the interesting cases where sizeof(int*) != sizeof(int)), I think the only way to do it is via explicit pre-processor typedefs based on the pointer size of the machine you’re compiling on.

Regards,

Dave

Hi Dave,

Thanks for the help.

I see my mistake. “*a+=*b” works great.

I was hoping there is a portable way of handling intptr_t ( hoping clang’s some hidden feature for the same).

Thanks again for the reply,

Ankur