free speedier betterer trueer fasterer gpu.cu true no scam#24
free speedier betterer trueer fasterer gpu.cu true no scam#24zajonc6118 wants to merge 0 commit into
Conversation
BoySanic
left a comment
BoySanic
left a comment
There was a problem hiding this comment.
Overall I hope you'll revert some of these formatting changes and do some analysis into whether the 1D array is really a performance win or not.
Welcome to github :)
| void *device_kernel_0A_addr; | ||
| TRY_CUDA(cudaGetSymbolAddress(&device_kernel_0A_addr, device_kernel_0A)); | ||
| TRY_CUDA(cudaMemcpy(device_kernel_0A_addr, temp_0A, sizeof(temp_0A), cudaMemcpyHostToDevice)); | ||
|
|
There was a problem hiding this comment.
Unnecessary formatting change
|
|
||
| __shared__ float conv_z0[513][6]; | ||
| __shared__ float conv_z1[513][6]; | ||
| __shared__ alignas(16) float conv_z0[512 * 6]; |
There was a problem hiding this comment.
I would like to see a performance comparison here because I kind of doubt that there's a functional difference between a flat array and a 2D array, but I'd be happy to be mistaken.
| p_z[dnz] = oct_0A.p[(nz + dnz) & 0xFF] & 0xF; | ||
| } | ||
|
|
||
| float* row0 = &conv_z0[nz * 6]; |
There was a problem hiding this comment.
This is the kind of complexity I think we can avoid with just using the 2D array again and idk that this really did anything for us.
| const uint16_t* cw1 = &w1[candidate]; | ||
|
|
||
| const float gate = score_center_2x2(conv_z0, conv_z1, cw0, cw1); | ||
| const float gate = score_center_2x2_flat(conv_z0, conv_z1, cw0, cw1); |
There was a problem hiding this comment.
You changed the name just because it's a 1D array instead of 2D?
| const float gate = score_center_2x2_flat(conv_z0, conv_z1, cw0, cw1); | ||
| if (gate >= kGradVecs1PrefilterThreshold) { | ||
| const float score = score_full_12(conv_z0, conv_z1, cw0, cw1); | ||
| const float score = score_full_12_flat(conv_z0, conv_z1, cw0, cw1); |
There was a problem hiding this comment.
Same comment about the name as on R912
| // cactus was here :) | ||
| namespace KernelFilter2_0A { | ||
| using T = KernelFilter2::Template<-5500, 3, 8 * 1024, 256, 27, false, false, true>; | ||
| using T = KernelFilter2::Template<-5500, 3, 8 * 1024, 256, 27, false, false, true>; |
There was a problem hiding this comment.
Something about how you reformatted this made the diff really weird.
Nothing seems to have changed except that you added an indent though?
There was a problem hiding this comment.
I think I would prefer to re-review this part after you revert the indentation change, and we can re-indent this section in its own PR to make it easier.
|
|
||
| namespace KernelFilter2_0B { | ||
| using T = KernelFilter2::Template<-5500, 3, 8 * 1024, 256, 20, false, false, false>; | ||
| using T = KernelFilter2::Template<-5500, 3, 8 * 1024, 256, 20, false, false, false>; |
There was a problem hiding this comment.
Same indentation comment here
| kernel<<<32 * 256, threads_per_block, 0, stream>>>(inputs, outputs, results); | ||
| TRY_CUDA(cudaGetLastError()); | ||
| } | ||
| void run(InputBuffer<SeedPos> inputs, OutputBuffer<SeedPos> outputs, |
There was a problem hiding this comment.
Why split into two lines? kinda doesn't seem like anything of substance changed here otherwise so this just adds overhead trying to review it.
| } | ||
|
|
||
| template <typename IndexT> | ||
| __device__ __forceinline__ float score_center_2x2_flat( |
There was a problem hiding this comment.
These new replacements for the old functions are a bit annoying because you left the old ones in instead of replacing them. So I can't really know what's changed between them, and I'm guessing it's mostly just that you made it a 1D array instead of 2D and I'm not convinced that's a big win.
|
Oh there were too many other issues but please also remove the log file you erroneously included |
2 participants
free speedier speedy betterer trueer no scam 1.4.3z for everyone