cuda-samples/Samples/6_Performance/UnifiedMemoryPerf/matrixMultiplyPerf.cu
2022-01-13 11:35:24 +05:30

698 lines
26 KiB
Plaintext

/* Copyright (c) 2022, 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 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 ``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 THE COPYRIGHT OWNER OR
* CONTRIBUTORS 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.
*/
#include <helper_cuda.h>
#include <helper_timer.h>
#include "commonDefs.hpp"
#include "commonKernels.hpp"
#define VERIFY_GPU_CORRECTNESS 0
size_t maxSampleSizeInMb = 64;
int numKernelRuns = 20;
int verboseResults = 0;
const char *memAllocTypeStr[MEMALLOC_TYPE_COUNT] = {
"Managed_Memory_With_Hints",
"Managed_Memory_With_Hints_FullyAsync",
"Managed_Memory_NoHints",
"Zero_Copy",
"Memcpy_HostMalloc_DeviceCudaMalloc",
"MemcpyAsync_HostMalloc_DeviceCudaMalloc",
"Memcpy_HostCudaHostAlloc_DeviceCudaMalloc",
"MemcpyAsync_HostCudaHostAlloc_DeviceCudaMalloc"};
const char *memAllocTypeShortStr[MEMALLOC_TYPE_COUNT] = {
"UMhint", // Managed Memory With Hints
"UMhntAs", // Managed Memory With_Hints Async
"UMeasy", // Managed_Memory with No Hints
"0Copy", // Zero Copy
"MemCopy", // USE HOST PAGEABLE AND DEVICE_MEMORY
"CpAsync", // USE HOST PAGEABLE AND DEVICE_MEMORY ASYNC
"CpHpglk", // USE HOST PAGELOCKED AND DEVICE MEMORY
"CpPglAs" // USE HOST PAGELOCKED AND DEVICE MEMORY ASYNC
};
static float RandFloat(float low, float high) {
float t = (float)rand() / (float)RAND_MAX;
return (1.0f - t) * low + t * high;
}
void fillMatrixWithRandomValues(float *matrix, unsigned int matrixDim) {
unsigned int i, j;
for (i = 0; i < matrixDim; ++i) {
for (j = 0; j < matrixDim; ++j) {
matrix[j + i * matrixDim] = RandFloat(0.0f, 10.0f);
}
}
}
#if VERIFY_GPU_CORRECTNESS
void verifyMatrixMultiplyCorrectness(float *C, float *A, float *B,
unsigned int matrixDim) {
unsigned int i, j, k, numErrors = 0;
for (i = 0; i < matrixDim; ++i) {
for (j = 0; j < matrixDim; ++j) {
float result = 0.0f;
for (k = 0; k < matrixDim; ++k) {
result += A[k + i * matrixDim] * B[j + k * matrixDim];
}
if (fabs(C[j + i * matrixDim] - result) > 0.001 * matrixDim) {
printf("At [%u, %u]: Expected %f, Found %f\n", i, j, result,
C[j + i * matrixDim]);
++numErrors;
}
}
}
if (numErrors != 0) {
printf("%d value mismatches occured\n", numErrors);
fflush(stdout);
exit(EXIT_FAILURE); // exit since value mismatches occured
}
}
#endif
void copyMatrix(float *dstMatrix, float *srcMatrix, unsigned int matrixDim) {
size_t size = matrixDim * matrixDim * sizeof(float);
memcpy(dstMatrix, srcMatrix, size);
}
void verifyMatrixData(float *expectedData, float *observedData,
unsigned int matrixDim) {
unsigned int i, j, numErrors = 0;
for (i = 0; i < matrixDim; ++i) {
for (j = 0; j < matrixDim; ++j) {
if (expectedData[j + i * matrixDim] != observedData[j + i * matrixDim]) {
++numErrors;
if (verboseResults) {
printf("At [%u, %u]: Expected %f, Found %f\n", i, j,
expectedData[j + i * matrixDim],
observedData[j + i * matrixDim]);
}
}
}
}
if (numErrors != 0) {
printf("%d value mismatches occured\n", numErrors);
fflush(stdout);
exit(EXIT_FAILURE); // exit since value mismatches occured
}
}
#define BLOCK_SIZE 32
__global__ void matrixMultiplyKernel(float *C, float *A, float *B,
unsigned int matrixDim) {
// Block index
int bx = blockIdx.x;
int by = blockIdx.y;
// Thread index
int tx = threadIdx.x;
int ty = threadIdx.y;
unsigned int wA = matrixDim;
unsigned int wB = matrixDim;
// Index of the first sub-matrix of A processed by the block
int aBegin = matrixDim * BLOCK_SIZE * by;
// Index of the last sub-matrix of A processed by the block
int aEnd = aBegin + wA - 1;
// Step size used to iterate through the sub-matrices of A
int aStep = BLOCK_SIZE;
// Index of the first sub-matrix of B processed by the block
int bBegin = BLOCK_SIZE * bx;
// Step size used to iterate through the sub-matrices of B
int bStep = BLOCK_SIZE * wB;
// Csub is used to store the element of the block sub-matrix
// that is computed by the thread
float Csub = 0;
// Loop over all the sub-matrices of A and B
// required to compute the block sub-matrix
for (int a = aBegin, b = bBegin; a <= aEnd; a += aStep, b += bStep) {
// Declaration of the shared memory array As used to
// store the sub-matrix of A
__shared__ float As[BLOCK_SIZE][BLOCK_SIZE];
// Declaration of the shared memory array Bs used to
// store the sub-matrix of B
__shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE];
// Load the matrices from device memory
// to shared memory; each thread loads
// one element of each matrix
As[ty][tx] = A[a + wA * ty + tx];
Bs[ty][tx] = B[b + wB * ty + tx];
// Synchronize to make sure the matrices are loaded
__syncthreads();
// Multiply the two matrices together;
// each thread computes one element
// of the block sub-matrix
#pragma unroll
for (int k = 0; k < BLOCK_SIZE; ++k) {
Csub += As[ty][k] * Bs[k][tx];
}
// Synchronize to make sure that the preceding
// computation is done before loading two new
// sub-matrices of A and B in the next iteration
__syncthreads();
}
// Write the block sub-matrix to device memory;
// each thread writes one element
int c = wB * BLOCK_SIZE * by + BLOCK_SIZE * bx;
C[c + wB * ty + tx] = Csub;
}
void runMatrixMultiplyKernel(unsigned int matrixDim, int allocType,
unsigned int numLoops, double *gpuLaunchCallsTimes,
double *gpuTransferToCallsTimes,
double *gpuTransferFromCallsTimes,
double *gpuLaunchAndTransferCallsTimes,
double *gpuLaunchTransferSyncTimes,
double *cpuAccessTimes, double *overallTimes,
int device_id) {
float *dptrA = NULL, *hptrA = NULL;
float *dptrB = NULL, *hptrB = NULL;
float *dptrC = NULL, *hptrC = NULL;
float *randValuesX = NULL, *randValuesY = NULL;
float *randValuesVerifyXmulY = NULL, *randValuesVerifyYmulX = NULL;
bool copyRequired = false, hintsRequired = false;
bool someTransferOpRequired;
bool isAsync = false;
cudaStream_t streamToRunOn;
unsigned int *latch;
size_t size = matrixDim * matrixDim * sizeof(float);
dim3 threads(32, 32);
dim3 grid(matrixDim / threads.x, matrixDim / threads.y);
StopWatchInterface *gpuLaunchCallsTimer = 0, *gpuTransferCallsTimer = 0;
StopWatchInterface *gpuSyncTimer = 0, *cpuAccessTimer = 0;
sdkCreateTimer(&gpuLaunchCallsTimer);
sdkCreateTimer(&gpuTransferCallsTimer);
sdkCreateTimer(&gpuSyncTimer);
sdkCreateTimer(&cpuAccessTimer);
unsigned int i;
cudaDeviceProp deviceProp;
checkCudaErrors(cudaGetDeviceProperties(&deviceProp, device_id));
checkCudaErrors(cudaStreamCreate(&streamToRunOn));
randValuesX = (float *)malloc(size);
if (!randValuesX) {
exit(EXIT_FAILURE); // exit since memory allocation error
}
randValuesY = (float *)malloc(size);
if (!randValuesY) {
exit(EXIT_FAILURE); // exit since memory allocation error
}
randValuesVerifyXmulY = (float *)malloc(size);
if (!randValuesVerifyXmulY) {
exit(EXIT_FAILURE); // exit since memory allocation error
}
randValuesVerifyYmulX = (float *)malloc(size);
if (!randValuesVerifyYmulX) {
exit(EXIT_FAILURE); // exit since memory allocation error
}
checkCudaErrors(cudaMalloc(&dptrA, size));
checkCudaErrors(cudaMalloc(&dptrB, size));
checkCudaErrors(cudaMalloc(&dptrC, size));
fillMatrixWithRandomValues(randValuesX, matrixDim);
fillMatrixWithRandomValues(randValuesY, matrixDim);
checkCudaErrors(
cudaMemcpyAsync(dptrA, randValuesX, size, cudaMemcpyHostToDevice));
checkCudaErrors(
cudaMemcpyAsync(dptrB, randValuesY, size, cudaMemcpyHostToDevice));
matrixMultiplyKernel<<<grid, threads>>>(dptrC, dptrA, dptrB, matrixDim);
checkCudaErrors(cudaMemcpyAsync(randValuesVerifyXmulY, dptrC, size,
cudaMemcpyDeviceToHost));
checkCudaErrors(cudaStreamSynchronize(NULL));
matrixMultiplyKernel<<<grid, threads>>>(dptrC, dptrB, dptrA, matrixDim);
checkCudaErrors(cudaMemcpyAsync(randValuesVerifyYmulX, dptrC, size,
cudaMemcpyDeviceToHost));
checkCudaErrors(cudaStreamSynchronize(NULL));
#if VERIFY_GPU_CORRECTNESS
verifyMatrixMultiplyCorrectness(randValuesVerifyXmulY, randValuesX,
randValuesY, matrixDim);
verifyMatrixMultiplyCorrectness(randValuesVerifyYmulX, randValuesY,
randValuesX, matrixDim);
#endif
checkCudaErrors(cudaFree(dptrA));
checkCudaErrors(cudaFree(dptrB));
checkCudaErrors(cudaFree(dptrC));
checkCudaErrors(cudaMallocHost(&latch, sizeof(unsigned int)));
switch (allocType) {
case USE_HOST_PAGEABLE_AND_DEVICE_MEMORY:
case USE_HOST_PAGEABLE_AND_DEVICE_MEMORY_ASYNC:
hptrA = (float *)malloc(size);
if (!hptrA) {
exit(EXIT_FAILURE); // exit since memory allocation error
}
hptrB = (float *)malloc(size);
if (!hptrB) {
exit(EXIT_FAILURE); // exit since memory allocation error
}
hptrC = (float *)malloc(size);
if (!hptrC) {
exit(EXIT_FAILURE); // exit since memory allocation error
}
checkCudaErrors(cudaMalloc(&dptrA, size));
checkCudaErrors(cudaMalloc(&dptrB, size));
checkCudaErrors(cudaMalloc(&dptrC, size));
copyRequired = true;
break;
case USE_HOST_PAGELOCKED_AND_DEVICE_MEMORY:
case USE_HOST_PAGELOCKED_AND_DEVICE_MEMORY_ASYNC:
checkCudaErrors(cudaMallocHost(&hptrA, size));
checkCudaErrors(cudaMallocHost(&hptrB, size));
checkCudaErrors(cudaMallocHost(&hptrC, size));
checkCudaErrors(cudaMalloc(&dptrA, size));
checkCudaErrors(cudaMalloc(&dptrB, size));
checkCudaErrors(cudaMalloc(&dptrC, size));
copyRequired = true;
break;
case USE_ZERO_COPY:
checkCudaErrors(cudaMallocHost(&hptrA, size));
checkCudaErrors(cudaMallocHost(&hptrB, size));
checkCudaErrors(cudaMallocHost(&hptrC, size));
checkCudaErrors(cudaHostGetDevicePointer(&dptrA, hptrA, 0));
checkCudaErrors(cudaHostGetDevicePointer(&dptrB, hptrB, 0));
checkCudaErrors(cudaHostGetDevicePointer(&dptrC, hptrC, 0));
break;
case USE_MANAGED_MEMORY:
checkCudaErrors(cudaMallocManaged(&dptrA, size));
checkCudaErrors(cudaMallocManaged(&dptrB, size));
checkCudaErrors(cudaMallocManaged(&dptrC, size));
hptrA = dptrA;
hptrB = dptrB;
hptrC = dptrC;
break;
case USE_MANAGED_MEMORY_WITH_HINTS:
case USE_MANAGED_MEMORY_WITH_HINTS_ASYNC:
if (deviceProp.concurrentManagedAccess) {
checkCudaErrors(cudaMallocManaged(&dptrA, size));
checkCudaErrors(cudaMallocManaged(&dptrB, size));
checkCudaErrors(cudaMallocManaged(&dptrC, size));
checkCudaErrors(cudaMemPrefetchAsync(dptrA, size, cudaCpuDeviceId));
checkCudaErrors(cudaMemPrefetchAsync(dptrB, size, cudaCpuDeviceId));
checkCudaErrors(cudaMemPrefetchAsync(dptrC, size, cudaCpuDeviceId));
} else {
checkCudaErrors(cudaMallocManaged(&dptrA, size, cudaMemAttachHost));
checkCudaErrors(cudaMallocManaged(&dptrB, size, cudaMemAttachHost));
checkCudaErrors(cudaMallocManaged(&dptrC, size, cudaMemAttachHost));
}
hptrA = dptrA;
hptrB = dptrB;
hptrC = dptrC;
hintsRequired = true;
break;
default:
exit(EXIT_FAILURE); // exit with error
}
if (allocType == USE_HOST_PAGEABLE_AND_DEVICE_MEMORY_ASYNC ||
allocType == USE_HOST_PAGELOCKED_AND_DEVICE_MEMORY_ASYNC ||
allocType == USE_MANAGED_MEMORY_WITH_HINTS_ASYNC) {
isAsync = true;
}
someTransferOpRequired = copyRequired || hintsRequired;
// fill buffers with 0 to avoid any first access page-fault overheads.
memset(hptrA, 0, size);
memset(hptrB, 0, size);
memset(hptrC, 0, size);
for (i = 0; i < numLoops; i++) {
cpuAccessTimes[i] = 0.0;
gpuLaunchCallsTimes[i] = 0.0;
gpuTransferToCallsTimes[i] = 0.0;
gpuTransferFromCallsTimes[i] = 0.0;
sdkStartTimer(&cpuAccessTimer);
{
copyMatrix(hptrA, (i & 0x1 == 0) ? randValuesX : randValuesY, matrixDim);
copyMatrix(hptrB, (i & 0x1 == 0) ? randValuesY : randValuesX, matrixDim);
}
sdkStopTimer(&cpuAccessTimer);
cpuAccessTimes[i] += sdkGetAverageTimerValue(&cpuAccessTimer);
sdkResetTimer(&cpuAccessTimer);
if (isAsync && hintsRequired) {
*latch = 0;
// Prevent any work on stream from starting until all work is pushed
spinWhileLessThanOne<<<1, 1, 0, streamToRunOn>>>(latch);
}
if (someTransferOpRequired) {
sdkStartTimer(&gpuTransferCallsTimer);
if (copyRequired) {
if (isAsync) {
checkCudaErrors(cudaMemcpyAsync(
dptrA, hptrA, size, cudaMemcpyHostToDevice, streamToRunOn));
checkCudaErrors(cudaMemcpyAsync(
dptrB, hptrB, size, cudaMemcpyHostToDevice, streamToRunOn));
} else {
checkCudaErrors(
cudaMemcpy(dptrA, hptrA, size, cudaMemcpyHostToDevice));
checkCudaErrors(
cudaMemcpy(dptrB, hptrB, size, cudaMemcpyHostToDevice));
}
}
if (hintsRequired) {
if (deviceProp.concurrentManagedAccess) {
checkCudaErrors(
cudaMemPrefetchAsync(dptrA, size, device_id, streamToRunOn));
checkCudaErrors(
cudaMemPrefetchAsync(dptrB, size, device_id, streamToRunOn));
checkCudaErrors(
cudaMemPrefetchAsync(dptrC, size, device_id, streamToRunOn));
} else {
checkCudaErrors(cudaStreamAttachMemAsync(streamToRunOn, dptrA, 0,
cudaMemAttachGlobal));
checkCudaErrors(cudaStreamAttachMemAsync(streamToRunOn, dptrB, 0,
cudaMemAttachGlobal));
checkCudaErrors(cudaStreamAttachMemAsync(streamToRunOn, dptrC, 0,
cudaMemAttachGlobal));
}
if (!isAsync) {
checkCudaErrors(cudaStreamSynchronize(streamToRunOn));
}
}
sdkStopTimer(&gpuTransferCallsTimer);
gpuTransferToCallsTimes[i] +=
sdkGetAverageTimerValue(&gpuTransferCallsTimer);
sdkResetTimer(&gpuTransferCallsTimer);
}
sdkStartTimer(&gpuLaunchCallsTimer);
{
matrixMultiplyKernel<<<grid, threads, 0, streamToRunOn>>>(
dptrC, dptrA, dptrB, matrixDim);
if (!isAsync) {
checkCudaErrors(cudaStreamSynchronize(streamToRunOn));
}
}
sdkStopTimer(&gpuLaunchCallsTimer);
gpuLaunchCallsTimes[i] += sdkGetAverageTimerValue(&gpuLaunchCallsTimer);
sdkResetTimer(&gpuLaunchCallsTimer);
if (someTransferOpRequired) {
sdkStartTimer(&gpuTransferCallsTimer);
if (hintsRequired) {
if (deviceProp.concurrentManagedAccess) {
checkCudaErrors(cudaMemPrefetchAsync(dptrA, size, cudaCpuDeviceId));
checkCudaErrors(cudaMemPrefetchAsync(dptrB, size, cudaCpuDeviceId));
checkCudaErrors(cudaMemPrefetchAsync(dptrC, size, cudaCpuDeviceId));
} else {
checkCudaErrors(cudaStreamAttachMemAsync(streamToRunOn, dptrA, 0,
cudaMemAttachHost));
checkCudaErrors(cudaStreamAttachMemAsync(streamToRunOn, dptrB, 0,
cudaMemAttachHost));
checkCudaErrors(cudaStreamAttachMemAsync(streamToRunOn, dptrC, 0,
cudaMemAttachHost));
}
if (!isAsync) {
checkCudaErrors(cudaStreamSynchronize(streamToRunOn));
}
}
if (copyRequired) {
if (isAsync) {
checkCudaErrors(cudaMemcpyAsync(
hptrC, dptrC, size, cudaMemcpyDeviceToHost, streamToRunOn));
} else {
checkCudaErrors(
cudaMemcpy(hptrC, dptrC, size, cudaMemcpyDeviceToHost));
}
}
sdkStopTimer(&gpuTransferCallsTimer);
gpuTransferFromCallsTimes[i] +=
sdkGetAverageTimerValue(&gpuTransferCallsTimer);
sdkResetTimer(&gpuTransferCallsTimer);
}
gpuLaunchAndTransferCallsTimes[i] = gpuLaunchCallsTimes[i] +
gpuTransferToCallsTimes[i] +
gpuTransferFromCallsTimes[i];
gpuLaunchTransferSyncTimes[i] = gpuLaunchAndTransferCallsTimes[i];
if (isAsync) {
sdkStartTimer(&gpuSyncTimer);
{
if (hintsRequired) {
*latch = 1;
}
checkCudaErrors(cudaStreamSynchronize(streamToRunOn));
}
sdkStopTimer(&gpuSyncTimer);
gpuLaunchTransferSyncTimes[i] += sdkGetAverageTimerValue(&gpuSyncTimer);
sdkResetTimer(&gpuSyncTimer);
}
sdkStartTimer(&cpuAccessTimer);
{
verifyMatrixData(
(i & 0x1 == 0) ? randValuesVerifyXmulY : randValuesVerifyYmulX, hptrC,
matrixDim);
}
sdkStopTimer(&cpuAccessTimer);
cpuAccessTimes[i] += sdkGetAverageTimerValue(&cpuAccessTimer);
sdkResetTimer(&cpuAccessTimer);
overallTimes[i] = cpuAccessTimes[i] + gpuLaunchTransferSyncTimes[i];
}
switch (allocType) {
case USE_HOST_PAGEABLE_AND_DEVICE_MEMORY:
case USE_HOST_PAGEABLE_AND_DEVICE_MEMORY_ASYNC:
free(hptrA);
free(hptrB);
free(hptrC);
checkCudaErrors(cudaFree(dptrA));
checkCudaErrors(cudaFree(dptrB));
checkCudaErrors(cudaFree(dptrC));
break;
case USE_HOST_PAGELOCKED_AND_DEVICE_MEMORY:
case USE_HOST_PAGELOCKED_AND_DEVICE_MEMORY_ASYNC:
checkCudaErrors(cudaFreeHost(hptrA));
checkCudaErrors(cudaFreeHost(hptrB));
checkCudaErrors(cudaFreeHost(hptrC));
checkCudaErrors(cudaFree(dptrA));
checkCudaErrors(cudaFree(dptrB));
checkCudaErrors(cudaFree(dptrC));
break;
case USE_ZERO_COPY:
checkCudaErrors(cudaFreeHost(hptrA));
checkCudaErrors(cudaFreeHost(hptrB));
checkCudaErrors(cudaFreeHost(hptrC));
break;
case USE_MANAGED_MEMORY:
case USE_MANAGED_MEMORY_WITH_HINTS:
case USE_MANAGED_MEMORY_WITH_HINTS_ASYNC:
checkCudaErrors(cudaFree(dptrA));
checkCudaErrors(cudaFree(dptrB));
checkCudaErrors(cudaFree(dptrC));
break;
default:
exit(EXIT_FAILURE); // exit due to error
}
checkCudaErrors(cudaStreamDestroy(streamToRunOn));
checkCudaErrors(cudaFreeHost(latch));
free(randValuesX);
free(randValuesY);
free(randValuesVerifyXmulY);
free(randValuesVerifyYmulX);
sdkDeleteTimer(&gpuLaunchCallsTimer);
sdkDeleteTimer(&gpuTransferCallsTimer);
sdkDeleteTimer(&gpuSyncTimer);
sdkDeleteTimer(&cpuAccessTimer);
}
void matrixMultiplyPerfRunner(bool reportAsBandwidth,
bool print_launch_transfer_results,
bool print_std_deviation, int device_id) {
int i;
unsigned int minMatrixDim = 32;
unsigned int multiplierDim = 2;
unsigned int matrixDim;
unsigned int minSize = minMatrixDim * minMatrixDim * sizeof(float);
unsigned int maxSize =
(maxSampleSizeInMb * ONE_MB) /
4; // 3 buffers are used, but dividing by 4 (power of 2)
unsigned int multiplier = multiplierDim * multiplierDim;
unsigned int numSizesToTest;
struct testResults *results;
struct resultsData *gpuLaunchCallsTimes;
struct resultsData *gpuTransferToCallsTimes;
struct resultsData *gpuTransferFromCallsTimes;
struct resultsData *gpuLaunchAndTransferCallsTimes;
struct resultsData *gpuLaunchTransferSyncTimes;
struct resultsData *cpuAccessTimes;
struct resultsData *overallTimes;
unsigned long *sizesToTest;
unsigned int j;
numSizesToTest = findNumSizesToTest(minSize, maxSize, multiplier);
createAndInitTestResults(&results, "matrixMultiplyPerf", numKernelRuns,
numSizesToTest);
sizesToTest = getPtrSizesToTest(results);
createResultDataAndAddToTestResults(&gpuLaunchCallsTimes, results,
"GPU Kernel Launch Call Time", false,
reportAsBandwidth);
createResultDataAndAddToTestResults(&gpuTransferToCallsTimes, results,
"CPU to GPU Transfer Calls Time", false,
reportAsBandwidth);
createResultDataAndAddToTestResults(&gpuTransferFromCallsTimes, results,
"GPU to CPU Transfer Calls Time", false,
reportAsBandwidth);
createResultDataAndAddToTestResults(&gpuLaunchAndTransferCallsTimes, results,
"GPU Launch and Transfer Calls Time",
false, reportAsBandwidth);
createResultDataAndAddToTestResults(&gpuLaunchTransferSyncTimes, results,
"GPU Launch Transfer and Sync Time",
false, reportAsBandwidth);
createResultDataAndAddToTestResults(
&cpuAccessTimes, results, "CPU Access Time", false, reportAsBandwidth);
createResultDataAndAddToTestResults(&overallTimes, results, "Overall Time",
false, reportAsBandwidth);
printf("Running ");
for (matrixDim = minMatrixDim, j = 0;
matrixDim * matrixDim <= maxSize / sizeof(float);
matrixDim *= multiplierDim, ++j) {
sizesToTest[j] = matrixDim * matrixDim * sizeof(float);
for (i = MEMALLOC_TYPE_START; i <= MEMALLOC_TYPE_END; i++) {
printf(".");
fflush(stdout);
runMatrixMultiplyKernel(
matrixDim, i, numKernelRuns,
getPtrRunTimesInMs(gpuLaunchCallsTimes, i, j),
getPtrRunTimesInMs(gpuTransferToCallsTimes, i, j),
getPtrRunTimesInMs(gpuTransferFromCallsTimes, i, j),
getPtrRunTimesInMs(gpuLaunchAndTransferCallsTimes, i, j),
getPtrRunTimesInMs(gpuLaunchTransferSyncTimes, i, j),
getPtrRunTimesInMs(cpuAccessTimes, i, j),
getPtrRunTimesInMs(overallTimes, i, j), device_id);
}
}
printf("\n");
printResults(results, print_launch_transfer_results, print_std_deviation);
freeTestResultsAndAllResultsData(results);
}
static void usage() {
printf(
"./cudaMemoryTypesPerf [-device=<device_id>] [-reportAsBandwidth] "
"[-print-launch-transfer-results] [-print-std-deviation] [-verbose]\n");
printf("Options:\n");
printf(
"-reportAsBandwidth: By default time taken is printed, this "
"option allows to instead print bandwidth.\n");
printf(
"-print-launch-transfer-results: By default overall results are printed, "
"this option allows to print data transfers and kernel time as well.\n");
printf(
"-print-std-deviation: Prints std deviation of the results.\n");
printf(
"-kernel-iterations=<num>: Number of times the kernel tests should "
"be run[default is 100 iterations].\n");
printf(
"-device=<device_id>: Allows to pass GPU Device ID on which "
"the tests will be run.\n");
printf("-verbose: Prints highly verbose output.\n");
}
int main(int argc, char **argv) {
bool reportAsBandwidth = false;
bool print_launch_transfer_results = false;
bool print_std_deviation = false;
if (checkCmdLineFlag(argc, (const char **)argv, "help") ||
checkCmdLineFlag(argc, (const char **)argv, "h")) {
usage();
printf("&&&& %s WAIVED\n", argv[0]);
exit(EXIT_WAIVED);
}
if (checkCmdLineFlag(argc, (const char **)argv, "reportAsBandwidth")) {
reportAsBandwidth = true;
}
if (checkCmdLineFlag(argc, (const char **)argv,
"print-launch-transfer-results")) {
print_launch_transfer_results = true;
}
if (checkCmdLineFlag(argc, (const char **)argv, "print-std-deviation")) {
print_std_deviation = true;
}
if (checkCmdLineFlag(argc, (const char **)argv, "kernel-iterations")) {
numKernelRuns =
getCmdLineArgumentInt(argc, (const char **)argv, "kernel-iterations");
}
if (checkCmdLineFlag(argc, (const char **)argv, "verbose")) {
verboseResults = 1;
}
int device_id = findCudaDevice(argc, (const char **)argv);
matrixMultiplyPerfRunner(reportAsBandwidth, print_launch_transfer_results,
print_std_deviation, device_id);
printf(
"\nNOTE: The CUDA Samples are not meant for performance measurements. "
"Results may vary when GPU Boost is enabled.\n");
exit(EXIT_SUCCESS);
}