Initial commit for counter correctness tests.
Change-Id: I74caa4ab89bd765e59686cfbaaf1ce44ab10fe19
- adds counter correctness tests for GRBM, SQWAVES and SQ INSTRUCTIONs
- pandas dependency
- made it compatible with test infra.
Change-Id: I74caa4ab89bd765e59686cfbaaf1ce44ab10fe19
[ROCm/rocprofiler commit: 3f5ddfa0ab]
This commit is contained in:
Giovanni LB
کامیت شده توسط
Gopesh Bhardwaj
Gopesh Bhardwaj
والد
45d34bb867
کامیت
976c0e6717
@@ -0,0 +1,281 @@
|
||||
/**********************************************************************
|
||||
Copyright �2015 Advanced Micro Devices, Inc. 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.
|
||||
|
||||
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 THE COPYRIGHT HOLDER 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 <math.h>
|
||||
#include "hip/hip_runtime.h"
|
||||
|
||||
#include <assert.h>
|
||||
#include <stdio.h>
|
||||
#include <algorithm>
|
||||
#include <stdlib.h>
|
||||
#include <iostream>
|
||||
#include <unistd.h>
|
||||
#include <vector>
|
||||
|
||||
#include "histogram.hpp"
|
||||
|
||||
#define LINEAR_MEM_ACCESS
|
||||
|
||||
#define BIN_SIZE 256
|
||||
#define SDK_SUCCESS 0
|
||||
#define SDK_FAILURE 1
|
||||
#define CHECK_ALLOCATION(x, msg) if(!(x)) { std::cout << __FILE__ << ' ' << __LINE__ << ' ' << msg << std::endl; }
|
||||
|
||||
|
||||
/**
|
||||
* @brief Calculates block-histogram bin whose bin size is 256
|
||||
* @param data input data pointer
|
||||
* @param sharedArray shared array for thread-histogram bins
|
||||
* @param binResult block-histogram array
|
||||
*/
|
||||
|
||||
__global__
|
||||
void histogram256(
|
||||
unsigned int* data,
|
||||
unsigned int* binResult)
|
||||
{
|
||||
HIP_DYNAMIC_SHARED(unsigned char, sharedArray);
|
||||
size_t localId = hipThreadIdx_x;
|
||||
size_t globalId = hipThreadIdx_x + hipBlockIdx_x*hipBlockDim_x;
|
||||
size_t groupId = hipBlockIdx_x;
|
||||
size_t groupSize = hipBlockDim_x;
|
||||
int offSet1 = localId & 31;
|
||||
int offSet2 = 4 * offSet1; //which element to access in one bank.
|
||||
int offSet3 = localId >> 5; //bank number
|
||||
/* initialize shared array to zero */
|
||||
uchar4 * input = (uchar4*)sharedArray;
|
||||
for(int i = 0; i < 64; ++i)
|
||||
input[groupSize * i + localId] = make_uchar4(0,0,0,0);
|
||||
|
||||
__syncthreads();
|
||||
|
||||
|
||||
/* calculate thread-histograms */
|
||||
//128 accumulations per thread
|
||||
for(int i = 0; i < 128; i++)
|
||||
{
|
||||
#ifdef LINEAR_MEM_ACCESS
|
||||
uint value = data[groupId * (groupSize * (BIN_SIZE/2)) + i * groupSize + localId];
|
||||
#else
|
||||
uint value = data[globalId + i*4096];
|
||||
|
||||
#endif // LINEAR_MEM_ACCESS
|
||||
sharedArray[value * 128 + offSet2 + offSet3]++;
|
||||
}
|
||||
__syncthreads();
|
||||
|
||||
/* merge all thread-histograms into block-histogram */
|
||||
|
||||
uint4 binCount;
|
||||
uint result;
|
||||
uchar4 binVal; //Introduced uint4 for summation to avoid overflows
|
||||
uint4 binValAsUint;
|
||||
for(int i = 0; i < BIN_SIZE / groupSize; ++i)
|
||||
{
|
||||
int passNumber = BIN_SIZE / 2 * 32 * i + localId * 32 ;
|
||||
binCount = make_uint4(0,0,0,0);
|
||||
result= 0;
|
||||
for(int j = 0; j < 32; ++j)
|
||||
{
|
||||
int bankNum = (j + offSet1) & 31; // this is bank number
|
||||
binVal = input[passNumber +bankNum];
|
||||
|
||||
binValAsUint.x = (unsigned int)binVal.x;
|
||||
binValAsUint.y = (unsigned int)binVal.y;
|
||||
binValAsUint.z = (unsigned int)binVal.z;
|
||||
binValAsUint.w = (unsigned int)binVal.w;
|
||||
|
||||
binCount.x += binValAsUint.x;
|
||||
binCount.y += binValAsUint.y;
|
||||
binCount.z += binValAsUint.z;
|
||||
binCount.w += binValAsUint.w;
|
||||
|
||||
}
|
||||
result = binCount.x + binCount.y + binCount.z + binCount.w;
|
||||
binResult[groupId * BIN_SIZE + groupSize * i + localId ] = result;
|
||||
}
|
||||
}
|
||||
|
||||
int
|
||||
Histogram::calculateHostBin()
|
||||
{
|
||||
for(int i = 0; i < height; ++i)
|
||||
{
|
||||
for(int j = 0; j < width; ++j)
|
||||
{
|
||||
hostBin[data[i * width + j]]++;
|
||||
}
|
||||
}
|
||||
|
||||
return SDK_SUCCESS;
|
||||
}
|
||||
|
||||
|
||||
int
|
||||
Histogram::setupHistogram()
|
||||
{
|
||||
int i = 0;
|
||||
|
||||
data = (unsigned int *)malloc(sizeof(unsigned int) * width * height);
|
||||
|
||||
for(i = 0; i < width * height; i++)
|
||||
{
|
||||
data[i] = rand() % (unsigned int)(binSize);
|
||||
}
|
||||
|
||||
hostBin = (unsigned int*)malloc(binSize * sizeof(unsigned int));
|
||||
CHECK_ALLOCATION(hostBin, "Failed to allocate host memory. (hostBin)");
|
||||
|
||||
memset(hostBin, 0, binSize * sizeof(unsigned int));
|
||||
|
||||
deviceBin = (unsigned int*)malloc(binSize * sizeof(unsigned int));
|
||||
CHECK_ALLOCATION(deviceBin, "Failed to allocate host memory. (deviceBin)");
|
||||
midDeviceBin = (unsigned int*)malloc(sizeof(unsigned int) * binSize * subHistgCnt);
|
||||
|
||||
memset(deviceBin, 0, binSize * sizeof(unsigned int));
|
||||
return SDK_SUCCESS;
|
||||
}
|
||||
|
||||
int
|
||||
Histogram::setupHIP(void)
|
||||
{
|
||||
hipDeviceProp_t devProp;
|
||||
hipGetDeviceProperties(&devProp, 0);
|
||||
cout << " System minor " << devProp.minor << endl;
|
||||
cout << " System major " << devProp.major << endl;
|
||||
cout << " agent prop name " << devProp.name << endl;
|
||||
|
||||
return SDK_SUCCESS;
|
||||
}
|
||||
|
||||
|
||||
int
|
||||
Histogram::runKernels(void)
|
||||
{
|
||||
groupSize = 128;
|
||||
globalThreads = (width * height) / (GROUP_ITERATIONS);
|
||||
|
||||
localThreads = groupSize;
|
||||
|
||||
|
||||
hipHostMalloc((void**)&dataBuf,sizeof(unsigned int) * width * height, hipHostMallocDefault);
|
||||
unsigned int *din;
|
||||
hipHostGetDevicePointer((void**)&din, dataBuf,0);
|
||||
hipMemcpy(din, data,sizeof(unsigned int) * width * height, hipMemcpyHostToDevice);
|
||||
|
||||
subHistgCnt = (width * height) / (groupSize * groupIterations);
|
||||
|
||||
hipHostMalloc((void**)&midDeviceBinBuf,sizeof(unsigned int) * binSize * subHistgCnt, hipHostMallocDefault);
|
||||
|
||||
hipLaunchKernelGGL(histogram256,
|
||||
dim3(globalThreads/localThreads),
|
||||
dim3(localThreads),
|
||||
groupSize * binSize * sizeof(unsigned char), 0,
|
||||
dataBuf ,midDeviceBinBuf);
|
||||
|
||||
hipDeviceSynchronize();
|
||||
|
||||
hipMemcpy(midDeviceBin, midDeviceBinBuf,sizeof(unsigned int) * binSize * subHistgCnt, hipMemcpyDeviceToHost);
|
||||
//printArray<unsigned int>("midDeviceBin", midDeviceBin, sizeof(unsigned int) * binSize * subHistgCnt, 1);
|
||||
// Clear deviceBin array
|
||||
memset(deviceBin, 0, binSize * sizeof(unsigned int));
|
||||
|
||||
// Calculate final histogram bin
|
||||
for(int i = 0; i < subHistgCnt; ++i)
|
||||
{
|
||||
for(int j = 0; j < binSize; ++j)
|
||||
{
|
||||
deviceBin[j] += midDeviceBin[i * binSize + j];
|
||||
}
|
||||
}
|
||||
|
||||
return SDK_SUCCESS;
|
||||
}
|
||||
|
||||
int
|
||||
Histogram::setup()
|
||||
{
|
||||
if(iterations < 1)
|
||||
{
|
||||
std::cout<<"Error, iterations cannot be 0 or negative. Exiting..\n";
|
||||
exit(0);
|
||||
}
|
||||
int status = 0;
|
||||
|
||||
/* width must be multiples of binSize and
|
||||
* height must be multiples of groupSize
|
||||
*/
|
||||
width = (width / binSize ? width / binSize: 1) * binSize;
|
||||
height = (height / groupSize ? height / groupSize: 1) * groupSize;
|
||||
|
||||
status = setupHIP();
|
||||
if(status != SDK_SUCCESS)
|
||||
return status;
|
||||
|
||||
status = setupHistogram();
|
||||
if(status != SDK_SUCCESS)
|
||||
return status;
|
||||
|
||||
return SDK_SUCCESS;
|
||||
}
|
||||
|
||||
|
||||
int Histogram::run()
|
||||
{
|
||||
for(int i = 0; i < 2 && iterations != 1; i++)
|
||||
if(runKernels() != SDK_SUCCESS)
|
||||
return SDK_FAILURE;
|
||||
|
||||
for(int i = 0; i < iterations; i++)
|
||||
if(runKernels() != SDK_SUCCESS)
|
||||
return SDK_FAILURE;
|
||||
|
||||
return SDK_SUCCESS;
|
||||
}
|
||||
|
||||
int Histogram::cleanup()
|
||||
{
|
||||
hipFree(dataBuf);
|
||||
hipFree(midDeviceBinBuf);
|
||||
|
||||
free(hostBin);
|
||||
free(deviceBin);
|
||||
|
||||
return SDK_SUCCESS;
|
||||
}
|
||||
|
||||
int
|
||||
main(int argc, char * argv[])
|
||||
{
|
||||
int status = 0;
|
||||
// Create MonteCalroAsian object
|
||||
Histogram hipHistogram;
|
||||
|
||||
// Setup
|
||||
status = hipHistogram.setup();
|
||||
if(status != SDK_SUCCESS)
|
||||
return status;
|
||||
|
||||
// Run
|
||||
if(hipHistogram.run() != SDK_SUCCESS)
|
||||
return SDK_FAILURE;
|
||||
|
||||
// Cleanup resources created
|
||||
if(hipHistogram.cleanup() != SDK_SUCCESS)
|
||||
return SDK_FAILURE;
|
||||
|
||||
return SDK_SUCCESS;
|
||||
}
|
||||
مرجع در شماره جدید
Block a user