mirror of
https://github.com/NVIDIA/cuda-samples.git
synced 2026-10-23 21:28:17 +08:00
add and update samples with CUDA 11.3 support
This commit is contained in:
Rutwik Choughule
@@ -30,7 +30,7 @@ cudaMalloc, cudaFree, cudaMemcpyAsync
|
||||
|
||||
## Prerequisites
|
||||
|
||||
Download and install the [CUDA Toolkit 11.2](https://developer.nvidia.com/cuda-downloads) for your corresponding platform.
|
||||
Download and install the [CUDA Toolkit 11.3](https://developer.nvidia.com/cuda-downloads) for your corresponding platform.
|
||||
Make sure the dependencies mentioned in [Dependencies]() section above are installed.
|
||||
|
||||
## Build and Run
|
||||
|
||||
@@ -25,7 +25,6 @@
|
||||
* OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
|
||||
*/
|
||||
|
||||
|
||||
// Includes, system
|
||||
#include <stdio.h>
|
||||
|
||||
@@ -35,229 +34,222 @@
|
||||
#include <cooperative_groups.h>
|
||||
|
||||
// Utilities and timing functions
|
||||
#include <helper_functions.h> // includes cuda.h and cuda_runtime_api.h
|
||||
#include <helper_functions.h> // includes cuda.h and cuda_runtime_api.h
|
||||
|
||||
// CUDA helper functions
|
||||
#include <helper_cuda.h> // helper functions for CUDA error check
|
||||
#include <helper_cuda.h> // helper functions for CUDA error check
|
||||
|
||||
namespace cg = cooperative_groups;
|
||||
|
||||
|
||||
#if __CUDA_ARCH__ >= 700
|
||||
template <bool writeSquareRoot> __device__ void reduceBlockData(cuda::barrier<cuda::thread_scope_block> &barrier,
|
||||
cg::thread_block_tile<32> &tile32, double &threadSum, double *result)
|
||||
{
|
||||
extern __shared__ double tmp[];
|
||||
template <bool writeSquareRoot>
|
||||
__device__ void reduceBlockData(
|
||||
cuda::barrier<cuda::thread_scope_block> &barrier,
|
||||
cg::thread_block_tile<32> &tile32, double &threadSum, double *result) {
|
||||
extern __shared__ double tmp[];
|
||||
|
||||
#pragma unroll
|
||||
for (int offset = tile32.size()/2; offset > 0; offset /= 2)
|
||||
{
|
||||
threadSum += tile32.shfl_down(threadSum, offset);
|
||||
}
|
||||
if (tile32.thread_rank() == 0)
|
||||
{
|
||||
tmp[tile32.meta_group_rank()] = threadSum;
|
||||
#pragma unroll
|
||||
for (int offset = tile32.size() / 2; offset > 0; offset /= 2) {
|
||||
threadSum += tile32.shfl_down(threadSum, offset);
|
||||
}
|
||||
if (tile32.thread_rank() == 0) {
|
||||
tmp[tile32.meta_group_rank()] = threadSum;
|
||||
}
|
||||
|
||||
auto token = barrier.arrive();
|
||||
|
||||
barrier.wait(std::move(token));
|
||||
|
||||
// The warp 0 will perform last round of reduction
|
||||
if (tile32.meta_group_rank() == 0) {
|
||||
double beta = tile32.thread_rank() < tile32.meta_group_size()
|
||||
? tmp[tile32.thread_rank()]
|
||||
: 0.0;
|
||||
|
||||
#pragma unroll
|
||||
for (int offset = tile32.size() / 2; offset > 0; offset /= 2) {
|
||||
beta += tile32.shfl_down(beta, offset);
|
||||
}
|
||||
|
||||
auto token = barrier.arrive();
|
||||
|
||||
barrier.wait(std::move(token));
|
||||
|
||||
// The warp 0 will perform last round of reduction
|
||||
if (tile32.meta_group_rank() == 0) {
|
||||
|
||||
double beta = tile32.thread_rank() < tile32.meta_group_size() ? tmp[tile32.thread_rank()] : 0.0;
|
||||
|
||||
#pragma unroll
|
||||
for (int offset = tile32.size()/2; offset > 0; offset /= 2)
|
||||
{
|
||||
beta += tile32.shfl_down(beta, offset);
|
||||
}
|
||||
|
||||
if (tile32.thread_rank() == 0)
|
||||
{
|
||||
if (writeSquareRoot)
|
||||
*result = sqrt(beta);
|
||||
else
|
||||
*result = beta;
|
||||
}
|
||||
if (tile32.thread_rank() == 0) {
|
||||
if (writeSquareRoot)
|
||||
*result = sqrt(beta);
|
||||
else
|
||||
*result = beta;
|
||||
}
|
||||
}
|
||||
}
|
||||
#endif
|
||||
|
||||
__global__ void normVecByDotProductAWBarrier(float *vecA, float *vecB, double *partialResults, int size)
|
||||
{
|
||||
__global__ void normVecByDotProductAWBarrier(float *vecA, float *vecB,
|
||||
double *partialResults, int size) {
|
||||
#if __CUDA_ARCH__ >= 700
|
||||
#pragma diag_suppress static_var_with_dynamic_init
|
||||
cg::thread_block cta = cg::this_thread_block();
|
||||
cg::grid_group grid = cg::this_grid();;
|
||||
cg::thread_block_tile<32> tile32 = cg::tiled_partition<32>(cta);
|
||||
cg::thread_block cta = cg::this_thread_block();
|
||||
cg::grid_group grid = cg::this_grid();
|
||||
;
|
||||
cg::thread_block_tile<32> tile32 = cg::tiled_partition<32>(cta);
|
||||
|
||||
__shared__ cuda::barrier<cuda::thread_scope_block> barrier;
|
||||
__shared__ cuda::barrier<cuda::thread_scope_block> barrier;
|
||||
|
||||
if (threadIdx.x == 0) {
|
||||
init(&barrier, blockDim.x);
|
||||
if (threadIdx.x == 0) {
|
||||
init(&barrier, blockDim.x);
|
||||
}
|
||||
|
||||
cg::sync(cta);
|
||||
|
||||
double threadSum = 0.0;
|
||||
for (int i = grid.thread_rank(); i < size; i += grid.size()) {
|
||||
threadSum += (double)(vecA[i] * vecB[i]);
|
||||
}
|
||||
|
||||
// Each thread block performs reduction of partial dotProducts and writes to
|
||||
// global mem.
|
||||
reduceBlockData<false>(barrier, tile32, threadSum,
|
||||
&partialResults[blockIdx.x]);
|
||||
|
||||
cg::sync(grid);
|
||||
|
||||
// One block performs the final summation of partial dot products
|
||||
// of all the thread blocks and writes the sqrt of final dot product.
|
||||
if (blockIdx.x == 0) {
|
||||
threadSum = 0.0;
|
||||
for (int i = cta.thread_rank(); i < gridDim.x; i += cta.size()) {
|
||||
threadSum += partialResults[i];
|
||||
}
|
||||
reduceBlockData<true>(barrier, tile32, threadSum, &partialResults[0]);
|
||||
}
|
||||
|
||||
cg::sync(cta);
|
||||
cg::sync(grid);
|
||||
|
||||
double threadSum = 0.0;
|
||||
for (int i = grid.thread_rank(); i < size; i += grid.size())
|
||||
{
|
||||
threadSum += (double) (vecA[i] * vecB[i]);
|
||||
}
|
||||
const double finalValue = partialResults[0];
|
||||
|
||||
// Each thread block performs reduction of partial dotProducts and writes to
|
||||
// global mem.
|
||||
reduceBlockData<false>(barrier, tile32, threadSum, &partialResults[blockIdx.x]);
|
||||
|
||||
cg::sync(grid);
|
||||
|
||||
// One block performs the final summation of partial dot products
|
||||
// of all the thread blocks and writes the sqrt of final dot product.
|
||||
if (blockIdx.x == 0)
|
||||
{
|
||||
threadSum = 0.0;
|
||||
for (int i = cta.thread_rank(); i < gridDim.x; i += cta.size())
|
||||
{
|
||||
threadSum += partialResults[i];
|
||||
}
|
||||
reduceBlockData<true>(barrier, tile32, threadSum, &partialResults[0]);
|
||||
}
|
||||
|
||||
cg::sync(grid);
|
||||
|
||||
const double finalValue = partialResults[0];
|
||||
|
||||
// Perform normalization of vecA & vecB.
|
||||
for (int i = grid.thread_rank(); i < size; i += grid.size())
|
||||
{
|
||||
vecA[i] = (float)vecA[i] / finalValue;
|
||||
vecB[i] = (float)vecB[i] / finalValue;
|
||||
}
|
||||
// Perform normalization of vecA & vecB.
|
||||
for (int i = grid.thread_rank(); i < size; i += grid.size()) {
|
||||
vecA[i] = (float)vecA[i] / finalValue;
|
||||
vecB[i] = (float)vecB[i] / finalValue;
|
||||
}
|
||||
#endif
|
||||
}
|
||||
|
||||
|
||||
int runNormVecByDotProductAWBarrier(int argc, char **argv, int deviceId);
|
||||
|
||||
////////////////////////////////////////////////////////////////////////////////
|
||||
// Program main
|
||||
////////////////////////////////////////////////////////////////////////////////
|
||||
int main(int argc, char **argv)
|
||||
{
|
||||
printf("%s starting...\n", argv[0]);
|
||||
int main(int argc, char **argv) {
|
||||
printf("%s starting...\n", argv[0]);
|
||||
|
||||
// This will pick the best possible CUDA capable device
|
||||
int dev = findCudaDevice(argc, (const char **)argv);
|
||||
// This will pick the best possible CUDA capable device
|
||||
int dev = findCudaDevice(argc, (const char **)argv);
|
||||
|
||||
int major = 0;
|
||||
checkCudaErrors(cudaDeviceGetAttribute(&major, cudaDevAttrComputeCapabilityMajor, dev));
|
||||
int major = 0;
|
||||
checkCudaErrors(
|
||||
cudaDeviceGetAttribute(&major, cudaDevAttrComputeCapabilityMajor, dev));
|
||||
|
||||
// Arrive-Wait Barrier require a GPU of Volta (SM7X) architecture or higher.
|
||||
if (major < 7) {
|
||||
printf("simpleAWBarrier requires SM 7.0 or higher. Exiting...\n");
|
||||
exit(EXIT_WAIVED);
|
||||
}
|
||||
// Arrive-Wait Barrier require a GPU of Volta (SM7X) architecture or higher.
|
||||
if (major < 7) {
|
||||
printf("simpleAWBarrier requires SM 7.0 or higher. Exiting...\n");
|
||||
exit(EXIT_WAIVED);
|
||||
}
|
||||
|
||||
int supportsCooperativeLaunch = 0;
|
||||
checkCudaErrors(cudaDeviceGetAttribute(&supportsCooperativeLaunch, cudaDevAttrCooperativeLaunch, dev));
|
||||
int supportsCooperativeLaunch = 0;
|
||||
checkCudaErrors(cudaDeviceGetAttribute(&supportsCooperativeLaunch,
|
||||
cudaDevAttrCooperativeLaunch, dev));
|
||||
|
||||
if (!supportsCooperativeLaunch)
|
||||
{
|
||||
printf("\nSelected GPU (%d) does not support Cooperative Kernel Launch, Waiving the run\n", dev);
|
||||
exit(EXIT_WAIVED);
|
||||
}
|
||||
if (!supportsCooperativeLaunch) {
|
||||
printf(
|
||||
"\nSelected GPU (%d) does not support Cooperative Kernel Launch, "
|
||||
"Waiving the run\n",
|
||||
dev);
|
||||
exit(EXIT_WAIVED);
|
||||
}
|
||||
|
||||
int testResult = runNormVecByDotProductAWBarrier(argc, argv, dev);
|
||||
int testResult = runNormVecByDotProductAWBarrier(argc, argv, dev);
|
||||
|
||||
printf("%s completed, returned %s\n", argv[0], testResult ? "OK" : "ERROR!");
|
||||
exit(testResult ? EXIT_SUCCESS : EXIT_FAILURE);
|
||||
printf("%s completed, returned %s\n", argv[0], testResult ? "OK" : "ERROR!");
|
||||
exit(testResult ? EXIT_SUCCESS : EXIT_FAILURE);
|
||||
}
|
||||
|
||||
int runNormVecByDotProductAWBarrier(int argc, char **argv, int deviceId)
|
||||
{
|
||||
float *vecA, *d_vecA;
|
||||
float *vecB, *d_vecB;
|
||||
double *d_partialResults;
|
||||
int size = 10000000;
|
||||
int runNormVecByDotProductAWBarrier(int argc, char **argv, int deviceId) {
|
||||
float *vecA, *d_vecA;
|
||||
float *vecB, *d_vecB;
|
||||
double *d_partialResults;
|
||||
int size = 10000000;
|
||||
|
||||
vecA = new float[size];
|
||||
vecB = new float[size];
|
||||
checkCudaErrors(cudaMallocHost(&vecA, sizeof(float) * size));
|
||||
checkCudaErrors(cudaMallocHost(&vecB, sizeof(float) * size));
|
||||
|
||||
checkCudaErrors(cudaMalloc(&d_vecA, sizeof(float)*size));
|
||||
checkCudaErrors(cudaMalloc(&d_vecB, sizeof(float)*size));
|
||||
checkCudaErrors(cudaMalloc(&d_vecA, sizeof(float) * size));
|
||||
checkCudaErrors(cudaMalloc(&d_vecB, sizeof(float) * size));
|
||||
|
||||
float baseVal = 2.0;
|
||||
for (int i = 0; i < size; i++)
|
||||
{
|
||||
vecA[i] = vecB[i] = baseVal;
|
||||
float baseVal = 2.0;
|
||||
for (int i = 0; i < size; i++) {
|
||||
vecA[i] = vecB[i] = baseVal;
|
||||
}
|
||||
|
||||
cudaStream_t stream;
|
||||
checkCudaErrors(cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking));
|
||||
|
||||
checkCudaErrors(cudaMemcpyAsync(d_vecA, vecA, sizeof(float) * size,
|
||||
cudaMemcpyHostToDevice, stream));
|
||||
checkCudaErrors(cudaMemcpyAsync(d_vecB, vecB, sizeof(float) * size,
|
||||
cudaMemcpyHostToDevice, stream));
|
||||
|
||||
// Kernel configuration, where a one-dimensional
|
||||
// grid and one-dimensional blocks are configured.
|
||||
int minGridSize = 0, blockSize = 0;
|
||||
checkCudaErrors(cudaOccupancyMaxPotentialBlockSize(
|
||||
&minGridSize, &blockSize, (void *)normVecByDotProductAWBarrier, 0, size));
|
||||
|
||||
int smemSize = ((blockSize / 32) + 1) * sizeof(double);
|
||||
|
||||
int numBlocksPerSm = 0;
|
||||
checkCudaErrors(cudaOccupancyMaxActiveBlocksPerMultiprocessor(
|
||||
&numBlocksPerSm, normVecByDotProductAWBarrier, blockSize, smemSize));
|
||||
|
||||
int multiProcessorCount = 0;
|
||||
checkCudaErrors(cudaDeviceGetAttribute(
|
||||
&multiProcessorCount, cudaDevAttrMultiProcessorCount, deviceId));
|
||||
|
||||
minGridSize = multiProcessorCount * numBlocksPerSm;
|
||||
checkCudaErrors(cudaMalloc(&d_partialResults, minGridSize * sizeof(double)));
|
||||
|
||||
printf(
|
||||
"Launching normVecByDotProductAWBarrier kernel with numBlocks = %d "
|
||||
"blockSize = %d\n",
|
||||
minGridSize, blockSize);
|
||||
|
||||
dim3 dimGrid(minGridSize, 1, 1), dimBlock(blockSize, 1, 1);
|
||||
|
||||
void *kernelArgs[] = {(void *)&d_vecA, (void *)&d_vecB,
|
||||
(void *)&d_partialResults, (void *)&size};
|
||||
|
||||
checkCudaErrors(
|
||||
cudaLaunchCooperativeKernel((void *)normVecByDotProductAWBarrier, dimGrid,
|
||||
dimBlock, kernelArgs, smemSize, stream));
|
||||
|
||||
checkCudaErrors(cudaMemcpyAsync(vecA, d_vecA, sizeof(float) * size,
|
||||
cudaMemcpyDeviceToHost, stream));
|
||||
checkCudaErrors(cudaStreamSynchronize(stream));
|
||||
|
||||
float expectedResult = (baseVal / sqrt(size * baseVal * baseVal));
|
||||
unsigned int matches = 0;
|
||||
for (int i = 0; i < size; i++) {
|
||||
if ((vecA[i] - expectedResult) > 0.00001) {
|
||||
printf("mismatch at i = %d\n", i);
|
||||
break;
|
||||
} else {
|
||||
matches++;
|
||||
}
|
||||
}
|
||||
|
||||
cudaStream_t stream;
|
||||
checkCudaErrors(cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking));
|
||||
printf("Result = %s\n", matches == size ? "PASSED" : "FAILED");
|
||||
checkCudaErrors(cudaFree(d_vecA));
|
||||
checkCudaErrors(cudaFree(d_vecB));
|
||||
checkCudaErrors(cudaFree(d_partialResults));
|
||||
|
||||
checkCudaErrors(cudaMemcpyAsync(d_vecA, vecA, sizeof(float)*size, cudaMemcpyHostToDevice, stream));
|
||||
checkCudaErrors(cudaMemcpyAsync(d_vecB, vecB, sizeof(float)*size, cudaMemcpyHostToDevice, stream));
|
||||
|
||||
// Kernel configuration, where a one-dimensional
|
||||
// grid and one-dimensional blocks are configured.
|
||||
int minGridSize = 0, blockSize = 0;
|
||||
checkCudaErrors(cudaOccupancyMaxPotentialBlockSize(
|
||||
&minGridSize,
|
||||
&blockSize,
|
||||
(void*)normVecByDotProductAWBarrier,
|
||||
0,
|
||||
size));
|
||||
|
||||
int smemSize = ((blockSize/32)+1) * sizeof(double);
|
||||
|
||||
int numBlocksPerSm = 0;
|
||||
checkCudaErrors(cudaOccupancyMaxActiveBlocksPerMultiprocessor(&numBlocksPerSm, normVecByDotProductAWBarrier, blockSize, smemSize));
|
||||
|
||||
int multiProcessorCount = 0;
|
||||
checkCudaErrors(cudaDeviceGetAttribute(&multiProcessorCount, cudaDevAttrMultiProcessorCount, deviceId));
|
||||
|
||||
minGridSize = multiProcessorCount * numBlocksPerSm;
|
||||
checkCudaErrors(cudaMalloc(&d_partialResults, minGridSize*sizeof(double)));
|
||||
|
||||
printf("Launching normVecByDotProductAWBarrier kernel with numBlocks = %d blockSize = %d\n", minGridSize, blockSize);
|
||||
|
||||
dim3 dimGrid(minGridSize, 1, 1), dimBlock(blockSize, 1, 1);
|
||||
|
||||
void *kernelArgs[] = {
|
||||
(void*)&d_vecA,
|
||||
(void*)&d_vecB,
|
||||
(void*)&d_partialResults,
|
||||
(void*)&size
|
||||
};
|
||||
|
||||
checkCudaErrors(cudaLaunchCooperativeKernel((void *)normVecByDotProductAWBarrier, dimGrid, dimBlock, kernelArgs, smemSize, stream));
|
||||
|
||||
checkCudaErrors(cudaMemcpyAsync(vecA, d_vecA, sizeof(float)*size, cudaMemcpyDeviceToHost, stream));
|
||||
checkCudaErrors(cudaStreamSynchronize(stream));
|
||||
|
||||
float expectedResult = (baseVal / sqrt(size*baseVal*baseVal));
|
||||
unsigned int matches = 0;
|
||||
for (int i=0; i < size; i++)
|
||||
{
|
||||
if ((vecA[i] - expectedResult) > 0.00001)
|
||||
{
|
||||
printf("mismatch at i = %d\n", i);
|
||||
break;
|
||||
}
|
||||
else
|
||||
{
|
||||
matches++;
|
||||
}
|
||||
}
|
||||
|
||||
printf("Result = %s\n", matches == size ? "PASSED" : "FAILED");
|
||||
checkCudaErrors(cudaFree(d_vecA));
|
||||
checkCudaErrors(cudaFree(d_vecB));
|
||||
checkCudaErrors(cudaFree(d_partialResults));
|
||||
|
||||
delete[] vecA;
|
||||
delete[] vecB;
|
||||
return matches == size;
|
||||
checkCudaErrors(cudaFreeHost(vecA));
|
||||
checkCudaErrors(cudaFreeHost(vecB));
|
||||
return matches == size;
|
||||
}
|
||||
|
||||
@@ -38,7 +38,7 @@
|
||||
</PropertyGroup>
|
||||
<Import Project="$(VCTargetsPath)\Microsoft.Cpp.props" />
|
||||
<ImportGroup Label="ExtensionSettings">
|
||||
<Import Project="$(CUDAPropsPath)\CUDA 11.2.props" />
|
||||
<Import Project="$(CUDAPropsPath)\CUDA 11.3.props" />
|
||||
</ImportGroup>
|
||||
<ImportGroup Label="PropertySheets">
|
||||
<Import Condition="exists('$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props')" Label="LocalAppDataPlatform" Project="$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props" />
|
||||
@@ -108,6 +108,6 @@
|
||||
</ItemGroup>
|
||||
<Import Project="$(VCTargetsPath)\Microsoft.Cpp.targets" />
|
||||
<ImportGroup Label="ExtensionTargets">
|
||||
<Import Project="$(CUDAPropsPath)\CUDA 11.2.targets" />
|
||||
<Import Project="$(CUDAPropsPath)\CUDA 11.3.targets" />
|
||||
</ImportGroup>
|
||||
</Project>
|
||||
|
||||
@@ -34,7 +34,7 @@
|
||||
</PropertyGroup>
|
||||
<Import Project="$(VCTargetsPath)\Microsoft.Cpp.props" />
|
||||
<ImportGroup Label="ExtensionSettings">
|
||||
<Import Project="$(CUDAPropsPath)\CUDA 11.2.props" />
|
||||
<Import Project="$(CUDAPropsPath)\CUDA 11.3.props" />
|
||||
</ImportGroup>
|
||||
<ImportGroup Label="PropertySheets">
|
||||
<Import Condition="exists('$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props')" Label="LocalAppDataPlatform" Project="$(UserRootDir)\Microsoft.Cpp.$(Platform).user.props" />
|
||||
@@ -104,6 +104,6 @@
|
||||
</ItemGroup>
|
||||
<Import Project="$(VCTargetsPath)\Microsoft.Cpp.targets" />
|
||||
<ImportGroup Label="ExtensionTargets">
|
||||
<Import Project="$(CUDAPropsPath)\CUDA 11.2.targets" />
|
||||
<Import Project="$(CUDAPropsPath)\CUDA 11.3.targets" />
|
||||
</ImportGroup>
|
||||
</Project>
|
||||
|
||||
Reference in New Issue
Block a user