How do you determine the OpenCL address space of the load/store instructions?

[zehanort]

This is more of a question than an issue... I am very interested on how oclgrind manages to distinguish the OpenCL address spaces regarding the operands of the load and store instructions. As I have seen, this is some info that clang does not keep when it compiles OpenCL to LLVM IR. The only info that clang keeps is the OpenCL address spaces of the kernel arguments and that 's pretty much it? I observed in my experiments that the getPointerAddressSpace method of the LoadInst and StoreInst classes return always 0 (i.e. global address space or uninitialized, I believe the later due to the many new variables that are introduced by LLVM between the actual kernel argument and where the load/store operation actually takes place). Can you please give me an explanation on what exactly oclgrind does to tackle this issue -which, as I 've seen, it actually does!-, or point me to the code that does it? It would be extremely helpful for my thesis. Thank you in advance!

[jrprice]

This information should be in the LLVM IR, but it only seems to be retained if you compile for the spir or spir64 targets.

Example kernel:

kernel void foo(constant int *a, global int *b, global int *c) {
  int i = get_global_id(0);
  c[i] = a[i] + b[i];
}

Compiling with bin/clang -c -x cl foo.cl -emit-llvm -include opencl-c.h -target spir64 produces this IR (extracted interesting part):

define dso_local spir_kernel void @foo(i32 addrspace(2)* nocapture readonly %0, i32 addrspace(1)* nocapture readonly %1, i32 addrspace(1)* nocapture %2) local_unnamed_addr #0 !kernel_arg_addr_space !4 !kernel_arg_access_qual !5 !kernel_arg_type !6 !kernel_arg_base_type !6 !kernel_arg_type_qual !7 {
  %4 = tail call spir_func i64 @_Z13get_global_idj(i32 0) #2
  %5 = shl i64 %4, 32
  %6 = ashr exact i64 %5, 32
  %7 = getelementptr inbounds i32, i32 addrspace(2)* %0, i64 %6
  %8 = load i32, i32 addrspace(2)* %7, align 4, !tbaa !8
  %9 = getelementptr inbounds i32, i32 addrspace(1)* %1, i64 %6
  %10 = load i32, i32 addrspace(1)* %9, align 4, !tbaa !8
  %11 = add nsw i32 %10, %8
  %12 = getelementptr inbounds i32, i32 addrspace(1)* %2, i64 %6
  store i32 %11, i32 addrspace(1)* %12, align 4, !tbaa !8
  ret void
}

The address spaces are carried through to the load/store instructions, and so should be returned by getPointerAddressSpace.