Add BLOCK_LOAD_STRIPED and BLOCK_STORE_STRIPED #274

mnicely · 2021-03-25T13:23:04Z

This PR adds the following to BlockLoadAlgorithm

BLOCK_LOAD_STRIPED
It's basically BLOCK_LOAD_TRANSPOSE without the BlockExchange

This PR adds the following to BlockStoreAlgorithm

BLOCK_STORE_STRIPED
It's basically BLOCK_STORE_TRANSPOSE without the BlockExchange

mnicely · 2021-03-25T13:24:40Z

@llukas & @jszuppe for vis

alliepiper · 2021-03-26T17:23:12Z

This looks good, I'll plan to take a closer look sometime next week.

Can you update the test_block_load_store.cuh test? It should just take an extra line in this block that exercises the new BlockAlgorithm.

mnicely · 2021-03-26T19:43:08Z

Thanks @allisonvacanti. Looks like everything passes.

...
      Start  3: cub.cpp14.test.block_load_store.thorough
 3/52 Test  #3: cub.cpp14.test.block_load_store.thorough ..............   Passed    0.14 sec
      Start  4: cub.cpp14.test.block_load_store.benchmark
 4/52 Test  #4: cub.cpp14.test.block_load_store.benchmark .............   Passed    0.10 sec
...

This PR covers our C2C scenarios. I need to think more about how R2C and C2R might work

alliepiper · 2021-05-10T20:59:07Z

LGTM -- Starting tests.

alliepiper · 2021-05-10T21:07:30Z

DVS CL: 29946764
gpuCI: NVIDIA/thrust#1427

alliepiper · 2021-05-12T20:59:02Z

gpuCI identified some issues building with nvcc+clang. @mnicely can you take a look? Looks like a missing member variable.

../dependencies/cub/cub/block/block_store.cuh(628): error #245: a nonstatic member reference must be relative to a specific object
          detected during:
            instantiation of "void cub::BlockStore<T, BLOCK_DIM_X, ITEMS_PER_THREAD, ALGORITHM, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH>::StoreInternal<cub::BLOCK_STORE_STRIPED, DUMMY>::Store(OutputIteratorT, T (&)[ITEMS_PER_THREAD], int) [with T=char, BLOCK_DIM_X=15, ITEMS_PER_THREAD=1, ALGORITHM=cub::BLOCK_STORE_STRIPED, BLOCK_DIM_Y=1, BLOCK_DIM_Z=1, PTX_ARCH=800, DUMMY=0, OutputIteratorT=cub::DiscardOutputIterator<std::ptrdiff_t>]" 
(1048): here
            instantiation of "void cub::BlockStore<T, BLOCK_DIM_X, ITEMS_PER_THREAD, ALGORITHM, BLOCK_DIM_Y, BLOCK_DIM_Z, PTX_ARCH>::Store(OutputIteratorT, T (&)[ITEMS_PER_THREAD], int) [with T=char, BLOCK_DIM_X=15, ITEMS_PER_THREAD=1, ALGORITHM=cub::BLOCK_STORE_STRIPED, BLOCK_DIM_Y=1, BLOCK_DIM_Z=1, PTX_ARCH=800, OutputIteratorT=cub::DiscardOutputIterator<std::ptrdiff_t>]" 
../dependencies/cub/test/test_block_load_store.cu(137): here
            instantiation of "void Kernel<BLOCK_THREADS,ITEMS_PER_THREAD,LOAD_ALGORITHM,STORE_ALGORITHM,InputIteratorT,OutputIteratorT>(InputIteratorT, OutputIteratorT, OutputIteratorT, int) [with BLOCK_THREADS=15, ITEMS_PER_THREAD=1, LOAD_ALGORITHM=cub::BLOCK_LOAD_STRIPED, STORE_ALGORITHM=cub::BLOCK_STORE_STRIPED, InputIteratorT=const char *, OutputIteratorT=cub::DiscardOutputIterator<std::ptrdiff_t>]" 
../dependencies/cub/test/test_block_load_store.cu(180): here
            instantiation of "void TestKernel<T,BLOCK_THREADS,ITEMS_PER_THREAD,LOAD_ALGORITHM,STORE_ALGORITHM,InputIteratorT,OutputIteratorT>(T *, InputIteratorT, OutputIteratorT, OutputIteratorT, T *, T *, int, int) [with T=char, BLOCK_THREADS=15, ITEMS_PER_THREAD=1, LOAD_ALGORITHM=cub::BLOCK_LOAD_STRIPED, STORE_ALGORITHM=cub::BLOCK_STORE_STRIPED, InputIteratorT=const char *, OutputIteratorT=char *]" 
../dependencies/cub/test/test_block_load_store.cu(261): here
            instantiation of "void TestNative<T,BLOCK_THREADS,ITEMS_PER_THREAD,LOAD_ALGORITHM,STORE_ALGORITHM>(int, float, cub::Int2Type<1>) [with T=char, BLOCK_THREADS=15, ITEMS_PER_THREAD=1, LOAD_ALGORITHM=cub::BLOCK_LOAD_STRIPED, STORE_ALGORITHM=cub::BLOCK_STORE_STRIPED]" 
../dependencies/cub/test/test_block_load_store.cu(403): here
            instantiation of "void TestPointerType<T,BLOCK_THREADS,ITEMS_PER_THREAD,LOAD_ALGORITHM,STORE_ALGORITHM>(int, float) [with T=char, BLOCK_THREADS=15, ITEMS_PER_THREAD=1, LOAD_ALGORITHM=cub::BLOCK_LOAD_STRIPED, STORE_ALGORITHM=cub::BLOCK_STORE_STRIPED]" 
../dependencies/cub/test/test_block_load_store.cu(442): here
            instantiation of "void TestStrategy<T,BLOCK_THREADS,ITEMS_PER_THREAD>(int, float, cub::Int2Type<0>) [with T=char, BLOCK_THREADS=15, ITEMS_PER_THREAD=1]" 
../dependencies/cub/test/test_block_load_store.cu(476): here
            instantiation of "void TestItemsPerThread<T,BLOCK_THREADS>(int, float) [with T=char, BLOCK_THREADS=15]" 
../dependencies/cub/test/test_block_load_store.cu(491): here
            instantiation of "void TestThreads<T>(int, float) [with T=char]" 
../dependencies/cub/test/test_block_load_store.cu(534): here

mnicely · 2021-06-08T17:15:22Z

@canonizer I didn't need those for my use case, but I might not be hard to add...
@allisonvacanti If you have not objections, I'll see if I can get everything in today.

mnicely · 2021-06-08T17:51:10Z

I'm getting a bunch of errors trying to build CUB and test this go-around. Maybe it's best we wait till after 1.13 has been released and revisit.

...
/home/belt/workStuff/git_examples/cub/thrust/dependencies/cub/cub/iterator/tex_ref_input_iterator.cuh(133): error #1215-D: function "tex1Dfetch(texture<T, 1, cudaReadModeElementType>, int) [with T=uint2]"
/usr/local/cuda/bin/../targets/x86_64-linux/include/texture_fetch_functions.h(132): here was declared deprecated
          detected during:
            instantiation of "T cub::<unnamed>::IteratorTexRef<T>::TexId<UNIQUE_ID>::Fetch(Distance) [with T=TestBar, UNIQUE_ID=554, Distance=ptrdiff_t]" 
(299): here
            instantiation of "cub::TexRefInputIterator<T, UNIQUE_ID, OffsetT>::reference cub::TexRefInputIterator<T, UNIQUE_ID, OffsetT>::operator*() const [with T=TestBar, UNIQUE_ID=554, OffsetT=ptrdiff_t]" 
/home/belt/workStuff/git_examples/cub/thrust/dependencies/cub/test/test_iterator.cu(113): here
            instantiation of "void Kernel(InputIteratorT, T *, InputIteratorT *) [with InputIteratorT=cub::TexRefInputIterator<TestBar, 554, ptrdiff_t>, T=TestBar]" 
/home/belt/workStuff/git_examples/cub/thrust/dependencies/cub/test/test_iterator.cu(160): here
            instantiation of "void Test(InputIteratorT, T (&)[TEST_VALUES]) [with InputIteratorT=cub::TexRefInputIterator<TestBar, 554, ptrdiff_t>, T=TestBar, TEST_VALUES=8]" 
/home/belt/workStuff/git_examples/cub/thrust/dependencies/cub/test/test_iterator.cu(561): here
            instantiation of "void TestTexRef<T,CastT>() [with T=TestBar, CastT=TestBar]" 
/home/belt/workStuff/git_examples/cub/thrust/dependencies/cub/test/test_iterator.cu(694): here
            instantiation of "void Test<T,CastT>(cub::Int2Type<0>) [with T=TestBar, CastT=TestBar]" 
/home/belt/workStuff/git_examples/cub/thrust/dependencies/cub/test/test_iterator.cu(726): here
            instantiation of "void Test<T>() [with T=TestBar]" 
/home/belt/workStuff/git_examples/cub/thrust/dependencies/cub/test/test_iterator.cu(797): here
...
60 errors detected in the compilation of "/home/belt/workStuff/git_examples/cub/thrust/dependencies/cub/test/test_iterator.cu".
make[2]: *** [dependencies/cub/test/CMakeFiles/cub.cpp14.test.iterator.dir/build.make:76: dependencies/cub/test/CMakeFiles/cub.cpp14.test.iterator.dir/test_iterator.cu.o] Error 1
make[1]: *** [CMakeFiles/Makefile2:18203: dependencies/cub/test/CMakeFiles/cub.cpp14.test.iterator.dir/all] Error 2

alliepiper · 2021-06-08T17:52:55Z

Hmm, strange -- that error should be suppressed by our build system. It's just a deprecation warning and can be ignored.

What platform are you seeing that on?

mnicely · 2021-06-08T17:54:15Z

I'm using CTK 11.3U1 and Ubuntu 20.04.

I've merged in the latest from main into my branch and I'm using the updated build instructions.

mnicely · 2021-06-08T17:55:54Z

Should one of these be ON?

THRUST_IGNORE_CUB_VERSION_CHECK:BOOL=OFF
THRUST_IGNORE_DEPRECATED_COMPILER:BOOL=OFF
THRUST_IGNORE_DEPRECATED_CPP_11:BOOL=OFF
THRUST_IGNORE_DEPRECATED_CPP_DIALECT:BOOL=OFF

alliepiper · 2021-06-08T18:00:03Z

No, those affect other things. I'll look into it, but that's a known issue that shouldn't hold things up, and 1.13 probably won't fix it.

Let me know if you are planning to add the other load algorithms Andy suggested, or if you want to leave that as future work. If you're finished with this patch for now, I'll restart the tests -- our CI seems to suppress the warning as intended, so we can just validate there.

mnicely · 2021-06-08T18:03:13Z

I think it'll be easier to hold off on BLOCK_*_WARP_STRIDED.

We can revisit it later. I'll create a new PR and branch off 1.13 when it's released.

alliepiper · 2021-06-08T18:15:57Z

Sounds good 👍🏻 I'll go ahead and start the tests.

alliepiper

Needs a bit more work -- @mnicely can you add the missing member variable?

TempStorage temp_storage;

needs to be added as a member to the new StoreInternal specialization.

alliepiper · 2021-06-08T18:28:57Z

cub/block/block_store.cuh

+            int                 valid_items)                ///< [in] Number of valid items to write
+        {
+            if (linear_tid == 0)
+                temp_storage.valid_items = valid_items;     // Move through volatile smem as a workaround to prevent RF spilling on subsequent loads


That last commit only added an initialization to the constructor, there still needs to be a temp_storage member variable defined.

alliepiper · 2021-06-09T16:50:00Z

Thanks! This looks good to me, I should be able to land it in time for 1.13 👍

alliepiper · 2021-06-09T19:29:19Z

DVS CL: 30058071
gpuCI: NVIDIA/thrust#1427

alliepiper · 2021-06-11T13:47:43Z

Tests look good, merging.

Add BLOCK_LOAD_STRIPED and BLOCK_STORE_STRIPED

1c689db

mnicely requested a review from alliepiper March 25, 2021 13:23

mnicely self-assigned this Mar 25, 2021

mnicely changed the title ~~Add BLOCK_LOAD_STRIPED and BLOCK_STORE_STRIPED~~ [WIP] Add BLOCK_LOAD_STRIPED and BLOCK_STORE_STRIPED Mar 25, 2021

alliepiper added the type: enhancement New feature or request. label Mar 26, 2021

alliepiper added this to the 1.13.0 milestone Mar 26, 2021

Added test

a5e97ed

mnicely changed the title ~~[WIP] Add BLOCK_LOAD_STRIPED and BLOCK_STORE_STRIPED~~ Add BLOCK_LOAD_STRIPED and BLOCK_STORE_STRIPED Mar 26, 2021

brycelelbach added the P1: should have Necessary, but not critical. label Mar 29, 2021

alliepiper assigned alliepiper and unassigned mnicely May 10, 2021

alliepiper added a commit to alliepiper/thrust that referenced this pull request May 10, 2021

Testing branch for NVIDIA/cub#274.

7ed6afa

alliepiper added testing: internal ci in progress Currently testing on internal NVIDIA CI (DVS). testing: gpuCI in progress Started gpuCI testing. labels May 10, 2021

alliepiper assigned mnicely and unassigned alliepiper May 12, 2021

alliepiper removed testing: internal ci in progress Currently testing on internal NVIDIA CI (DVS). testing: gpuCI in progress Started gpuCI testing. labels May 12, 2021

alliepiper suggested changes Jun 8, 2021

View reviewed changes

mnicely added 2 commits June 8, 2021 15:20

Add temp_storage as member

1ed7ee6

Changes per Allison's suggestions

f622bf3

alliepiper assigned alliepiper and unassigned mnicely Jun 9, 2021

Fix some compile issues.

f9fbcbe

alliepiper added a commit to alliepiper/thrust that referenced this pull request Jun 9, 2021

Testing branch for NVIDIA/cub#274.

7fe901c

alliepiper added testing: internal ci in progress Currently testing on internal NVIDIA CI (DVS). testing: gpuCI in progress Started gpuCI testing. labels Jun 9, 2021

alliepiper added testing: internal ci passed Passed internal NVIDIA CI (DVS). testing: gpuCI passed Passed gpuCI testing. and removed testing: internal ci in progress Currently testing on internal NVIDIA CI (DVS). testing: gpuCI in progress Started gpuCI testing. labels Jun 11, 2021

alliepiper approved these changes Jun 11, 2021

View reviewed changes

alliepiper merged commit d056a9a into NVIDIA:main Jun 11, 2021

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Add BLOCK_LOAD_STRIPED and BLOCK_STORE_STRIPED #274

Add BLOCK_LOAD_STRIPED and BLOCK_STORE_STRIPED #274

mnicely commented Mar 25, 2021 •

edited

Loading

mnicely commented Mar 25, 2021 •

edited

Loading

alliepiper commented Mar 26, 2021

mnicely commented Mar 26, 2021

alliepiper commented May 10, 2021

alliepiper commented May 10, 2021

alliepiper commented May 12, 2021

mnicely commented Jun 8, 2021

mnicely commented Jun 8, 2021

alliepiper commented Jun 8, 2021

mnicely commented Jun 8, 2021

mnicely commented Jun 8, 2021

alliepiper commented Jun 8, 2021

mnicely commented Jun 8, 2021

alliepiper commented Jun 8, 2021

alliepiper left a comment

alliepiper Jun 8, 2021

alliepiper commented Jun 9, 2021

alliepiper commented Jun 9, 2021 •

edited

Loading

alliepiper commented Jun 11, 2021

Add BLOCK_LOAD_STRIPED and BLOCK_STORE_STRIPED #274

Add BLOCK_LOAD_STRIPED and BLOCK_STORE_STRIPED #274

Conversation

mnicely commented Mar 25, 2021 • edited Loading

mnicely commented Mar 25, 2021 • edited Loading

alliepiper commented Mar 26, 2021

mnicely commented Mar 26, 2021

alliepiper commented May 10, 2021

alliepiper commented May 10, 2021

alliepiper commented May 12, 2021

mnicely commented Jun 8, 2021

mnicely commented Jun 8, 2021

alliepiper commented Jun 8, 2021

mnicely commented Jun 8, 2021

mnicely commented Jun 8, 2021

alliepiper commented Jun 8, 2021

mnicely commented Jun 8, 2021

alliepiper commented Jun 8, 2021

alliepiper left a comment

Choose a reason for hiding this comment

alliepiper Jun 8, 2021

Choose a reason for hiding this comment

alliepiper commented Jun 9, 2021

alliepiper commented Jun 9, 2021 • edited Loading

alliepiper commented Jun 11, 2021

mnicely commented Mar 25, 2021 •

edited

Loading

mnicely commented Mar 25, 2021 •

edited

Loading

alliepiper commented Jun 9, 2021 •

edited

Loading