Skip to content

Commit

Permalink
Reclaiming hipLaunchKernel : Incorporated review comments
Browse files Browse the repository at this point in the history
  • Loading branch information
Sarbojit Sarkar authored and Sarbojit Sarkar committed Aug 22, 2019
1 parent 02de329 commit f01587a
Showing 4 changed files with 65 additions and 82 deletions.
1 change: 0 additions & 1 deletion include/hip/hcc_detail/hip_runtime_api.h
View file
Original file line number Diff line number Diff line change
@@ -1538,7 +1538,6 @@ extern "C" {
*
* @returns #hipSuccess, #hipErrorInvalidValue, hipInvalidDevice
*
* @warning should be used from gcc compiled host code only
*/

hipError_t hipLaunchKernel(const void* function_address,
15 changes: 7 additions & 8 deletions src/hip_module.cpp
View file
Original file line number Diff line number Diff line change
@@ -1153,10 +1153,14 @@ hipError_t hipLaunchKernel(

hipFunction_t kd = hip_impl::get_program_state().kernel_descriptor((std::uintptr_t)func_addr,
hip_impl::target_agent(stream));
size_t szKernArg = 0;

if(kd != nullptr && kd->_header != nullptr)
szKernArg = kd->_header->kernarg_segment_byte_size;
if(kd == nullptr || kd->_header == nullptr)
return ihipLogStatus(hipErrorInvalidValue);

size_t szKernArg = kd->_header->kernarg_segment_byte_size;

if(args == NULL && szKernArg != 0)
return ihipLogStatus(hipErrorInvalidValue);

void* config[]{
HIP_LAUNCH_PARAM_BUFFER_POINTER,
@@ -1165,11 +1169,6 @@ hipError_t hipLaunchKernel(
&szKernArg,
HIP_LAUNCH_PARAM_END};

if(args == NULL)
{
config[1] = NULL;
}

return ihipLogStatus(ihipModuleLaunchKernel(tls, kd, numBlocks.x * dimBlocks.x, numBlocks.y * dimBlocks.y, numBlocks.z * dimBlocks.z,
dimBlocks.x, dimBlocks.y, dimBlocks.z, sharedMemBytes, stream, nullptr, (void**)&config, nullptr, nullptr, 0));
}
129 changes: 57 additions & 72 deletions tests/src/gcc/LaunchKernel.c
View file
Original file line number Diff line number Diff line change
@@ -28,56 +28,56 @@


#include "hip/hip_runtime.h"
#include "hip/hip_runtime_api.h"
#include <stdio.h>
#include "LaunchKernel.h"

void LaunchKernelArg()
#define HIPCHECK(error) \

This comment has been minimized.

Sign in to view

Copy link
@gargrahul

gargrahul Aug 22, 2019

Contributor

You don't need to define this , you can use https://github.com/ROCm-Developer-Tools/HIP/blob/master/tests/src/test_common.h#L72
{ \
hipError_t localError = error; \
if ((localError != hipSuccess) && (localError != hipErrorPeerAccessAlreadyEnabled)) { \
printf("%serror: '%s'(%d) from %s at %s:%d%s\n", "\x1B[31m", hipGetErrorString(localError), \
localError, #error, __FILE__, __LINE__, "\x1B[0m"); \
return false; \
} \
}


bool LaunchKernelArg()
{
hipError_t status = hipSuccess;
dim3 blocks = {1,1,1};
dim3 threads = {1,1,1};

printf(" -----: Launch Zero Argument Kernel :------ \n");
status = hipLaunchKernel(kernel, blocks, threads,NULL, 0, 0);
printf("hipLaunchKernel status %s\n",hipGetErrorString(status));
HIPCHECK(hipLaunchKernel(kernel, blocks, threads,NULL, 0, 0));

printf("Test Passed!!\n\n");
return true;
}

void LaunchKernelArg1()
bool LaunchKernelArg1()
{
hipError_t status = hipSuccess;
int A = 0;
int *A_d = NULL;
dim3 blocks = {1,1,1};
dim3 threads = {1,1,1};

printf(" -----: Launch One Argument Kernel :------ \n");

// Allocate Device memory
status = hipMalloc((void**)&A_d, sizeof(int));
printf("hipMalloc status %s\n", hipGetErrorString(status));
HIPCHECK(hipMalloc((void**)&A_d, sizeof(int)));

void* Args[]={A_d};
hipLaunchKernel(kernel1, blocks, threads, Args,0,0);
printf("hipLaunchKernel status %s\n",hipGetErrorString(status));
HIPCHECK(hipLaunchKernel(kernel1, blocks, threads, Args,0,0));

// Get the result back to host memory
status = hipMemcpy(&A, A_d, sizeof(int), hipMemcpyDeviceToHost);
printf("hipMalloc status %s\n", hipGetErrorString(status));

if(A == 333)
printf("Test Passed!!\n\n");
else
printf("Test Failed :Value of A is %d\n\n",A);
HIPCHECK(hipMemcpy(&A, A_d, sizeof(int), hipMemcpyDeviceToHost));

hipFree(A_d);

if(A != 333)
return false;

return true;
}

void LaunchKernelArg2()
bool LaunchKernelArg2()
{
hipError_t status = hipSuccess;
int A = 0;
int B = 123;
int *A_d = NULL;
@@ -86,40 +86,31 @@ void LaunchKernelArg2()
dim3 blocks = {1,1,1};
dim3 threads = {1,1,1};

printf(" -----: Launch Two Argument Kernel :------ \n");

// Allocate Device memory
status = hipMalloc((void**)&A_d, sizeof(int));
printf("hipMalloc status %s\n", hipGetErrorString(status));
HIPCHECK(hipMalloc((void**)&A_d, sizeof(int)));

status = hipMalloc((void**)&B_d, sizeof(int));
printf("hipMalloc status %s\n", hipGetErrorString(status));
HIPCHECK(hipMalloc((void**)&B_d, sizeof(int)));

// Copy data from host memory to device memory
status = hipMemcpy(B_d,&B, sizeof(int), hipMemcpyHostToDevice);
printf("hipMalloc status %s\n", hipGetErrorString(status));

HIPCHECK(hipMemcpy(B_d,&B, sizeof(int), hipMemcpyHostToDevice));

void* Args[]={A_d,B_d};
hipLaunchKernel(kernel2, blocks, threads, Args,0,0);
printf("hipLaunchKernel status %s\n", hipGetErrorString(status));
HIPCHECK(hipLaunchKernel(kernel2, blocks, threads, Args,0,0));

// Get the result back to host memory
status = hipMemcpy(&A, A_d, sizeof(int), hipMemcpyDeviceToHost);
printf("hipMalloc status %s\n", hipGetErrorString(status));

if(A == 123)
printf("Test Passed!! \n\n");
else
printf("Test Failed : Value of A is %d\n\n",A);
HIPCHECK(hipMemcpy(&A, A_d, sizeof(int), hipMemcpyDeviceToHost));

hipFree(A_d);

This comment has been minimized.

Sign in to view

Copy link
@gargrahul

gargrahul Aug 22, 2019

Contributor

Add HIPCHECK to hipFree too.
hipFree(B_d);

if(A != 123)
return false;

return true;
}

void LaunchKernelArg3()
bool LaunchKernelArg3()
{
hipError_t status = hipSuccess;
int A = 321;
int B = 123;
int C = 0;
@@ -130,51 +121,45 @@ void LaunchKernelArg3()
dim3 blocks = {1,1,1};
dim3 threads = {1,1,1};

printf(" -----: Launch Three Argument Kernel :------ \n");

// Allocate Device memory
status = hipMalloc((void**)&A_d, sizeof(int));
printf("hipMalloc status %s\n", hipGetErrorString(status));
HIPCHECK(hipMalloc((void**)&A_d, sizeof(int)));

status = hipMalloc((void**)&B_d, sizeof(int));
printf("hipMalloc status %s\n", hipGetErrorString(status));
HIPCHECK(hipMalloc((void**)&B_d, sizeof(int)));

status = hipMalloc((void**)&C_d, sizeof(int));
printf("hipMalloc status %s\n", hipGetErrorString(status));
HIPCHECK(hipMalloc((void**)&C_d, sizeof(int)));

// Copy data from host memory to device memory
status = hipMemcpy(A_d,&A, sizeof(int), hipMemcpyHostToDevice);
printf("hipMalloc status %s\n", hipGetErrorString(status));

status = hipMemcpy(B_d,&B, sizeof(int), hipMemcpyHostToDevice);
printf("hipMalloc status %s\n", hipGetErrorString(status));
HIPCHECK(hipMemcpy(A_d,&A, sizeof(int), hipMemcpyHostToDevice));

HIPCHECK(hipMemcpy(B_d,&B, sizeof(int), hipMemcpyHostToDevice));

void* Args[]={A_d,B_d,C_d};
hipLaunchKernel(kernel3, blocks, threads, Args,0,0);
printf("hipLaunchKernel status %s\n", hipGetErrorString(status));
HIPCHECK(hipLaunchKernel(kernel3, blocks, threads, Args,0,0));

// Get the result back to host memory
status = hipMemcpy(&C, C_d, sizeof(int), hipMemcpyDeviceToHost);
printf("hipMalloc status %s\n", hipGetErrorString(status));

if(C == 444)
printf("Test Passed\n\n");
else
printf("Test Failed : Value of C is %d\n\n",C);
HIPCHECK(hipMemcpy(&C, C_d, sizeof(int), hipMemcpyDeviceToHost));

hipFree(A_d);
hipFree(B_d);
hipFree(C_d);

if(C != 444)
return false;

return true;
}


int main()
{
LaunchKernelArg();
LaunchKernelArg1();
LaunchKernelArg2();
LaunchKernelArg3();

printf("PASSED!\n");
}
if( LaunchKernelArg() &&
LaunchKernelArg1() &&
LaunchKernelArg2() &&
LaunchKernelArg3())
{
printf("PASSED!\n");

This comment has been minimized.

Sign in to view

Copy link
@gargrahul

gargrahul Aug 22, 2019

Contributor

Please use https://github.com/ROCm-Developer-Tools/HIP/blob/master/tests/src/test_common.h#L52
}
else
printf("FAILED\n");
}
2 changes: 1 addition & 1 deletion tests/src/gcc/LaunchKernel.h
View file
Original file line number Diff line number Diff line change
@@ -1,5 +1,5 @@

/* Copyright (c) 2019-2020 Advanced Micro Devices, Inc. All rights reserved.
/* Copyright (c) 2019-Present Advanced Micro Devices, Inc. All rights reserved.
* Permission is hereby granted, free of charge, to any person obtaining a copy
* of this software and associated documentation files (the "Software"), to deal
* in the Software without restriction, including without limitation the rights

0 comments on commit f01587a

Please sign in to comment.