Comments (1)
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
.
from oclgrind.
Related Issues (20)
- Include a Linux binary in oclgrind releases HOT 2
- cmake fails on "Clang libraries not found (set CLANG_ROOT)" HOT 5
- Question about optimisation flags HOT 2
- Any future plan to support the atomic_* data types and their related API functions?
- Oclgrind segfaults when running program HOT 2
- Add support to LLVM/Clang 11 HOT 4
- Unable to treat llvm 'freeze' instruction. HOT 3
- New git tag/release with LLVM 11 support HOT 2
- Cannot configure device type
- Change endianness of data buffers
- Detailed description of each command and usage
- clCreateBuffer thread safety
- "Unsupported constant int size: 5 bytes" HOT 2
- Still get 'Undefined external function: llvm.fshl.i32' error without inline keyword HOT 2
- How to stop any error which would be print? HOT 1
- Project status HOT 2
- Compile Error: ‘const class llvm::ConstantExpr’ has no member named ‘getIndices’ HOT 1
- Support for LLVM 16 HOT 1
- SPIR-V Instruction reordering resulting in false-positive read-write data race? HOT 1
Recommend Projects
-
React
A declarative, efficient, and flexible JavaScript library for building user interfaces.
-
Vue.js
🖖 Vue.js is a progressive, incrementally-adoptable JavaScript framework for building UI on the web.
-
Typescript
TypeScript is a superset of JavaScript that compiles to clean JavaScript output.
-
TensorFlow
An Open Source Machine Learning Framework for Everyone
-
Django
The Web framework for perfectionists with deadlines.
-
Laravel
A PHP framework for web artisans
-
D3
Bring data to life with SVG, Canvas and HTML. 📊📈🎉
-
Recommend Topics
-
javascript
JavaScript (JS) is a lightweight interpreted programming language with first-class functions.
-
web
Some thing interesting about web. New door for the world.
-
server
A server is a program made to process requests and deliver data to clients.
-
Machine learning
Machine learning is a way of modeling and interpreting data that allows a piece of software to respond intelligently.
-
Visualization
Some thing interesting about visualization, use data art
-
Game
Some thing interesting about game, make everyone happy.
Recommend Org
-
Facebook
We are working to build community through open source technology. NB: members must have two-factor auth.
-
Microsoft
Open source projects and samples from Microsoft.
-
Google
Google ❤️ Open Source for everyone.
-
Alibaba
Alibaba Open Source for everyone
-
D3
Data-Driven Documents codes.
-
Tencent
China tencent open source team.
from oclgrind.