Comments (6)
total_size += sizeof(int) * some_other_size * 3;
- this might be the problem. Alignment might be wrong. hipcub_buffer
should have the same alignment as if it was allocated directly with hipMalloc
.
from rocprim.
Then it should work, if the hipcub_buffer
is the first chunk of buffer data in the row, right?
from rocprim.
I guess. I don't see any other problem right now, but only looked at your code. If you are allocating multiple buffers with one hipMalloc
I'd recommend aligning every buffer to 256 bytes. It's not much and it helps.
from rocprim.
Works, thanks!
Btw: Buffer size returned by rocprim is not a multiple of 256.
I think I remember that for some nvidia cards pointer alignment is not always 256, do you happen to have any information on that or am I wrong.
from rocprim.
Btw: Buffer size returned by rocprim is not a multiple of 256.
That's true. Temporary buffer size does not have to be a multiple of 256. However, its pointer may need to be aligned to 256 bytes.
I think I remember that for some nvidia cards pointer alignment is not always 256, do you happen to have any information on that or am I wrong.
From NVIDIA docs: "Any address of a variable residing in global memory or returned by one of the memory allocation routines from the driver or runtime API is always aligned to at least 256 bytes." It might be more on newer GPUs, but 256 bytes is good enough for everything to work correctly.
from rocprim.
Great, thanks a lot!
from rocprim.
Related Issues (20)
- Block_reduce fails to distribute correct answer to all lanes when hipBlockDim > 64 HOT 5
- Invalid use of inline assembly for GFX1010 target HOT 14
- Compiler Build issue HOT 1
- ROCm 3.9.1 fails to compile benchmark_warp_scan HOT 2
- error: reference to __host__ function 'inclusive_scan<rocprim::default_config, double *, double *, thrust::plus<double>>' in __host__ __device__ function HOT 2
- Ignoring return value warning HOT 2
- error: invalid operands to binary expression HOT 6
- Please enable two factor authentication in your github account
- Alternative documentation HOT 2
- error: no matching function for call to 'ceiling_div' HOT 2
- Add support for NAVI22 and NAVI23 i.e. gfx1031 and gfx1032 HOT 1
- rocPRIM reduction (block_reduce_int) issue HOT 2
- Follow the example of rocPRIM custom types example but compile failed. HOT 2
- rocPRIM on NAVI22 = gfx1031 only one test fails - can we fix this? HOT 6
- rocPRIM -- rocThrust dependency: error: no member named 'init_offset_scan_state_kernel' in rocprim::detail HOT 2
- rocPRIM 5.4.3 failed device_adjacent_difference on gfx906 HOT 4
- rocPRIM selects wrong code paths for warp reductions and scas on gfx1036 HOT 3
- Page fault during `rocprim::select` for no obvious reason. HOT 2
- `rocprim::block_load` fails to instantiate HOT 8
- first installation failed
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 rocprim.