Given the following opencl code: __kernel void ldiv(__global long * out, long arg0, long arg1) { out[0] = arg0 / arg1; } Mesa generates the following code: ; ModuleID = 'radeon' target datalayout = "e-p:32:32:32-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-v16:16:16-v24:32:32-v32:32:32-v48:64:64-v64:64:64-v96:128:128-v128:128:128-v192:256:256-v256:256:256-v512:512:512-v1024:1024:1024-v2048:2048:2048-n32:64" target triple = "r600--" ; Function Attrs: nounwind define void @ldiv(i32 addrspace(1)* nocapture %out, i32 %arg0, i32 %arg1) #0 { %1 = sdiv i32 %arg0, %arg1 store i32 %1, i32 addrspace(1)* %out, align 4, !tbaa !1 ret void } attributes #0 = { nounwind "less-precise-fpmad"="false" "no-frame-pointer-elim"="false" "no-frame-pointer-elim-non-leaf"="false" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "unsafe-fp-math"="false" "use-soft-float"="false" } !opencl.kernels = !{!0} !0 = metadata !{void (i32 addrspace(1)*, i32, i32)* @ldiv} !1 = metadata !{metadata !"long", metadata !2} !2 = metadata !{metadata !"omnipotent char", metadata !3} !3 = metadata !{metadata !"Simple C/C++ TBAA"} This is obviously wrong since according to the opencl specification (Section 6.1.1 "Built-in Scalar Data Types") the long and ulong types are 64 bit long whilst i32 is 32 bit long.
I have been checking a bit around this issue, whilst modifying clc/clctypes.h to add "typedef unsigned __INT64_TYPE__ ulong;" works, doing the same with long doesn't, quite likely because it is a standard C type. I'm unsure on how to modify libclc or clang to fix this issue.
The type sizes are supposed to be controlled by the TargetInfo class. TargetInfo::setForcedLangOptions() in clang/lib/basic/TargetInfo.cpp is where the type sizes are set for OpenCL. R600's TargetInfo is defined in clang/lib/basic/Target.cpp Maybe there is a bug in one of these two places. Also, if you invoke clang like this: clang -cc1 -E -dM -ffreestanding -triple=r600-- -std=cl -o - < /dev/null It will dump all of the default definitions for macros like __INT64_TYPE__
Yeah I found it, the issue is basically that the special case for OpenCL isn't defined at clang/lib/basic/TargetInfo.cpp on clang-3.3 I suppose it's a thing of waiting for clang-3.4 for this issue to get fixed.
Use of freedesktop.org services, including Bugzilla, is subject to our Code of Conduct. How we collect and use information is described in our Privacy Policy.