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

Merge branch 'master' into amd-master-next

Browse Source 
Change-Id: Ibb7775e7acd263d2ece40a241517bbd15976fdd4


[ROCm/hip commit: 2cf3257795]
This commit is contained in:
Laurent Morichetti
2020-02-12 00:14:07 -08:00
parent 652e82f726 2868c785d8
commit 471b697a05
38 changed files with 602 additions and 1010 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip/tests/src/gcc/LaunchKernel.c
+41 -12
View File
@@ -36,7 +36,7 @@ bool LaunchKernelArg()
dim3 blocks = {1,1,1};
dim3 threads = {1,1,1};
HIPCHECK(hipLaunchKernel(kernel, blocks, threads,NULL, 0, 0));
HIPCHECK(hipLaunchKernel(kernel, blocks, threads, NULL, 0, 0));
return true;
}
@@ -50,9 +50,9 @@ bool LaunchKernelArg1()
// Allocate Device memory
HIPCHECK(hipMalloc((void**)&A_d, sizeof(int)));
void* Args[]={A_d};
HIPCHECK(hipLaunchKernel(kernel1, blocks, threads, Args,0,0));
void* Args[]={&A_d};
HIPCHECK(hipLaunchKernel(kernel1, blocks, threads, Args, 0, 0));
// Get the result back to host memory
HIPCHECK(hipMemcpy(&A, A_d, sizeof(int), hipMemcpyDeviceToHost));
@@ -60,7 +60,7 @@ bool LaunchKernelArg1()
HIPCHECK(hipFree(A_d));
if(A != 333)
return false;
return false;
return true;
}
@@ -81,9 +81,9 @@ bool LaunchKernelArg2()
HIPCHECK(hipMalloc((void**)&B_d, sizeof(int)));
// Copy data from host memory to device memory
HIPCHECK(hipMemcpy(B_d,&B, sizeof(int), hipMemcpyHostToDevice));
HIPCHECK(hipMemcpy(B_d, &B, sizeof(int), hipMemcpyHostToDevice));
void* Args[]={A_d,B_d};
void* Args[]={&A_d, &B_d};
HIPCHECK(hipLaunchKernel(kernel2, blocks, threads, Args,0,0));
// Get the result back to host memory
@@ -118,11 +118,11 @@ bool LaunchKernelArg3()
HIPCHECK(hipMalloc((void**)&C_d, sizeof(int)));
// Copy data from host memory to device memory
HIPCHECK(hipMemcpy(A_d,&A, sizeof(int), hipMemcpyHostToDevice));
HIPCHECK(hipMemcpy(A_d, &A, sizeof(int), hipMemcpyHostToDevice));
HIPCHECK(hipMemcpy(B_d,&B, sizeof(int), hipMemcpyHostToDevice));
HIPCHECK(hipMemcpy(B_d, &B, sizeof(int), hipMemcpyHostToDevice));
void* Args[]={A_d,B_d,C_d};
void* Args[]={&A_d, &B_d, &C_d};
HIPCHECK(hipLaunchKernel(kernel3, blocks, threads, Args,0,0));
// Get the result back to host memory
@@ -138,14 +138,43 @@ bool LaunchKernelArg3()
return true;
}
bool LaunchKernelArg4()
{
int A = 0;
int *A_d = NULL;
dim3 blocks = {1,1,1};
dim3 threads = {1,1,1};
// Allocate Device memory
HIPCHECK(hipMalloc((void**)&A_d, sizeof(int)));
char c = 1;
short s = 10;
int i = 100;
struct things t = {2,20,200};
void* Args[]={&A_d, &c, &s, &i, &t};
HIPCHECK(hipLaunchKernel(kernel4, blocks, threads, Args, 0, 0));
// Get the result back to host memory
HIPCHECK(hipMemcpy(&A, A_d, sizeof(int), hipMemcpyDeviceToHost));
HIPCHECK(hipFree(A_d));
if (A != (c + s + i + t.c + t.s + t.i))
return false;
return true;
}
int main()
{
if( LaunchKernelArg() &&
LaunchKernelArg1() &&
LaunchKernelArg2() &&
LaunchKernelArg3())
LaunchKernelArg3() &&
LaunchKernelArg4())
{
printf("PASSED!\n");
}
projects/hip/tests/src/gcc/LaunchKernel.h
+16
View File
@@ -17,7 +17,23 @@
* THE SOFTWARE.
*/
#ifdef __cplusplus
extern "C"
{
#endif
extern __global__ void kernel();
extern __global__ void kernel1(int*);
extern __global__ void kernel2(int*,int*);
extern __global__ void kernel3(int*,int*,int*);
struct things {
char c;
short s;
int i;
};
extern __global__ void kernel4(int*, char, short, int, struct things);
#ifdef __cplusplus
}
#endif
projects/hip/tests/src/gcc/gpu.cpp
+6
View File
@@ -19,6 +19,7 @@
#include<hip/hip_runtime.h>
#include "LaunchKernel.h"
extern "C"
{
@@ -43,4 +44,9 @@ __global__ void kernel3(int *a, int*b, int* c)
*c = *a+*b;
}
__global__ void kernel4(int *a, char c, short s, int i, struct things t)
{
*a = c + s + i + t.c + t.s + t.i;
}
}//extern "C"
projects/hip/tests/src/runtimeApi/module/hipLaunchCoopMultiKernel.cpp
+1 -1
View File
@@ -156,7 +156,7 @@ int main() {
for (int i = 0; i < nGpu; i++) {
HIPCHECK(hipSetDevice(i));
dimBlock.x = workgroups[set];
HIPCHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlocks,
HIPCHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(reinterpret_cast<uint32_t*>(&numBlocks),
(hipFunction_t)test_gws, dimBlock.x * dimBlock.y * dimBlock.z, dimBlock.x * sizeof(long)));
std::cout << "GPU(" << i << ") Block size: " << dimBlock.x << " Num blocks per CU: " << numBlocks << "\n";
projects/hip/tests/src/runtimeApi/module/hipLaunchCooperativeKernel.cpp
+1 -1
View File
@@ -116,7 +116,7 @@ int main() {
dimBlock.x = workgroups[i];
// Calculate the device occupancy to know how many blocks can be run concurrently
hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlocks,
hipOccupancyMaxActiveBlocksPerMultiprocessor(reinterpret_cast<uint32_t*>(&numBlocks),
test_gws, dimBlock.x * dimBlock.y * dimBlock.z, dimBlock.x * sizeof(long));
dimGrid.x = deviceProp.multiProcessorCount * std::min(numBlocks, 32);
projects/hip/tests/src/runtimeApi/occupancy/hipOccupancyMaxActiveBlocksPerMultiprocessor.cpp → projects/hip/tests/src/runtimeApi/module/hipOccupancyMaxActiveBlocksPerMultiprocessor.cpp
+18 -3
View File
@@ -30,6 +30,10 @@ THE SOFTWARE.
#include "hip/hip_runtime.h"
#include "test_common.h"
#define fileName "vcpy_kernel.code"
#define kernel_name "hello_world"
__global__ void f1(float *a) { *a = 1.0; }
template <typename T>
@@ -45,10 +49,11 @@ int main(int argc, char* argv[]) {
hipOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, f1, 0, 0);
assert(gridSize != 0 && blockSize != 0);
int numBlock = 0;
hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlock, f1, (int)blockSize, 0);
uint32_t numBlock = 0;
hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlock, f1, blockSize, 0);
assert(numBlock != 0);
// test case for using kernel function pointer with template
gridSize = 0;
blockSize = 0;
@@ -56,7 +61,17 @@ int main(int argc, char* argv[]) {
assert(gridSize != 0 && blockSize != 0);
numBlock = 0;
hipOccupancyMaxActiveBlocksPerMultiprocessor<void(*)(int *)>(&numBlock, f2, (int)blockSize, 0);
hipOccupancyMaxActiveBlocksPerMultiprocessor<void(*)(int *)>(&numBlock, f2, blockSize, 0);
assert(numBlock != 0);
// test case for using kernel with hipFunction_t type
numBlock = 0;
hipModule_t Module;
hipFunction_t Function;
HIPCHECK(hipModuleLoad(&Module, fileName));
HIPCHECK(hipModuleGetFunction(&Function, Module, kernel_name));
HIPCHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlock, Function, blockSize, 0));
assert(numBlock != 0);
passed();
projects/hip/tests/src/runtimeApi/occupancy/hipOccupancyMaxPotentialBlockSize.cpp → projects/hip/tests/src/runtimeApi/module/hipOccupancyMaxPotentialBlockSize.cpp
+1
View File
@@ -33,6 +33,7 @@ THE SOFTWARE.
#define fileName "vcpy_kernel.code"
#define kernel_name "hello_world"
__global__ void f1(float *a) { *a = 1.0; }
template <typename T>
projects/hip/tests/src/texture/hipNormalizedFloatValueTex.cpp
+19 -3
View File
@@ -27,8 +27,24 @@ THE SOFTWARE.
*/
#include "test_common.h"
#define SIZE 10
static float getNormalizedValue(const float value,
const enum hipArray_Format texFormat) {
switch (texFormat) {
case HIP_AD_FORMAT_SIGNED_INT8:
return (value / SCHAR_MAX);
case HIP_AD_FORMAT_UNSIGNED_INT8:
return (value / UCHAR_MAX);
case HIP_AD_FORMAT_SIGNED_INT16:
return (value / SHRT_MAX);
case HIP_AD_FORMAT_UNSIGNED_INT16:
return (value / USHRT_MAX);
default:
return value;
}
}
texture<float, hipTextureType1D, hipReadModeElementType> textureNormalizedVal_1D;
__global__ void normalizedValTextureTest(unsigned int numElements, float* pDst)
@@ -47,7 +63,6 @@ bool textureTest(enum hipArray_Format texFormat)
T *dData = NULL;
HIPCHECK(hipMalloc((void **) &dData, sizeof(T)*SIZE));
HIPCHECK(hipMemcpyHtoD((hipDeviceptr_t)dData, hData, sizeof(T)*SIZE));
textureReference* texRef = &textureNormalizedVal_1D;
HIPCHECK(hipTexRefSetAddressMode(texRef, 0, hipAddressModeClamp));
HIPCHECK(hipTexRefSetAddressMode(texRef, 1, hipAddressModeClamp));
@@ -73,7 +88,8 @@ bool textureTest(enum hipArray_Format texFormat)
for(int i = 0; i < SIZE; i++)
{
if((float)hData[i]/texFormatToSize[texFormat] != hOutputData[i])
float expected = getNormalizedValue(float(hData[i]), texFormat);
if(expected != hOutputData[i])
{
printf("mismatch at index:%d for texType:%d output:%f\n",i,texFormat,hOutputData[i]);
testResult = false;