NVIDIA / NVIDIA/cccl

[DOC]: code example from https://nvidia.github.io/cccl/cub/ doesn't compile

Open
#1,542 0 comments 0 reactions 0 assignees View on GitHub
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

Open the contributing guide

Assessment

This issue has not been assessed yet.

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.