From af57c166c063cf253a003c50739c826917a51a27 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Wed, 28 Mar 2018 15:36:57 +0530 Subject: [PATCH 1/4] Remove textureObj kernel argument for HIP/HCC path [ROCm/hip commit: 017a6fb0590d2d3e94b9e649e269fa01d6a9d99d] --- projects/hip/tests/src/texture/hipTextureRef2D.cpp | 12 ------------ 1 file changed, 12 deletions(-) diff --git a/projects/hip/tests/src/texture/hipTextureRef2D.cpp b/projects/hip/tests/src/texture/hipTextureRef2D.cpp index 6f879f80c8..c4c0b9e2fe 100644 --- a/projects/hip/tests/src/texture/hipTextureRef2D.cpp +++ b/projects/hip/tests/src/texture/hipTextureRef2D.cpp @@ -14,17 +14,10 @@ texture tex; bool testResult = true; __global__ void tex2DKernel(float* outputData, -#ifdef __HIP_PLATFORM_HCC__ - hipTextureObject_t textureObject, -#endif int width, int height) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; -#ifdef __HIP_PLATFORM_HCC__ - outputData[y * width + x] = tex2D(tex, textureObject, x, y); -#else outputData[y * width + x] = tex2D(tex, x, y); -#endif } void runTest(int argc, char** argv); @@ -73,12 +66,7 @@ void runTest(int argc, char** argv) { dim3 dimBlock(16, 16, 1); dim3 dimGrid(width / dimBlock.x, height / dimBlock.y, 1); -#ifdef __HIP_PLATFORM_HCC__ - hipLaunchKernelGGL(tex2DKernel, dim3(dimGrid), dim3(dimBlock), 0, 0, dData, tex.textureObject, - width, height); -#else hipLaunchKernelGGL(tex2DKernel, dim3(dimGrid), dim3(dimBlock), 0, 0, dData, width, height); -#endif hipDeviceSynchronize(); float* hOutputData = (float*)malloc(size); From d0e8f219a93378de20055a55729f07067b568970 Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Mon, 2 Apr 2018 11:58:32 -0400 Subject: [PATCH 2/4] Update HIP language spec to support both <<< >>> kernel launching mechanism and hipLaunchKernel [ROCm/hip commit: cf78d85638b22a586317dfee4514d4590e7b2eec] --- projects/hip/docs/markdown/hip_kernel_language.md | 15 +++++++++------ 1 file changed, 9 insertions(+), 6 deletions(-) diff --git a/projects/hip/docs/markdown/hip_kernel_language.md b/projects/hip/docs/markdown/hip_kernel_language.md index 094d7531e8..7c209acadf 100644 --- a/projects/hip/docs/markdown/hip_kernel_language.md +++ b/projects/hip/docs/markdown/hip_kernel_language.md @@ -98,8 +98,8 @@ HIP parses the `__noinline__` and `__forceinline__` keywords and converts them t ## Calling `__global__` Functions -`__global__` functions are often referred to as *kernels,* and calling one is termed *launching the kernel.* These functions require the caller to specify an "execution configuration" that includes the grid and block dimensions. The execution configuration can also include other information for the launch, such as the amount of additional shared memory to allocate and the stream where the kernel should execute. HIP introduces a standard C++ calling convention to pass the execution configuration to the kernel (this convention replaces the Cuda <<< >>> syntax). In HIP, -- Kernels launch with the "hipLaunchKernel" function +`__global__` functions are often referred to as *kernels,* and calling one is termed *launching the kernel.* These functions require the caller to specify an "execution configuration" that includes the grid and block dimensions. The execution configuration can also include other information for the launch, such as the amount of additional shared memory to allocate and the stream where the kernel should execute. HIP introduces a standard C++ calling convention to pass the execution configuration to the kernel in addition to the Cuda <<< >>> syntax. In HIP, +- Kernels launch with either <<< >>> syntax or the "hipLaunchKernel" function - The first five parameters to hipLaunchKernel are the following: - **symbol kernelName**: the name of the kernel to launch. To support template kernels which contains "," use the HIP_KERNEL_NAME macro. The hipify tools insert this automatically. - **dim3 gridDim**: 3D-grid dimensions specifying the number of blocks to launch. @@ -116,12 +116,13 @@ __global__ MyKernel(hipLaunchParm lp, float *A, float *B, float *C, size_t N) ... } -// Replace MyKernel<<>> (a,b,c,n); -hipLaunchKernel(MyKernel, dim3(gridDim), dim3(groupDim), 0/*dynamicShared*/, 0/*stream), a, b, c, n); +MyKernel<<>> (a,b,c,n); +// Alternatively, kernel can be launched by +// hipLaunchKernel(MyKernel, dim3(gridDim), dim3(groupDim), 0/*dynamicShared*/, 0/*stream), a, b, c, n); ``` -The hipLaunchKernel macro always starts with the five parameters specified above, followed by the kernel arguments. The Hipify script automatically converts Cuda launch syntax to hipLaunchKernel, including conversion of optional arguments in <<< >>> to the five required hipLaunchKernel parameters. The dim3 constructor accepts zero to three arguments and will by default initialize unspecified dimensions to 1. See [dim3](#dim3). The kernel uses the coordinate built-ins (hipThread*, hipBlock*, hipGrid*) to determine coordinate index and coordinate bounds of the work item that’s currently executing. See [Coordinate Built-Ins](#coordinate-builtins). +The hipLaunchKernel macro always starts with the five parameters specified above, followed by the kernel arguments. The Hipify script optionally converts Cuda launch syntax to hipLaunchKernel, including conversion of optional arguments in <<< >>> to the five required hipLaunchKernel parameters. The dim3 constructor accepts zero to three arguments and will by default initialize unspecified dimensions to 1. See [dim3](#dim3). The kernel uses the coordinate built-ins (hipThread*, hipBlock*, hipGrid*) to determine coordinate index and coordinate bounds of the work item that’s currently executing. See [Coordinate Built-Ins](#coordinate-builtins). ## Kernel-Launch Example @@ -149,7 +150,9 @@ void callMyKernel() unsigned N = 1000000; const unsigned blockSize = 256; - hipLaunchKernel(MyKernel, dim3(N/blockSize), dim3(blockSize), 0, 0, a,b,c,N); + MyKernel<<>> (a,b,c,n); + // Alternatively, kernel can be launched by + // hipLaunchKernel(MyKernel, dim3(N/blockSize), dim3(blockSize), 0, 0, a,b,c,N); } ``` From e5e0efe002ed174578212fb959d9c933b201d5c6 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Mon, 2 Apr 2018 21:47:20 +0530 Subject: [PATCH 3/4] Fix texture driver api TRFS flags [ROCm/hip commit: 62fc6c85e02d0a1e932546f84085c0b6bb76239b] --- projects/hip/include/hip/hcc_detail/driver_types.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/hip/include/hip/hcc_detail/driver_types.h b/projects/hip/include/hip/hcc_detail/driver_types.h index a0d2e5bd77..c25053ec32 100644 --- a/projects/hip/include/hip/hcc_detail/driver_types.h +++ b/projects/hip/include/hip/hcc_detail/driver_types.h @@ -39,8 +39,8 @@ struct hipChannelFormatDesc { enum hipChannelFormatKind f; }; -#define HIP_TRSF_NORMALIZED_COORDINATES 0x02 -#define HIP_TRSF_READ_AS_INTEGER 0x01 +#define HIP_TRSF_NORMALIZED_COORDINATES 0x01 +#define HIP_TRSF_READ_AS_INTEGER 0x00 #define HIP_TRSA_OVERRIDE_FORMAT 0x01 enum hipArray_Format { From 2de854324a72fd24ee7e8d4e8e7f93113441691c Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Tue, 3 Apr 2018 14:04:33 +0530 Subject: [PATCH 4/4] [ci] Re-enable testing against hcc_1.7 [ROCm/hip commit: d9cb38c7642989bb9ab04f1b86e3b14663340309] --- projects/hip/Jenkinsfile | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/projects/hip/Jenkinsfile b/projects/hip/Jenkinsfile index cfbe6cd883..b110474329 100644 --- a/projects/hip/Jenkinsfile +++ b/projects/hip/Jenkinsfile @@ -443,6 +443,7 @@ hcc_1_6: // docker_clean_images( job_name, hip_image_name ) } }, +*/ hcc_1_7: { node('docker && rocm && dkms') @@ -478,7 +479,6 @@ hcc_1_7: // docker_clean_images( job_name, hip_image_name ) } }, -*/ nvcc: { node('docker && cuda')