[DOC]: code example from https://nvidia.github.io/cccl/cub/ doesn't compile
- Dominant language
- C++
- Stars
- 2.5k
- Forks
- 487
- Avg merge
- 2d 7h
- Merged PRs (30d)
- 296
Description
### Is this a duplicate?
- [X] I confirmed there appear to be no [duplicate issues](https://github.com/NVIDIA/cccl/issues) for this bug and that I agree to the [Code of Conduct](CODE_OF_CONDUCT.md)
### Is this for new documentation, or an update to existing docs?
Update
### Describe the incorrect/future/missing documentation
The code provided in https://nvidia.github.io/cccl/cub/ doesn't compile:
```cuda
// main.cu
#include
//
// Block-sorting CUDA kernel
//
template
__global__ void BlockSortKernel(int *d_in, int *d_out)
{
// Specialize BlockLoad, BlockStore, and BlockRadixSort collective types
typedef cub::BlockLoad<
int*, BLOCK_THREADS, ITEMS_PER_THREAD, BLOCK_LOAD_TRANSPOSE> BlockLoadT;
typedef cub::BlockStore<
int*, BLOCK_THREADS, ITEMS_PER_THREAD, BLOCK_STORE_TRANSPOSE> BlockStoreT;
typedef cub::BlockRadixSort<
int, BLOCK_THREADS, ITEMS_PER_THREAD> BlockRadixSortT;
// Allocate type-safe, repurposable shared memory for collectives
__shared__ union {
typename BlockLoadT::TempStorage load;
typename BlockStoreT::TempStorage store;
typename BlockRadixSortT::TempStorage sort;
} temp_storage;
// Obtain this block's segment of consecutive keys (blocked across threads)
int thread_keys[ITEMS_PER_THREAD];
int block_offset = blockIdx.x * (BLOCK_THREADS * ITEMS_PER_THREAD);
BlockLoadT(temp_storage.load).Load(d_in + block_offset, thread_keys);
__syncthreads(); // Barrier for smem reuse
// Collectively sort the keys
BlockRadixSortT(temp_storage.sort).Sort(thread_keys);
__syncthreads(); // Barrier for smem reuse
// Store the sorted segment
BlockStoreT(temp_storage.store).Store(d_out + block_offset, thread_keys);
}
int main() {
return 0;
}
```
When compiling, I got error:
1. `identifier "BLOCK_LOAD_TRANSPOSE" is undefined` from line
```cuda
int*, BLOCK_THREADS, ITEMS_PER_THREAD, BLOCK_LOAD_TRANSPOSE> BlockLoadT;
```
2. `identifier "BLOCK_STORE_TRANSPOSE" is undefined` from line
```cuda
int*, BLOCK_THREADS, ITEMS_PER_THREAD, BLOCK_STORE_TRANSPOSE> BlockStoreT;
```
CMakeList.txt:
```
cmake_minimum_required(VERSION 3.18 FATAL_ERROR)
if (NOT CMAKE_CUDA_COMPILER)
set(CMAKE_CUDA_COMPILER "/usr/local/cuda/bin/nvcc")
# required by CLion
endif ()
set(CMAKE_CXX_STANDARD 17)
set(CMAKE_CXX_STANDARD_REQUIRED ON)
set(PROJ_NAME "cub-test")
project(${PROJ_NAME} LANGUAGES CUDA CXX)
add_executable(${PROJ_NAME} main.cu)
set_property(TARGET ${PROJ_NAME} PROPERTY CUDA_ARCHITECTURES native)
target_compile_options(${PROJ_NAME} PRIVATE $<$:
--expt-relaxed-constexpr
>)
```
hardware:
```
NVIDIA GeForce GTX 1650
compute capability: 7.5
SM arch: 750
runtime: 12.3.0
```
### If this is a correction, please provide a link to the incorrect documentation. If this is a new documentation request, please link to where you have looked.
https://nvidia.github.io/cccl/cub/
Contributor guide
Assessment
This issue has not been assessed yet.