Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

add testing for SPIR-V 1.5 #2208

Open
wants to merge 3 commits into
base: main
Choose a base branch
from
Open
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
1 change: 1 addition & 0 deletions test_conformance/spirv_new/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -28,6 +28,7 @@ set(${MODULE_NAME}_SOURCES
test_op_vector_insert.cpp
test_op_vector_times_scalar.cpp
test_spirv_14.cpp
test_spirv_15.cpp
)

set(TEST_HARNESS_SOURCES
Expand Down
Original file line number Diff line number Diff line change
@@ -0,0 +1,37 @@
; SPIR-V
; Version: 1.5
; Reference:
; kernel void non_uniform_broadcast_dynamic_index_test(global uint* dst_base) {
; uint id = get_global_id(0);
; uint index = get_group_id(0);
svenvh marked this conversation as resolved.
Show resolved Hide resolved
; uint value = sub_group_non_uniform_broadcast(id, index);
; dst_base[id] = value;
; }
OpCapability Addresses
OpCapability Kernel
OpCapability GroupNonUniformBallot
OpMemoryModel Physical32 OpenCL
OpEntryPoint Kernel %kernel "non_uniform_broadcast_dynamic_index_test" %pglobalid %pgroupid
OpDecorate %pglobalid BuiltIn GlobalInvocationId
OpDecorate %pgroupid BuiltIn WorkgroupId
%uint = OpTypeInt 32 0
%sg_scope = OpConstant %uint 3
%uint3 = OpTypeVector %uint 3
%void = OpTypeVoid
%iptr_uint3 = OpTypePointer Input %uint3
%gptr_uint = OpTypePointer CrossWorkgroup %uint
%kernel_sig = OpTypeFunction %void %gptr_uint
%pglobalid = OpVariable %iptr_uint3 Input
%pgroupid = OpVariable %iptr_uint3 Input
%kernel = OpFunction %void None %kernel_sig
%dst_base = OpFunctionParameter %gptr_uint
%entry = OpLabel
%globalid = OpLoad %uint3 %pglobalid Aligned 32
%id = OpCompositeExtract %uint %globalid 0
%groupid = OpLoad %uint3 %pgroupid Aligned 32
%index = OpCompositeExtract %uint %groupid 0
%value = OpGroupNonUniformBroadcast %uint %sg_scope %id %index
%dst = OpInBoundsPtrAccessChain %gptr_uint %dst_base %id
OpStore %dst %value
OpReturn
OpFunctionEnd
Original file line number Diff line number Diff line change
@@ -0,0 +1,41 @@
; SPIR-V
; Version: 1.5
; Reference:
; kernel void non_uniform_broadcast_dynamic_index_test(global uint* dst_base) {
; uint id = get_global_id(0);
; uint index = get_group_id(0);
; uint value = sub_group_non_uniform_broadcast(id, index);
; dst_base[id] = value;
; }
OpCapability Addresses
OpCapability Kernel
OpCapability Int64
OpCapability GroupNonUniformBallot
OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %kernel "non_uniform_broadcast_dynamic_index_test" %pglobalid %pgroupid
OpDecorate %pglobalid BuiltIn GlobalInvocationId
OpDecorate %pgroupid BuiltIn WorkgroupId
%uint = OpTypeInt 32 0
%sg_scope = OpConstant %uint 3
%ulong = OpTypeInt 64 0
%ulong3 = OpTypeVector %ulong 3
%void = OpTypeVoid
%iptr_ulong3 = OpTypePointer Input %ulong3
%gptr_uint = OpTypePointer CrossWorkgroup %uint
%kernel_sig = OpTypeFunction %void %gptr_uint
%pglobalid = OpVariable %iptr_ulong3 Input
%pgroupid = OpVariable %iptr_ulong3 Input
%kernel = OpFunction %void None %kernel_sig
%dst_base = OpFunctionParameter %gptr_uint
%entry = OpLabel
%globalid = OpLoad %ulong3 %pglobalid Aligned 32
%globalid0 = OpCompositeExtract %ulong %globalid 0
%id = OpUConvert %uint %globalid0
%groupid = OpLoad %ulong3 %pgroupid Aligned 32
%groupid0 = OpCompositeExtract %ulong %groupid 0
%index = OpUConvert %uint %groupid0
%value = OpGroupNonUniformBroadcast %uint %sg_scope %id %index
%dst = OpInBoundsPtrAccessChain %gptr_uint %dst_base %globalid0
OpStore %dst %value
OpReturn
OpFunctionEnd
22 changes: 22 additions & 0 deletions test_conformance/spirv_new/spirv_asm/spv1.5/ptr_bitcast.spvasm32
Original file line number Diff line number Diff line change
@@ -0,0 +1,22 @@
; SPIR-V
; Version: 1.5
OpCapability Addresses
OpCapability Kernel
OpMemoryModel Physical32 OpenCL
OpEntryPoint Kernel %kernel "ptr_bitcast_test"
%uint = OpTypeInt 32 0
%void = OpTypeVoid
%pptr_int = OpTypePointer Function %uint
%gptr_uint = OpTypePointer CrossWorkgroup %uint
%kernel_sig = OpTypeFunction %void %gptr_uint %gptr_uint
%uint_42 = OpConstant %uint 42
%kernel = OpFunction %void None %kernel_sig
%dst_uint0 = OpFunctionParameter %gptr_uint
%dst_uint1 = OpFunctionParameter %gptr_uint
%entry = OpLabel
%pvalue = OpVariable %pptr_int Function %uint_42
%uint_ptr = OpBitcast %uint %pvalue
OpStore %dst_uint0 %uint_ptr
OpStore %dst_uint1 %uint_ptr
OpReturn
OpFunctionEnd
27 changes: 27 additions & 0 deletions test_conformance/spirv_new/spirv_asm/spv1.5/ptr_bitcast.spvasm64
Original file line number Diff line number Diff line change
@@ -0,0 +1,27 @@
; SPIR-V
; Version: 1.5
OpCapability Addresses
OpCapability Kernel
OpCapability Int64
OpMemoryModel Physical64 OpenCL
OpEntryPoint Kernel %kernel "ptr_bitcast_test"
%uint = OpTypeInt 32 0
%ulong = OpTypeInt 64 0
%uint2 = OpTypeVector %uint 2
%void = OpTypeVoid
%pptr_int = OpTypePointer Function %uint
%gptr_ulong = OpTypePointer CrossWorkgroup %ulong
%gptr_uint2 = OpTypePointer CrossWorkgroup %uint2
%kernel_sig = OpTypeFunction %void %gptr_ulong %gptr_uint2
%uint_42 = OpConstant %uint 42
%kernel = OpFunction %void None %kernel_sig
%dst_ulong = OpFunctionParameter %gptr_ulong
%dst_uint2 = OpFunctionParameter %gptr_uint2
%entry = OpLabel
%pvalue = OpVariable %pptr_int Function %uint_42
%ulong_ptr = OpBitcast %ulong %pvalue
OpStore %dst_ulong %ulong_ptr
%uint2_ptr = OpBitcast %uint2 %pvalue
OpStore %dst_uint2 %uint2_ptr
OpReturn
OpFunctionEnd
162 changes: 162 additions & 0 deletions test_conformance/spirv_new/test_spirv_15.cpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,162 @@
//
// Copyright (c) 2024 The Khronos Group Inc.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
//

#include "testBase.h"
#include "spirvInfo.hpp"
#include "types.hpp"

#include <algorithm>
#include <cinttypes>
#include <vector>

REGISTER_TEST(spirv15_ptr_bitcast)
{
if (!is_spirv_version_supported(device, "SPIR-V_1.5"))
{
log_info("SPIR-V 1.5 not supported; skipping tests.\n");
return TEST_SKIPPED_ITSELF;
}

cl_int error = CL_SUCCESS;

cl_uint address_bits;
error = clGetDeviceInfo(device, CL_DEVICE_ADDRESS_BITS, sizeof(cl_uint),
&address_bits, NULL);
SPIRV_CHECK_ERROR(error, "Failed to get address bits");

clProgramWrapper prog;
error = get_program_with_il(prog, device, context, "spv1.5/ptr_bitcast");
SPIRV_CHECK_ERROR(error, "Failed to compile spv program");

clKernelWrapper kernel = clCreateKernel(prog, "ptr_bitcast_test", &error);
SPIRV_CHECK_ERROR(error, "Failed to create spv kernel");

cl_ulong result_ulong =
address_bits == 32 ? 0xAAAAAAAAUL : 0xAAAAAAAAAAAAAAAAUL;
cl_ulong result_uint2 =
address_bits == 32 ? 0x55555555UL : 0x5555555555555555UL;

clMemWrapper dst_ulong =
clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,
sizeof(result_ulong), &result_ulong, &error);
SPIRV_CHECK_ERROR(error, "Failed to create dst_ulong buffer");

clMemWrapper dst_uint2 =
clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,
sizeof(result_uint2), &result_uint2, &error);
SPIRV_CHECK_ERROR(error, "Failed to create dst_uint2 buffer");

error |= clSetKernelArg(kernel, 0, sizeof(dst_ulong), &dst_ulong);
error |= clSetKernelArg(kernel, 1, sizeof(dst_uint2), &dst_uint2);
SPIRV_CHECK_ERROR(error, "Failed to set kernel args");

size_t global = 1;
error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &global, NULL, 0,
NULL, NULL);
SPIRV_CHECK_ERROR(error, "Failed to enqueue kernel");

error =
clEnqueueReadBuffer(queue, dst_ulong, CL_TRUE, 0, sizeof(result_ulong),
&result_ulong, 0, NULL, NULL);
SPIRV_CHECK_ERROR(error, "Unable to read dst_ulong buffer");

error =
clEnqueueReadBuffer(queue, dst_uint2, CL_TRUE, 0, sizeof(result_uint2),
&result_uint2, 0, NULL, NULL);
SPIRV_CHECK_ERROR(error, "Unable to read dst_uint2 buffer");

if (result_ulong != result_uint2)
{
log_error("Results mismatch! ulong = 0x%016" PRIx64
" vs. uint2 = 0x%016" PRIx64 "\n",
result_ulong, result_uint2);
return TEST_FAIL;
}

return TEST_PASS;
}

REGISTER_TEST(spirv15_non_uniform_broadcast)
{
if (!is_spirv_version_supported(device, "SPIR-V_1.5"))
{
log_info("SPIR-V 1.5 not supported; skipping tests.\n");
return TEST_SKIPPED_ITSELF;
}

if (!is_extension_available(device, "cl_khr_subgroup_ballot"))
{
log_info("cl_khr_subgroup_ballot is not supported; skipping tests.\n");
return TEST_SKIPPED_ITSELF;
}

cl_int error = CL_SUCCESS;

clProgramWrapper prog;
error = get_program_with_il(prog, device, context,
"spv1.5/non_uniform_broadcast_dynamic_index");
SPIRV_CHECK_ERROR(error, "Failed to compile spv program");

clKernelWrapper kernel = clCreateKernel(
prog, "non_uniform_broadcast_dynamic_index_test", &error);
SPIRV_CHECK_ERROR(error, "Failed to create spv kernel");

// Get the local work-group size for one sub-group per work-group.
size_t lws = 0;
size_t one = 1;
error = clGetKernelSubGroupInfo(
kernel, device, CL_KERNEL_LOCAL_SIZE_FOR_SUB_GROUP_COUNT,
sizeof(size_t), &one, sizeof(size_t), &lws, NULL);
SPIRV_CHECK_ERROR(error, "Failed to get local work size for one sub-group");

// Use four work-groups, unless the local-group size is less than four.
size_t wgcount = std::min<size_t>(lws, 4);
bashbaug marked this conversation as resolved.
Show resolved Hide resolved
size_t gws = wgcount * lws;
clMemWrapper dst = clCreateBuffer(context, CL_MEM_READ_WRITE,
sizeof(cl_int) * gws, NULL, &error);
SPIRV_CHECK_ERROR(error, "Failed to create dst buffer");

error |= clSetKernelArg(kernel, 0, sizeof(dst), &dst);
SPIRV_CHECK_ERROR(error, "Failed to set kernel args");

error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &gws, &lws, 0, NULL,
NULL);
SPIRV_CHECK_ERROR(error, "Failed to enqueue kernel");

std::vector<cl_int> results(gws);
bashbaug marked this conversation as resolved.
Show resolved Hide resolved
error = clEnqueueReadBuffer(queue, dst, CL_TRUE, 0, sizeof(cl_int) * gws,
results.data(), 0, NULL, NULL);
SPIRV_CHECK_ERROR(error, "Unable to read destination buffer");

// Remember: the test kernel did:
// sub_group_non_uniform_broadcast(get_global_id(0), get_group_id(0))
for (size_t g = 0; g < wgcount; g++)
{
for (size_t l = 0; l < lws; l++)
{
size_t index = g * lws + l;
size_t check = g * lws + g;
if (results[index] != static_cast<cl_int>(check))
{
log_error("Result mismatch at index %zu! Got %d, Wanted %zu\n",
index, results[index], check);
return TEST_FAIL;
}
}
}

return TEST_PASS;
}
Loading