Initial Commit
Contributors:
Ammar ELWazir <aelwazir@amd.com>
AravindanC <aravindan.cheruvally@amd.com>
Benjamin Welton <bewelton@amd.com>
Ma, Bing <Bing.Ma@amd.com>
Chun Yang <chun.yang@amd.com>
Cole Nelson <cole.nelson@amd.com>
Ethan Stewart <ethan.stewart@amd.com>
Evgeny <evgeny.shcherbakov@amd.com>
Freddy Paul <Freddy.paul@amd.com>
Giovanni Baraldi <gbaraldi@amd.com>
Gopesh Bhardwaj <Gopesh.Bhardwaj@amd.com>
Icarus Sparry <icarus.sparry@amd.com>
itrowbri <Ian.Trowbridge@amd.com>
James Edwards <JamesAdrian.Edwards@amd.com>
jatang <jatang@amd.com>
Jeremy Newton <Jeremy.Newton@amd.com>
Jonathan Kim <jonathan.kim@amd.com>
Kent Russell <kent.russell@amd.com>
Kiumars Sabeti <kiumars.sabeti@amd.com>
Lang Yu <lang.yu@amd.com>
Laurent Morichetti <laurent.morichetti@amd.com>
Mallya, Ameya Keshava <AmeyaKeshava.Mallya@amd.com>
Manjunath Jakaraddi <manjunath.jakaraddi@amd.com>
Mark Laws <markdavid.laws@amd.com>
Mohan Kumar Mithur <Mohan.KumarMithur@amd.com>
Nicholas Curtis <nicurtis@amd.com>
Nirmal Unnikrishnan <Nirmal.Unnikrishnan@amd.com>
Parag Bhandari <parag.bhandari@amd.com>
Ranjith Ramakrishnan <Ranjith.Ramakrishnan@amd.com>
Robert Gregory <Robert.Gregory@amd.com>
Saravanan Solaiyappan <saravanan.solaiyappan@amd.com>
Saurabh Verma <saurabh.verma@amd.com>
Srihari Uttanur <srihari.u@amd.com>
Srinivasan Subramanian <srinivasan.subramanian@amd.com>
Sriraksha Nagaraj <Sriraksha.Nagaraj@amd.com>
Sushma Vaddireddy <svaddire@amd.com>
Xianwei Zhang <Xianwei.Zhang@amd.com>
[ROCm/aqlprofile commit: 1ed169e30c]
This commit is contained in:
Evgeny
Ammar ELWazir
Binary file not shown.
@@ -0,0 +1,56 @@
|
||||
|
||||
|
||||
/**
|
||||
* simple_convolution is where each pixel of the output image
|
||||
* is the weighted sum of the neighborhood pixels of the input image
|
||||
* The neighborhood is defined by the dimensions of the mask and
|
||||
* weight of each neighbor is defined by the mask itself.
|
||||
* @param output Output matrix after performing convolution
|
||||
* @param input Input matrix on which convolution is to be performed
|
||||
* @param mask mask matrix using which convolution was to be performed
|
||||
* @param inputDimensions dimensions of the input matrix
|
||||
* @param maskDimensions dimensions of the mask matrix
|
||||
*/
|
||||
__kernel void simple_convolution(__global uint * output,
|
||||
__global uint * input,
|
||||
__global float * mask,
|
||||
const uint2 inputDimensions,
|
||||
const uint2 maskDimensions) {
|
||||
|
||||
uint tid = get_global_id(0);
|
||||
|
||||
uint width = inputDimensions.x;
|
||||
uint height = inputDimensions.y;
|
||||
|
||||
uint x = tid%width;
|
||||
uint y = tid/width;
|
||||
|
||||
uint maskWidth = maskDimensions.x;
|
||||
uint maskHeight = maskDimensions.y;
|
||||
|
||||
uint vstep = (maskWidth -1)/2;
|
||||
uint hstep = (maskHeight -1)/2;
|
||||
|
||||
// find the left, right, top and bottom indices such that
|
||||
// the indices do not go beyond image boundaires
|
||||
uint left = (x < vstep) ? 0 : (x - vstep);
|
||||
uint right = ((x + vstep) >= width) ? width - 1 : (x + vstep);
|
||||
uint top = (y < hstep) ? 0 : (y - hstep);
|
||||
uint bottom = ((y + hstep) >= height)? height - 1: (y + hstep);
|
||||
|
||||
// initializing wighted sum value
|
||||
float sumFX = 0;
|
||||
|
||||
for(uint i = left; i <= right; ++i) {
|
||||
for(uint j = top; j <= bottom; ++j) {
|
||||
// performing wighted sum within the mask boundaries
|
||||
uint maskIndex = (j - (y - hstep)) * maskWidth + (i - (x - vstep));
|
||||
uint index = j * width + i;
|
||||
sumFX += ((float)input[index] * mask[maskIndex]);
|
||||
}
|
||||
}
|
||||
|
||||
// To round to the nearest integer
|
||||
sumFX += 0.5f;
|
||||
output[tid] = (uint)sumFX;
|
||||
}
|
||||
@@ -0,0 +1,163 @@
|
||||
// MIT License
|
||||
//
|
||||
// Copyright (c) 2017-2025 Advanced Micro Devices, Inc.
|
||||
//
|
||||
// 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 WARRANTY OF ANY KIND, EXPRESS OR
|
||||
// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
// THE SOFTWARE.
|
||||
|
||||
|
||||
#include "simple_convolution/simple_convolution.h"
|
||||
|
||||
#include <assert.h>
|
||||
#include <string.h>
|
||||
#include <random>
|
||||
#include <iostream>
|
||||
|
||||
#include "util/helper_funcs.h"
|
||||
#include "util/test_assert.h"
|
||||
|
||||
#define IMAGE_SIZE 128
|
||||
#define MASK_SIZE 3
|
||||
|
||||
std::vector<uint32_t> simple_convolution::get_input_data(size_t width_, size_t height_) {
|
||||
std::random_device dev;
|
||||
std::mt19937 rng(dev());
|
||||
// Low range to avoid floating point issues
|
||||
std::uniform_int_distribution<uint32_t> rsampler(0, INT8_MAX);
|
||||
|
||||
std::vector<uint32_t> data;
|
||||
data.reserve(width_ * height_);
|
||||
|
||||
for (size_t idy = 0; idy < height_; idy++)
|
||||
for (size_t idx = 0; idx < width_; idx++) data.push_back(rsampler(rng));
|
||||
|
||||
return data;
|
||||
}
|
||||
|
||||
simple_convolution::simple_convolution()
|
||||
: width_(IMAGE_SIZE), height_(IMAGE_SIZE), mask_width_(MASK_SIZE), mask_height_(MASK_SIZE) {
|
||||
assert(IsPowerOf2(width_));
|
||||
assert(IsPowerOf2(height_));
|
||||
// Ensure symmetric mask (odd number)
|
||||
assert(mask_width_ % 2 == 1);
|
||||
assert(mask_height_ % 2 == 1);
|
||||
|
||||
const uint32_t input_size_bytes = width_ * height_ * sizeof(uint32_t);
|
||||
const uint32_t mask_size_bytes =
|
||||
static_cast<uint64_t>(mask_width_) * mask_height_ * sizeof(float);
|
||||
|
||||
SetInDescr(KERNARG_BUF_ID, KERNARG_DES_ID, sizeof(kernel_args_t));
|
||||
SetInDescr(INPUT_BUF_ID, SYS_DES_ID, input_size_bytes);
|
||||
SetInDescr(MASK_BUF_ID, SYS_DES_ID, mask_size_bytes);
|
||||
SetOutDescr(LOCAL_BUF_ID, LOCAL_DES_ID, input_size_bytes);
|
||||
SetHostDescr(REFOUT_BUF_ID, REFOUT_DES_ID, input_size_bytes);
|
||||
|
||||
input_data_ = get_input_data(width_, height_);
|
||||
TEST_ASSERT(input_data_.size() == static_cast<uint64_t>(width_) * height_);
|
||||
}
|
||||
|
||||
void simple_convolution::Init() {
|
||||
std::clog << "simple_convolution::init :" << std::endl;
|
||||
|
||||
mem_descr_t kernarg_des = GetDescr(KERNARG_BUF_ID);
|
||||
mem_descr_t input_des = GetDescr(INPUT_BUF_ID);
|
||||
mem_descr_t mask_des = GetDescr(MASK_BUF_ID);
|
||||
mem_descr_t output_des = GetDescr(LOCAL_BUF_ID);
|
||||
#if 0
|
||||
printf("kernarg_des %p 0x%x\n", kernarg_des.ptr, kernarg_des.size);
|
||||
printf("input_des %p 0x%x\n", input_des.ptr, input_des.size);
|
||||
printf("mask_des %p 0x%x\n", mask_des.ptr, mask_des.size);
|
||||
printf("output_des %p 0x%x\n", output_des.ptr, output_des.size);
|
||||
#endif
|
||||
uint32_t* input = reinterpret_cast<uint32_t*>(input_des.ptr);
|
||||
uint32_t* output_local = reinterpret_cast<uint32_t*>(output_des.ptr);
|
||||
float* mask = reinterpret_cast<float*>(mask_des.ptr);
|
||||
kernel_args_t* kernel_args = reinterpret_cast<kernel_args_t*>(kernarg_des.ptr);
|
||||
|
||||
// Fill a blurr filter or some other filter of your choice
|
||||
const float val = 1.0f / (mask_width_ * 2.0f - 1.0f);
|
||||
for (uint32_t i = 0; i < (mask_width_ * mask_height_); i++) {
|
||||
mask[i] = 0;
|
||||
}
|
||||
for (uint32_t i = 0; i < mask_width_; i++) {
|
||||
uint32_t y = mask_height_ / 2;
|
||||
mask[y * mask_width_ + i] = val;
|
||||
}
|
||||
for (uint32_t i = 0; i < mask_height_; i++) {
|
||||
uint32_t x = mask_width_ / 2;
|
||||
mask[i * mask_width_ + x] = val;
|
||||
}
|
||||
|
||||
// Print the INPUT array.
|
||||
std::clog << std::dec;
|
||||
PrintArray<uint32_t>("> Input[0]", input, width_, 1);
|
||||
PrintArray<float>("> Mask", mask, mask_width_, mask_height_);
|
||||
|
||||
// Fill the kernel args
|
||||
kernel_args->arg1 = output_local;
|
||||
kernel_args->arg2 = input;
|
||||
kernel_args->arg3 = mask;
|
||||
kernel_args->arg4 = width_;
|
||||
kernel_args->arg41 = height_;
|
||||
kernel_args->arg5 = mask_width_;
|
||||
kernel_args->arg51 = mask_height_;
|
||||
|
||||
// Calculate the reference output
|
||||
ReferenceImplementation(reinterpret_cast<uint32_t*>(GetRefOut()), input, mask, width_, height_,
|
||||
mask_width_, mask_height_);
|
||||
}
|
||||
|
||||
void simple_convolution::PrintOutput(const void* ptr) const {
|
||||
PrintArray<uint32_t>("> Output[0]", reinterpret_cast<const uint32_t*>(ptr), width_, 1);
|
||||
}
|
||||
|
||||
bool simple_convolution::ReferenceImplementation(uint32_t* output, const uint32_t* input,
|
||||
const float* mask, const uint32_t width,
|
||||
const uint32_t height, const uint32_t mask_width,
|
||||
const uint32_t mask_height) {
|
||||
const uint32_t vstep = (mask_width - 1) / 2;
|
||||
const uint32_t hstep = (mask_height - 1) / 2;
|
||||
|
||||
// for each pixel in the input
|
||||
for (uint32_t x = 0; x < width; x++) {
|
||||
for (uint32_t y = 0; y < height; y++) {
|
||||
// find the left, right, top and bottom indices such that
|
||||
// the indices do not go beyond image boundaires
|
||||
const uint32_t left = (x < vstep) ? 0 : (x - vstep);
|
||||
const uint32_t right = ((x + vstep) >= width) ? width - 1 : (x + vstep);
|
||||
const uint32_t top = (y < hstep) ? 0 : (y - hstep);
|
||||
const uint32_t bottom = ((y + hstep) >= height) ? height - 1 : (y + hstep);
|
||||
|
||||
// initializing wighted sum value
|
||||
float sum_fx = 0;
|
||||
for (uint32_t i = left; i <= right; ++i) {
|
||||
for (uint32_t j = top; j <= bottom; ++j) {
|
||||
// performing wighted sum within the mask boundaries
|
||||
uint32_t mask_idx = (j - (y - hstep)) * mask_width + (i - (x - vstep));
|
||||
uint32_t index = j * width + i;
|
||||
|
||||
// to round to the nearest integer
|
||||
sum_fx += ((float)input[index] * mask[mask_idx]);
|
||||
}
|
||||
}
|
||||
output[y * width + x] = uint32_t(sum_fx + 0.5f);
|
||||
}
|
||||
}
|
||||
|
||||
return true;
|
||||
}
|
||||
@@ -0,0 +1,94 @@
|
||||
// MIT License
|
||||
//
|
||||
// Copyright (c) 2017-2025 Advanced Micro Devices, Inc.
|
||||
//
|
||||
// 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 WARRANTY OF ANY KIND, EXPRESS OR
|
||||
// IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
|
||||
// FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
|
||||
// AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
|
||||
// LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
|
||||
// OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
|
||||
// THE SOFTWARE.
|
||||
|
||||
|
||||
#ifndef TEST_SIMPLE_CONVOLUTION_SIMPLE_CONVOLUTION_H_
|
||||
#define TEST_SIMPLE_CONVOLUTION_SIMPLE_CONVOLUTION_H_
|
||||
|
||||
#include <map>
|
||||
#include <string>
|
||||
#include <vector>
|
||||
|
||||
#include "ctrl/test_kernel.h"
|
||||
|
||||
// Class implements simple_convolution kernel parameters
|
||||
class simple_convolution : public TestKernel {
|
||||
public:
|
||||
// Kernel buffers IDs
|
||||
enum { INPUT_BUF_ID, LOCAL_BUF_ID, MASK_BUF_ID, KERNARG_BUF_ID, REFOUT_BUF_ID };
|
||||
|
||||
// Constructor
|
||||
simple_convolution();
|
||||
|
||||
// Initialize method
|
||||
void Init();
|
||||
|
||||
// Return compute grid size
|
||||
uint32_t GetGridSize() const { return width_ * height_; }
|
||||
|
||||
// Print output
|
||||
void PrintOutput(const void* ptr) const;
|
||||
|
||||
// Return name
|
||||
std::string Name() const { return std::string("simple_convolution"); }
|
||||
|
||||
private:
|
||||
// Local kernel arguments declaration
|
||||
struct kernel_args_t {
|
||||
void* arg1;
|
||||
void* arg2;
|
||||
void* arg3;
|
||||
uint32_t arg4;
|
||||
uint32_t arg41;
|
||||
uint32_t arg5;
|
||||
uint32_t arg51;
|
||||
};
|
||||
|
||||
// Reference CPU implementation of Simple Convolution
|
||||
// @param output Output matrix after performing convolution
|
||||
// @param input Input matrix on which convolution is to be performed
|
||||
// @param mask mask matrix using which convolution was to be performed
|
||||
// @param input_dimensions dimensions of the input matrix
|
||||
// @param mask_dimensions dimensions of the mask matrix
|
||||
// @return bool true on success and false on failure
|
||||
bool ReferenceImplementation(uint32_t* output, const uint32_t* input, const float* mask,
|
||||
const uint32_t width, const uint32_t height,
|
||||
const uint32_t maskWidth, const uint32_t maskHeight);
|
||||
|
||||
// Width of the Input array
|
||||
uint32_t width_;
|
||||
|
||||
// Height of the Input array
|
||||
uint32_t height_;
|
||||
|
||||
// Mask dimensions
|
||||
uint32_t mask_width_;
|
||||
|
||||
// Mask dimensions
|
||||
uint32_t mask_height_;
|
||||
|
||||
// Input data
|
||||
std::vector<uint32_t> input_data_;
|
||||
static std::vector<uint32_t> get_input_data(size_t width, size_t height);
|
||||
};
|
||||
|
||||
#endif // TEST_SIMPLE_CONVOLUTION_SIMPLE_CONVOLUTION_H_
|
||||
@@ -0,0 +1,154 @@
|
||||
module &m:1:0:$full:$large:$default;
|
||||
extension "amd:gcn";
|
||||
extension "IMAGE";
|
||||
|
||||
decl prog function &abort()();
|
||||
|
||||
prog kernel &__OpenCL_simple_convolution(kernarg_u64 %__global_offset_0,
|
||||
kernarg_u64 %output,
|
||||
kernarg_u64 %input,
|
||||
kernarg_u64 %mask,
|
||||
kernarg_u32 %inputDimensions[2],
|
||||
kernarg_u32 %maskDimensions[2]) {
|
||||
|
||||
pragma "AMD RTI", "ARGSTART:__OpenCL_simple_convolution";
|
||||
pragma "AMD RTI", "version:3:1:104";
|
||||
pragma "AMD RTI", "device:generic";
|
||||
pragma "AMD RTI", "uniqueid:1024";
|
||||
pragma "AMD RTI", "memory:private:0";
|
||||
pragma "AMD RTI", "memory:region:0";
|
||||
pragma "AMD RTI", "memory:local:0";
|
||||
pragma "AMD RTI", "value:__global_offset_0:u64:1:1:0";
|
||||
pragma "AMD RTI", "pointer:output:u32:1:1:96:uav:7:4:RW:0:0:0";
|
||||
pragma "AMD RTI", "pointer:input:u32:1:1:112:uav:7:4:RW:0:0:0";
|
||||
pragma "AMD RTI", "pointer:mask:float:1:1:128:uav:7:4:RW:0:0:0";
|
||||
pragma "AMD RTI", "value:inputDimensions:u32:2:1:144";
|
||||
pragma "AMD RTI", "constarg:4:inputDimensions";
|
||||
pragma "AMD RTI", "value:maskDimensions:u32:2:1:160";
|
||||
pragma "AMD RTI", "constarg:5:maskDimensions";
|
||||
pragma "AMD RTI", "function:1:0";
|
||||
pragma "AMD RTI", "memory:64bitABI";
|
||||
pragma "AMD RTI", "privateid:8";
|
||||
pragma "AMD RTI", "enqueue_kernel:0";
|
||||
pragma "AMD RTI", "kernel_index:0";
|
||||
pragma "AMD RTI", "reflection:0:size_t";
|
||||
pragma "AMD RTI", "reflection:1:uint*";
|
||||
pragma "AMD RTI", "reflection:2:uint*";
|
||||
pragma "AMD RTI", "reflection:3:float*";
|
||||
pragma "AMD RTI", "reflection:4:uint2";
|
||||
pragma "AMD RTI", "reflection:5:uint2";
|
||||
pragma "AMD RTI", "ARGEND:__OpenCL_simple_convolution";
|
||||
|
||||
@__OpenCL_simple_convolution_Entry:
|
||||
|
||||
// BB#0: // %entry
|
||||
|
||||
workitemabsid_u32 $s6, 0;
|
||||
cvt_u64_u32 $d0, $s6;
|
||||
ld_kernarg_align(8)_width(all)_u64 $d4, [%__global_offset_0];
|
||||
add_u64 $d0, $d0, $d4;
|
||||
cvt_u32_u64 $s5, $d0;
|
||||
ld_v2_kernarg_align(4)_width(all)_u32 ($s0, $s4), [%inputDimensions];
|
||||
ld_v2_kernarg_align(4)_width(all)_u32 ($s1, $s9), [%maskDimensions];
|
||||
rem_u32 $s7, $s5, $s0;
|
||||
add_u32 $s2, $s1, 4294967295;
|
||||
shr_u32 $s8, $s2, 1;
|
||||
add_u32 $s2, $s7, $s8;
|
||||
add_u32 $s3, $s0, 4294967295;
|
||||
cmp_ge_b1_u32 $c0, $s2, $s0;
|
||||
cmov_b32 $s2, $c0, $s3, $s2;
|
||||
sub_u32 $s3, $s7, $s8;
|
||||
cmp_lt_b1_u32 $c0, $s7, $s8;
|
||||
cmov_b32 $s3, $c0, 0, $s3;
|
||||
ld_kernarg_align(8)_width(all)_u64 $d1, [%output];
|
||||
cmp_le_b1_u32 $c0, $s3, $s2;
|
||||
cbr_b1 $c0, @BB0_2;
|
||||
|
||||
// BB#1:
|
||||
|
||||
mov_b32 $s6, 0;
|
||||
br @BB0_6;
|
||||
|
||||
// @BB0_2: // %for.cond32.preheader.lr.ph
|
||||
|
||||
@BB0_2:
|
||||
|
||||
div_u32 $s5, $s5, $s0;
|
||||
add_u32 $s9, $s9, 4294967295;
|
||||
shr_u32 $s9, $s9, 1;
|
||||
add_u32 $s10, $s5, $s9;
|
||||
add_u32 $s11, $s4, 4294967295;
|
||||
cmp_ge_b1_u32 $c0, $s10, $s4;
|
||||
cmov_b32 $s4, $c0, $s11, $s10;
|
||||
sub_u32 $s10, $s5, $s9;
|
||||
cmp_lt_b1_u32 $c0, $s5, $s9;
|
||||
cmov_b32 $s5, $c0, 0, $s10;
|
||||
ld_kernarg_align(8)_width(all)_u64 $d2, [%mask];
|
||||
ld_kernarg_align(8)_width(all)_u64 $d3, [%input];
|
||||
cvt_u64_u32 $d5, $s6;
|
||||
add_u64 $d4, $d4, $d5;
|
||||
cvt_u32_u64 $s6, $d4;
|
||||
div_u32 $s6, $s6, $s0;
|
||||
max_u32 $s10, $s9, $s6;
|
||||
sub_u32 $s12, $s10, $s6;
|
||||
max_u32 $s11, $s7, $s8;
|
||||
mov_b32 $s6, 0;
|
||||
mad_u32 $s12, $s1, $s12, $s11;
|
||||
sub_u32 $s7, $s12, $s7;
|
||||
sub_u32 $s9, $s10, $s9;
|
||||
mad_u32 $s9, $s0, $s9, $s11;
|
||||
sub_u32 $s8, $s9, $s8;
|
||||
|
||||
// @BB0_3: // %for.cond32.preheader
|
||||
|
||||
@BB0_3:
|
||||
|
||||
cmp_gt_b1_u32 $c0, $s5, $s4;
|
||||
mov_b32 $s9, $s7;
|
||||
mov_b32 $s10, $s8;
|
||||
mov_b32 $s11, $s5;
|
||||
cbr_b1 $c0, @BB0_5;
|
||||
|
||||
// @BB0_4: // %for.body35
|
||||
|
||||
@BB0_4:
|
||||
|
||||
cvt_u64_u32 $d4, $s9;
|
||||
shl_u64 $d4, $d4, 2;
|
||||
add_u64 $d4, $d2, $d4;
|
||||
ld_global_align(4)_f32 $s12, [$d4];
|
||||
cvt_u64_u32 $d4, $s10;
|
||||
shl_u64 $d4, $d4, 2;
|
||||
add_u64 $d4, $d3, $d4;
|
||||
ld_global_align(4)_u32 $s13, [$d4];
|
||||
cvt_f32_u32 $s13, $s13;
|
||||
mul_ftz_f32 $s12, $s13, $s12;
|
||||
add_u32 $s9, $s9, $s1;
|
||||
add_u32 $s10, $s10, $s0;
|
||||
add_u32 $s11, $s11, 1;
|
||||
add_ftz_f32 $s6, $s6, $s12;
|
||||
cmp_le_b1_u32 $c0, $s11, $s4;
|
||||
cbr_b1 $c0, @BB0_4;
|
||||
|
||||
// @BB0_5: // %for.inc48
|
||||
|
||||
@BB0_5:
|
||||
|
||||
add_u32 $s7, $s7, 1;
|
||||
add_u32 $s8, $s8, 1;
|
||||
add_u32 $s3, $s3, 1;
|
||||
cmp_le_b1_u32 $c0, $s3, $s2;
|
||||
cbr_b1 $c0, @BB0_3;
|
||||
|
||||
// @BB0_6: // %for.end50
|
||||
|
||||
@BB0_6:
|
||||
|
||||
and_b64 $d0, $d0, 4294967295;
|
||||
shl_u64 $d0, $d0, 2;
|
||||
add_u64 $d0, $d1, $d0;
|
||||
add_ftz_f32 $s0, $s6, 0F3f000000;
|
||||
cvt_ftz_u32_f32 $s0, $s0;
|
||||
st_global_align(4)_u32 $s0, [$d0];
|
||||
ret;
|
||||
};
|
||||
Reference in New Issue
Block a user