AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Add more apps to 2_Cookbook

Browse Source 
Change-Id: Iafe462df9726a32f450bd240a2de3eaa73a10057
This commit is contained in:
Sandeep Kumar
2016-10-14 18:00:26 +05:30
committed by Maneesh Gupta
parent 5f2bf5219a
commit d025ed980b
21 changed files with 686 additions and 94 deletions
Download Patch File Download Diff File Expand all files Collapse all files
hipamd/samples/2_Cookbook/4_shfl/shfl.cpp
+10 -13
View File
@@ -27,9 +27,8 @@ THE SOFTWARE.
#define WIDTH 4
#define HEIGHT 4
#define NUM (WIDTH*HEIGHT)
#define NUM (WIDTH*WIDTH)
#define THREADS_PER_BLOCK_X 4
#define THREADS_PER_BLOCK_Y 4
@@ -40,17 +39,16 @@ THE SOFTWARE.
__global__ void matrixTranspose(hipLaunchParm lp,
float *out,
float *in,
const int width,
const int height)
const int width)
{
int x = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
//int y = hipBlockDim_y * hipBlockIdx_y + hipThreadIdx_y;
float val = in[x];
for(int i=0;i<width;i++)
{
for(int j=0;j<width;j++)
out[i*height + j] = __shfl(val,j*width + i);
out[i*width + j] = __shfl(val,j*width + i);
}
}
@@ -58,14 +56,13 @@ __global__ void matrixTranspose(hipLaunchParm lp,
void matrixTransposeCPUReference(
float * output,
float * input,
const unsigned int width,
const unsigned int height)
const unsigned int width)
{
for(unsigned int j=0; j < height; j++)
for(unsigned int j=0; j < width; j++)
{
for(unsigned int i=0; i < width; i++)
{
output[i*height + j] = input[j*width + i];
output[i*width + j] = input[j*width + i];
}
}
}
@@ -108,19 +105,19 @@ int main() {
dim3(1),
dim3(THREADS_PER_BLOCK_X * THREADS_PER_BLOCK_Y),
0, 0,
gpuTransposeMatrix , gpuMatrix, WIDTH ,HEIGHT);
gpuTransposeMatrix , gpuMatrix, WIDTH);
// Memory transfer from device to host
hipMemcpy(TransposeMatrix, gpuTransposeMatrix, NUM*sizeof(float), hipMemcpyDeviceToHost);
// CPU MatrixTranspose computation
matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH, HEIGHT);
matrixTransposeCPUReference(cpuTransposeMatrix, Matrix, WIDTH);
// verify the results
errors = 0;
double eps = 1.0E-6;
for (i = 0; i < NUM; i++) {
if (std::abs(TransposeMatrix[i] - cpuTransposeMatrix[i]) > 0 ) {
if (std::abs(TransposeMatrix[i] - cpuTransposeMatrix[i]) > eps ) {
printf("%d cpu: %f gpu %f\n",i,cpuTransposeMatrix[i],TransposeMatrix[i]);
errors++;
}