Siu Chi Chan
bff8e15e13
* Fix bug in LaunchKernel test Instead of passing the address of the gpu buffer, pass the address of the pointer that holds the address of the gpu buffer * Fix hipLaunchKernel's kernarg buffer construction. The hipLaunchKernel implementation should rely on ihipModuleLaunchKernel to construct the kernarg buffer correctly based on kernel metadata. * Fix a bug in get_functions where the Kernel_descriptor wasn't constructed with the correct kernarg layout information. * Fix a bug in kernarg layout parsing dealing with kernel without any arg * teach ihipModuleLaunchKernel to handle kernel without any arg * Add a more interesting test
53 righe
1.5 KiB
C++
53 righe
1.5 KiB
C++
/*
|
|
* 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
|
|
* to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
|
|
* copies of the Software, and to permit persons to whom the Software is
|
|
* furnished to do so, subject to the following conditions:
|
|
* The above copyright notice and this permission notice shall be included in
|
|
* all copies or substantial portions of the Software.
|
|
* THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANNTY OF ANY KIND, EXPRESS OR
|
|
* IMPLIED, INNCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
|
* FITNNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
|
* AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANNY CLAIM, DAMAGES OR OTHER
|
|
* LIABILITY, WHETHER INN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
|
* OUT OF OR INN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
|
* THE SOFTWARE.
|
|
* */
|
|
|
|
|
|
#include<hip/hip_runtime.h>
|
|
#include "LaunchKernel.h"
|
|
|
|
extern "C"
|
|
{
|
|
|
|
__global__ void kernel()
|
|
{
|
|
int a = 20;
|
|
}
|
|
|
|
__global__ void kernel1(int *a)
|
|
{
|
|
*a = 333;
|
|
}
|
|
|
|
__global__ void kernel2(int *a, int*b)
|
|
{
|
|
*a = *b;
|
|
}
|
|
|
|
__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"
|