Timeline for Computational verification of Collatz conjecture using OpenCL
Current License: CC BY-SA 4.0
17 events
| when toggle format | what | by | license | comment | |
|---|---|---|---|---|---|
| Oct 24, 2019 at 7:30 | comment | added | tmlen | GPUs may have instructions for 64bit operations (even if they operate on two registers). Maybe also change the kernel so that 2 work items (or 4) are used for each task, processing the lo and hi parts of the 128bit integer, and using subgroup operations to carry over data. | |
| Oct 24, 2019 at 7:16 | history | bounty ended | DaBler | ||
| Oct 24, 2019 at 7:16 | vote | accept | DaBler | ||
| Oct 24, 2019 at 7:16 | comment | added | DaBler | Either way, this answer helped me a lot. | |
| Oct 24, 2019 at 7:14 | comment | added | DaBler | I believe I can get a much bigger acceleration if I rewrite the whole program only to use only 32×32 multiplication (as the GPUs are 32-bit machines). However, this would not be useful in practice, as such small numbers have been verified for a long time. So working with the 128-bit type cannot be avoided here. | |
| Oct 24, 2019 at 7:14 | comment | added | DaBler | I rewrote the multiplication from 128×128 to 128×64 and 128×32, and it really helped a bit again. On my GeForce GTX 1050 Ti, the speedup is about 6 % (comparing the 64-bit implementation to 128-bit) and 31 % (32-bit to 128-bit). On the other hand, on my older machine with GeForce GT 730, the spedup is about 33% (64-bit to 128-bit) and 43 % (32-bit to 128-bit). | |
| Oct 19, 2019 at 14:40 | comment | added | DaBler | Profiling is surely a good idea. Does also __local memory access require coalescing? Note that the __global memory is only accessed using unsigned long integers, having the access pattern I mentioned. | |
| Oct 18, 2019 at 18:49 | comment | added | tmlen |
Yes but not sure if GPUs still coalesce the access for 128bit values. Maybe separate it into 2 64bit arrays for the hi and lo part. Or replace the LUT by a function with a large switch statement with cases for all 81 values. But it would surely be good to use a profiler (such as AMD CodeXL, or with CUDA the NVVP).
|
|
| Oct 18, 2019 at 13:49 | comment | added | DaBler |
I also guess that the kernel always access the global memory in a coalesced manner. The access pattern is like array[get_global_id(0)]. Is that right?
|
|
| Oct 18, 2019 at 13:40 | comment | added | DaBler |
Placing the LUT into __local memory really helps a bit. The acceleration is about 3% at Tesla K20Xm and about 1% at GeForce GTX 1050 Ti. It is not much, but after a long time finally something that helped. However, initializing the LUT from global __constant memory slows the program down significantly. I guess that calculating the pow3 is much faster than accessing global memory.
|
|
| Oct 16, 2019 at 8:28 | history | edited | tmlen | CC BY-SA 4.0 |
deleted 1 character in body
|
| Oct 16, 2019 at 8:21 | comment | added | tmlen | @DaBler updated the answer | |
| Oct 16, 2019 at 8:20 | history | edited | tmlen | CC BY-SA 4.0 |
added 1318 characters in body
|
| Oct 16, 2019 at 8:04 | history | edited | Toby Speight | CC BY-SA 4.0 |
Spelling and markdown
|
| Oct 16, 2019 at 7:34 | comment | added | DaBler | The __constant space leads to the most significant slowdown. This is probably because all threads access the cache hierarchy at the same time. Leaving the LUT in on-chip memory is much faster. Maybe there's something else I could try? | |
| Oct 16, 2019 at 7:31 | comment | added | DaBler | Thank you for the answer. This is exactly what I needed! Unfortunately, neither of these points led to acceleration (all resulted either in a small or significant slowdown). Probably it is important to note that the most inner do-while loop is executed only once in most cases, in less cases twice, in even less case three times, etc. So the diverge after the first iteration is probably (?) quite quickly resolved. Also rewriting the code so that the for-loop is controlled by exactly same values for the counter does not help. A particular value probably doesn't matter. | |
| Oct 14, 2019 at 14:12 | history | answered | tmlen | CC BY-SA 4.0 |