CUB มอบส่วนประกอบซอฟต์แวร์ที่ล้ำสมัยและสามารถนำมาใช้ซ้ำได้สำหรับทุกเลเยอร์ของโมเดลการเขียนโปรแกรม CUDA:
CUB รวมอยู่ใน NVIDIA HPC SDK และ CUDA Toolkit
เราขอแนะนำเว็บไซต์โครงการ CUB สำหรับข้อมูลเพิ่มเติมและตัวอย่าง
# include < cub/cub.cuh >
// Block-sorting CUDA kernel
__global__ void BlockSortKernel ( int *d_in, int *d_out)
{
using namespace cub ;
// Specialize BlockRadixSort, BlockLoad, and BlockStore for 128 threads
// owning 16 integer items each
typedef BlockRadixSort< int , 128 , 16 > BlockRadixSort;
typedef BlockLoad< int , 128 , 16 , BLOCK_LOAD_TRANSPOSE> BlockLoad;
typedef BlockStore< int , 128 , 16 , BLOCK_STORE_TRANSPOSE> BlockStore;
// Allocate shared memory
__shared__ union {
typename BlockRadixSort::TempStorage sort;
typename BlockLoad::TempStorage load;
typename BlockStore::TempStorage store;
} temp_storage;
int block_offset = blockIdx. x * ( 128 * 16 ); // OffsetT for this block's ment
// Obtain a segment of 2048 consecutive keys that are blocked across threads
int thread_keys[ 16 ];
BlockLoad (temp_storage. load ). Load (d_in + block_offset, thread_keys);
__syncthreads ();
// Collectively sort the keys
BlockRadixSort (temp_storage. sort ). Sort (thread_keys);
__syncthreads ();
// Store the sorted segment
BlockStore (temp_storage. store ). Store (d_out + block_offset, thread_keys);
}
แต่ละบล็อกเธรดใช้ cub::BlockRadixSort
เพื่อเรียงลำดับส่วนอินพุตของตนเองโดยรวม คลาสนี้มีความเชี่ยวชาญเฉพาะตามประเภทข้อมูลที่ถูกจัดเรียง ตามจำนวนเธรดต่อบล็อก ตามจำนวนคีย์ต่อเธรด และโดยนัยตามสถาปัตยกรรมการคอมไพล์เป้าหมาย
คลาส cub::BlockLoad
และ cub::BlockStore
มีความเชี่ยวชาญในทำนองเดียวกัน นอกจากนี้ เพื่อให้การเข้าถึงหน่วยความจำอุปกรณ์แบบรวม พื้นฐานเหล่านี้ได้รับการกำหนดค่าให้เข้าถึงหน่วยความจำโดยใช้รูปแบบการเข้าถึงแบบสไทรพ์ (โดยที่เธรดที่ต่อเนื่องกันเข้าถึงรายการต่อเนื่องกัน) จากนั้นจึง ย้าย คีย์เป็นการ จัดเรียงองค์ประกอบที่ถูกบล็อก ข้ามเธรด
เมื่อเชี่ยวชาญแล้ว คลาสเหล่านี้จะเปิดเผยประเภทสมาชิก TempStorage
ที่ทึบแสง บล็อกเธรดใช้ประเภทการจัดเก็บข้อมูลเหล่านี้เพื่อจัดสรรการรวมหน่วยความจำที่ใช้ร่วมกันแบบคงที่ซึ่งบล็อกเธรดต้องการ (หรืออีกทางหนึ่งประเภทการจัดเก็บข้อมูลเหล่านี้อาจเป็นนามแฝงของการจัดสรรหน่วยความจำส่วนกลาง)
CUB ได้รับการทดสอบเป็นประจำโดยใช้เวอร์ชันที่ระบุของคอมไพเลอร์ต่อไปนี้ เวอร์ชันที่ไม่รองรับอาจปล่อยคำเตือนการเลิกใช้งาน ซึ่งสามารถปิดเสียงได้โดยการกำหนด CUB_IGNORE_DEPRECATED_COMPILER ระหว่างการคอมไพล์
CUB ได้รับการเผยแพร่พร้อมกับ NVIDIA HPC SDK และ CUDA Toolkit นอกเหนือจาก GitHub
ดูบันทึกการเปลี่ยนแปลงสำหรับรายละเอียดเกี่ยวกับรุ่นเฉพาะ
การเปิดตัวของ CUB | รวมไว้ใน |
---|---|
2.0.1 | ชุดเครื่องมือ CUDA 12.0 |
2.0.0 | จะแจ้งภายหลัง |
1.17.2 | จะแจ้งภายหลัง |
1.17.1 | จะแจ้งภายหลัง |
1.17.0 | จะแจ้งภายหลัง |
1.16.0 | จะแจ้งภายหลัง |
1.15.0 | NVIDIA HPC SDK 22.1 และชุดเครื่องมือ CUDA 11.6 |
1.14.0 | NVIDIA HPC SDK 21.9 |
1.13.1 | ชุดเครื่องมือ CUDA 11.5 |
1.13.0 | NVIDIA HPC SDK 21.7 |
1.12.1 | ชุดเครื่องมือ CUDA 11.4 |
1.12.0 | NVIDIA HPC SDK21.3 |
1.11.0 | ชุดเครื่องมือ CUDA 11.3 |
1.10.0 | NVIDIA HPC SDK 20.9 และชุดเครื่องมือ CUDA 11.2 |
1.9.10-1 | NVIDIA HPC SDK 20.7 และชุดเครื่องมือ CUDA 11.1 |
1.9.10 | NVIDIA HPC SDK20.5 |
1.9.9 | ชุดเครื่องมือ CUDA 11.0 |
1.9.8-1 | NVIDIA HPC SDK 20.3 |
1.9.8 | ชุดเครื่องมือ CUDA 11.0 การเข้าถึงก่อน |
1.9.8 | CUDA 11.0 การเข้าถึงล่วงหน้า |
1.8.0 | |
1.7.5 | แรงขับ 1.9.2 |
1.7.4 | แรงขับ 1.9.1-2 |
1.7.3 | |
1.7.2 | |
1.7.1 | |
1.7.0 | แรงขับ 1.9.0-5 |
1.6.4 | |
1.6.3 | |
1.6.2 (ก่อนหน้า 1.5.5) | |
1.6.1 (ก่อนหน้า 1.5.4) | |
1.6.0 (ก่อนหน้า 1.5.3) | |
1.5.2 | |
1.5.1 | |
1.5.0 | |
1.4.1 | |
1.4.0 | |
1.3.2 | |
1.3.1 | |
1.3.0 | |
1.2.3 | |
1.2.2 | |
1.2.0 | |
1.1.1 | |
1.0.2 | |
1.0.1 | |
0.9.4 | |
0.9.2 | |
0.9.1 | |
0.9.0 |
CUB และ Thrust ขึ้นอยู่กับแต่ละอื่น ๆ ขอแนะนำให้โคลน Thrust และสร้าง CUB เป็นส่วนประกอบของ Thrust
CUB ใช้ระบบการสร้าง CMake เพื่อสร้างการทดสอบหน่วย ตัวอย่าง และการทดสอบส่วนหัว หากต้องการสร้าง CUB ในฐานะนักพัฒนา ควรปฏิบัติตามสูตรต่อไปนี้:
# Clone Thrust and CUB from Github. CUB is located in Thrust's
# `dependencies/cub` submodule.
git clone --recursive https://github.com/NVIDIA/thrust.git
cd thrust
# Create build directory:
mkdir build
cd build
# Configure -- use one of the following:
cmake -DTHRUST_INCLUDE_CUB_CMAKE=ON .. # Command line interface.
ccmake -DTHRUST_INCLUDE_CUB_CMAKE=ON .. # ncurses GUI (Linux only)
cmake-gui # Graphical UI, set source/build directories and options in the app
# Build:
cmake --build . -j < num jobs > # invokes make (or ninja, etc)
# Run tests and examples:
ctest
ตามค่าเริ่มต้น มาตรฐาน C++14 จะถูกกำหนดเป้าหมาย แต่สามารถเปลี่ยนแปลงได้ใน CMake ข้อมูลเพิ่มเติมเกี่ยวกับการกำหนดค่ารุ่น CUB ของคุณและการสร้างคำขอดึงมีอยู่ใน CONTRIBUTING.md
CUB มีให้บริการภายใต้ใบอนุญาตโอเพ่นซอร์ส "New BSD":
Copyright (c) 2010-2011, Duane Merrill. All rights reserved.
Copyright (c) 2011-2018, NVIDIA CORPORATION. All rights reserved.
Redistribution and use in source and binary forms, with or without
modification, are permitted provided that the following conditions are met:
* Redistributions of source code must retain the above copyright
notice, this list of conditions and the following disclaimer.
* Redistributions in binary form must reproduce the above copyright
notice, this list of conditions and the following disclaimer in the
documentation and/or other materials provided with the distribution.
* Neither the name of the NVIDIA CORPORATION nor the
names of its contributors may be used to endorse or promote products
derived from this software without specific prior written permission.
THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
(INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
(INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.